58df22546b
SWDEV-532868 - Disable implicit wait in hipFree for async allocations ( #323 )
...
* SWDEV-532868 - Disable implicit wait in hipFree for async allocations
* Fix compilation error
2025-05-21 21:14:38 +05:30
7a357800dd
SWDEV-508979 - Match hipModuleLoad negative return with Cuda ( #326 )
2025-05-21 21:09:07 +05:30
a46ed60e2e
SWDEV-508970 - Match hipBindTextureToArray negative return with Cuda ( #325 )
2025-05-21 20:57:09 +05:30
78f3caba0b
SWDEV-519346 - Fix __clock64() compiling issue in SPIRV ( #207 )
2025-05-20 16:57:17 -04:00
12131de4a9
SWDEV-529929 - hipMemGetHandleForAddressRange implementation. ( #245 )
2025-05-20 15:56:04 -04:00
f5b8db33f1
SWDEV-531711 - Report correct error code based on device failure. ( #286 )
2025-05-17 06:33:13 -04:00
c50610b44d
Add __syncwarp operation ( #160 )
...
Change-Id: I6a3783beafdbb9f11a3b37333f4ff3f5be27ea54
2025-05-15 14:20:13 +05:30
32eb6a5d89
SWDEV-530803 - User current device id while cloning graph node. ( #313 )
2025-05-15 09:06:15 +05:30
0457b634f8
SWDEV-527781 - Remove Stream Validation in HIP APIs
2025-05-13 13:45:27 -07:00
f7482ef0a6
SWDEV-529449 - Bug fix when retrieving a memobj from the IPC mem handle
2025-05-13 19:18:22 +01:00
5606debd8e
SWDEV-491314 - Re-enable cross-lane sync builtins ( #94 )
...
* Enables warp sync builtins by default
* Removes HIP_ENABLE_WARP_SYNC_BUILTINS; that macro will no longer have an effect. Instead, we will now be able to disable the builtins with the macro: HIP_DISABLE_WARP_SYNC_BUILTINS
2025-05-13 16:35:58 +01:00
f9199ac205
SWDEV-528683 - Hardcode valid wavefront compile time options ( #306 )
2025-05-12 19:29:39 +02:00
acb1f7e8d5
SWDEV-526855 - Modify the SIMDPerCU calculation for gfx1250/1. ( #275 )
2025-05-12 11:09:03 -04:00
c5ced8c3a2
SWDEV-512344 - Unmap all subbuffers ( #214 )
2025-05-12 16:56:10 +02:00
a7492c516d
SWDEV-511204 - Mapped virtual memory should use device instead of host context ( #213 )
...
Since the sub-buffer(virtual memory that is mapped to device memory) is associated with device memory, it should utilize the device context instead of the host context. The original implementation caused hipMemcpyPeer to not take the P2P path, as the memory object was treated as host memory.
2025-05-12 16:55:25 +02:00
b434fbe2bd
SWDEV-527299 - Support HIP_POINTER_ATTRIBUTE_CONTEXT ( #180 )
...
* SWDEV-527299 - Support HIP_POINTER_ATTRIBUTE_CONTEXT
As HIP enables UVA by default, it seems we can simply expose the context to support this feature.
2025-05-09 17:34:16 +08:00
2f73e1385b
SWDEV-525933 - add constexpr operators for fp16/bf16 ( #199 )
2025-05-09 09:53:58 +01:00
fa55557f46
SWDEV-493805 - Cleaning up launch parameters arguments. ( #241 )
2025-05-06 15:06:13 -04:00
607f8f26fd
SWDEV-529831 - Return error if the program is empty ( #257 )
2025-05-06 15:12:12 +05:30
a71c6eb1a0
SWDEV-529854 - __hmax/__hmin should handle nan's ( #246 )
2025-05-06 09:42:15 +01:00
b1ebf33850
SWDEV-529927 - add missing operations for fp16/bf16 ( #238 )
2025-05-06 09:41:21 +01:00
3775298655
Disable HIP_PLATFORM auto-detect if already defined ( #254 )
...
Co-authored-by: Stella Laurenzo <stellaraccident@gmail.com >
2025-05-05 15:37:53 -04:00
1db9a7d48b
SWDEV-1 - Stop using ocml rounding functions ( #228 )
...
Directly use the builtins. Use the elementwise versions since there's
no implied errno, regardless of -f[no]-math-errno.
I didn't change the cases unnecessarily casting. The bfloat and vector
cases should work directly.
2025-05-05 19:35:12 +02:00
cd9bc61559
Fix typos in warning msgs ( #231 )
2025-05-02 14:31:42 -07:00
12febe6782
SWDEV-514560 - add fp6 header implementation ( #54 )
...
Co-authored-by: rahul manocha <rmanocha_amdeng>
2025-05-01 15:17:38 +01:00
2538d7f02b
SWDEV-522841 - Graph nodes must be created/launched on device where they are captured/created ( #108 )
2025-04-29 22:20:39 +05:30
eb62fe9f62
SWDEV-522634 - Fix device properties in hipInfo ( #203 )
2025-04-29 11:29:47 -04:00
96cadbc9e9
SWDEV-520352 - Remove HostThread and legacy monitor ( #230 )
...
* SWDEV-520352 - Remove HostThread and legacy monitor
Remove HostThread, semaphore and legacy monitor.
Make original logics of thread and command queue stricker.
Add more comments to make logics clearer.
Some other minor improvement.
Also part of SWDEV-458943.
2025-04-29 09:55:24 -04:00
b2388dfb88
SWDEV-506467 - Skip Abort in case of crash from the device. ( #60 )
...
Change-Id: I964b2f2647d068202e9c38fcddb1337da754df8d
2025-04-29 11:19:02 +05:30
9cf3f1e461
SWDEV-528351 - Removes unused code and asserts to improve coverage ( #219 )
2025-04-28 14:40:35 -07:00
fb92683d86
SWDEV-469422 - Avoid using of hipStream_t in internal methods ( #69 )
...
Change-Id: Ifd5362f371c846a88241927383cb95cf046548ef
2025-04-28 15:09:11 +05:30
bbcb1f9c70
SWDEV-469422 - hipGraphNodeDOTAttribute change std::string members to const char* ( #70 )
...
Compiler creates global variables for every unique string
Change-Id: I4cf8dd3e763d16740096e345da67a7ef72f61515
2025-04-28 14:57:36 +05:30
1099e0a131
SWDEV-526188 - Fix race condition in StatCO::getStatFunc()
...
Make sure that a newly created FatBinaryInfo is assigned to modules only after extractFatBinary has been called for the object.
2025-04-27 21:14:01 +01:00
1113eff3f9
SWDEV-493275 - Support scratch limit ( #20 )
...
Support programmatic query and change of scratch limit on
AMD devices.
Change-Id: Id5da355a77366f97868e462847f3916e87fd2af6
2025-04-24 17:15:25 -04:00
e5ce544c45
SWDEV-469423 - hipStreamEndCapture graph* can be nullptr ( #170 )
2025-04-24 13:57:09 +05:30
27aad09bd4
SWDEV-518831 - fix streams' sync issue in mthreads ( #123 )
...
* SWDEV-518831 - fix streams' sync issue in mthreads
1. Fix sync issue of null stream and non-null streams in
multithreads.
2. Remove assert(GetSubmissionBatch() == nullptr) as it
is invalid in multithreads.
3. Update getActiveQueues() to deal with the state of
being terminated.
2025-04-23 15:08:07 -04:00
64d6f5714a
SWDEV-523281 - CHANGELOG.md and negative test return values : hipLaunchKernelEx, hipLaunchKernelExC, hipDrvLaunchKernelEx ( #155 )
2025-04-22 21:47:37 +05:30
1a8d766836
SWDEV-420237 - Fix reduce sync operations when masks are divergent ( #181 )
...
Do not use __ockl_activelane_u32() to calculate the index of the lane within the mask, as that would not work with divergent masks that have other bits on before the associated lane.
2025-04-22 19:47:58 +05:30
bf28bbd9ab
SWDEV-508538 - Optimize mem access and pack structure ( #71 )
...
Change-Id: Ib05b8891a6d228fc3266918a000d332fddc7438b
2025-04-21 13:43:25 +05:30
99142c3dd9
SWDEV-526612 - Add missing copyright notices ( #201 )
2025-04-18 20:54:27 +05:30
d00b2a0953
SWDEV-525423 - In COMGR Loader don't open file if image is already mapped ( #193 )
2025-04-16 11:00:54 +02:00
5fe080fd67
SWDEV-523137 - function ptrs should match across all devices ( #171 )
2025-04-16 10:35:48 +02:00
5d638d831c
SWDEV-512924 - add fp4 API ( #52 )
...
* Remove C-style include guard
* clean up issues in the PR
2025-04-15 17:53:50 +01:00
e92ea151b2
SWDEV-518317 - Remove Redundant Error Message in removeFatBinary ( #164 )
2025-04-15 09:00:39 -07:00
fcaefe97b8
SWDEV-509213 - make cmake_minimum_required consistent across clr ( #51 )
...
Change-Id: Ib0b1df7af8984a37d6bf7ca68ec99597d5978821
2025-04-15 15:23:41 +05:30
588cf0fc69
SWDEV-520627 - include warp functions header for warpSize ( #177 )
...
Change-Id: Id3fff8f2722d521071ef0ff71b09fc365ef6fa82
2025-04-15 14:40:27 +05:30
07e57a1f0d
SWDEV-517941 - use device bitcode before spirv ( #95 )
...
Also add flag: HIP_FORCE_SPIRV_CODEOBJECT to allow override to force use
SPIRV.
* use cache for already compiled code objects
* address review comments and use the two spirv isa names
2025-04-14 23:40:52 +01:00
c4fa3ef927
SWDEV-526208 - Fix miscalculation of coalesced tiled partition mask ( #162 )
2025-04-11 19:40:26 +02:00
66496258b4
SWDEV-521920 - Fix compilation issues introduced by the reduce sync operations - 2 ( #167 )
...
Fix pytorch 2.5 issues, by defining reduce sync operations for type __half in amd_hip_fp16.h and not in
amd_warp_sync_functions.h which is problematic in case __half does not get included before that header.
Only define types not supported by cuda if HIP_ENABLE_EXTRA_WARP_SYNC_TYPES is defined, to avoid portability issues
2025-04-11 17:00:59 +05:30
199b0f1086
Report null stream creation failure ( #152 )
...
Explicitly nulling the pointer causes us to report the error below
instead of keeping a dangling pointer around that will most likely lead
to a subsequent segfault.
2025-04-10 11:40:05 -07:00
Andryeyev, German
Brzak, Branislav
Sang, Tao
Jayaprakash, Karthik
Belton-Schure, Aidan
Patel, Jaydeepkumar
Xie, Pengda
Assiouras, Ioannis
Hernandez, Gerardo
Arandjelovic, Marko
Huang, AnZhong
Chaudhary, Jatin Jaikishan
Dittakavi, Satyanvesh
Guan, Zichuan
Arsenault, Matthew
Searles, Mark
Godavarthy Surya, Anusha
Jiang, Julia
Betigeri, Sourabh
GunaShekar, Ajay
Milanov, Aleksandar
Haehnle, Nicolai