5e77c58bd0
SWDEV-536813 - Make alignment 16 for special cases ( #609 )
2025-06-27 14:20:16 +05:30
a4631b5700
SWDEV-533232 Add num_threads API in cooperative_groups ( #437 )
...
* Add num_threads API in cooperative_groups
* Reimplementsize API as its alias.
* Update changelog
2025-06-26 15:05:17 +05:30
49690b8b38
Fix rocWMMA specialization ( #532 )
...
Co-authored-by: Aidan Belton <abeltons@amd.com >
2025-06-23 18:22:37 +05:30
3faece02dc
SWDEV-536813 - Fix HIP_vector_base alignment ( #501 )
...
* SWDEV-536813 - Fix HIP_vector_base alignment
Change-Id: I6ee211437768bbdf42186bcd0805b8c2290373da
* Remove alignas for vec3
Change-Id: If93c55327cad38cf179bdd2a89faa725b0a0fc15
* align array to match native_vector
Change-Id: I54fbdba01c19241a4c4945ba4c35ac6d40f17ea6
* update
Change-Id: Ib7de5a7ebde0b2a5c18fe5c70f7329a90807abfb
2025-06-23 18:21:51 +05:30
428fce9fd6
SWDEV-519446 Part 2: Update unsafe atomic function file header ( #394 )
...
* SWDEV-519446 Part 2: Update unsafe atomic function file header
* Use __HIP_ATOMICS_IGNORE_DENORMAL_MODE to denote [[clang::atomic(ignore_denormal_mode)]] when condition met.
2025-06-19 19:12:59 +05:30
adda2602e1
SWDEV-535738 - Properly define size_t for windows ( #481 )
2025-06-18 19:32:11 +05:30
e88d9fabe0
SWDEV-537656 - Added c++11 include guard for constexpr HIP_vector_type operator== ( #517 )
2025-06-13 12:05:41 +05:30
67f6409eb1
SWDEV-519446 - Part 1: update normal atomic function file header ( #345 )
...
Remove the scope of __has_builtin(__hip_atomic_compare_exchange_strong) == false
2025-06-12 09:16:51 +05:30
b7778ecd65
SWDEV-537139 - Add missing typename qualifier to permuteType ( #505 )
2025-06-11 08:03:10 +05:30
10ea744996
SWDEV-511020 - fix bugs in surface function header ( #449 )
2025-06-11 08:00:42 +05:30
543925e048
SWDEV-529899 - [ABI Break] remove unnamed namespace from fp16 operations ( #244 )
...
Change-Id: I4feb9089b4cee145300dc78fc43a4f42473f5b1c
2025-06-06 21:22:37 +05:30
a01134a1ab
SWDEV-536048 - Add explicit cast to int when warpSize is used with == operator ( #473 )
2025-06-06 12:22:31 +05:30
1bce89ca51
SWDEV-535738 - Use the type traits implementation from __hip_internal namespace in hip headers ( #471 )
2025-06-05 00:14:56 +05:30
c2f6c98006
Add HIP_vector_base constructor ( #456 )
...
Change-Id: I57571f6bf9876d3a9ff312b952858527ec6abc7b
2025-06-03 12:23:38 +05:30
54a2e943d7
SWDEV-535184 - Fix to switch HIP_RUNTIME_API_TABLE_STEP_VERSION 9 and 10 ( #448 )
2025-06-02 21:24:40 +05:30
a018d4278a
SWDEV-535738 - Use __hip_internal instead of std namespace for device code. ( #455 )
2025-06-02 20:09:17 +05:30
9533ceb8aa
Fix vec operator[] ( #419 )
...
Change-Id: I1b6767b8b7f7dbcf9998662dd9f8ae3a96e1deca
2025-05-29 13:09:34 +05:30
c0c3aa3637
SWDEV-508963 - Change hipMemcpy subset of methods to use const qualifier for src ( #139 ) ( #418 )
...
Change-Id: Ia5b68bacff634d3d9e5757f515bac50887895bc3
2025-05-29 13:07:51 +05:30
54503f0d67
SWDEV-530921 - Remove the usage of __AMDGCN_WAVEFRONT_SIZE as compile time constant ( #330 )
...
* SWDEV-530921 - Remove the usage of __AMDGCN_WAVEFRONT_SIZE as compile time constant
* wavefrontsize builtin to be used only in device compilation
* SWDEV-530921 - Remove the usage of __AMDGCN_WAVEFRONT_SIZE as compile time constant
2025-05-29 12:52:10 +05:30
0f46bd5131
Add operator[] to vec class ( #402 )
...
Change-Id: I452cd5d1b258ccbc354a2f424d062f0fa91100ea
2025-05-28 13:38:38 +05:30
ee2d4a0358
SWDEV-518962 - Match hipCtxGetApiVersion signature with Cuda ( #327 )
2025-05-26 20:28:22 +05:30
ad44dd527f
[CLR ABI] step runtime api version hip memset node params ( #373 )
...
* increase hip api table version for HIP_MEMSET_NODE_PARAMS removed.
* update version in static_assert.
---------
Co-authored-by: Venkateshwar Reddy Kandula <vkandula@amd.com >
2025-05-22 15:09:31 +05:30
946ea72c94
SWDEV-513341 - Remove datatypes from HIPRTC ( #377 )
...
Defining int64_t, uint64_t, int32_t, uint32_t in HIPRTC
seem to result in conflicts with some apps as they use
their own definitions for these types. NVRTC also doesn't
define these. Hence remove them and replace their usage
with hip internal datatypes such as __hip_int64_t
Change-Id: Id314730bbd61c05f126f915c546052f956550a3f
2025-05-22 13:28:31 +05:30
b27e2abc41
SWDEV-482125 - HIPRTC implementation of standard headers moved from std to __hip_internal namespace ( #376 )
...
Change-Id: Ia359ed45e0916ad397401113a69a6b26a590d807
2025-05-22 13:27:19 +05:30
3cd3b3ffc5
SWDEV-527851 SWDEV-527890 SWDEV-529456 - Make HIP vector simple struct ( #356 )
...
* Make HIP vector simple struct
Change-Id: I8442c2cc9af26b2a3c7d6719e3348df1593e83b3
* Update make_vector_type
Change-Id: Ic5060994a08baa4262c2a4b09fcbe6bc74276720
2025-05-21 21:30:55 +05:30
968f5599a8
SWDEV-508965 - [6.4 Preview] Remove HIP_MEMSET_NODE_PARAMS struct ( #121 ) ( #364 )
2025-05-21 21:16:42 +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
c50610b44d
Add __syncwarp operation ( #160 )
...
Change-Id: I6a3783beafdbb9f11a3b37333f4ff3f5be27ea54
2025-05-15 14:20:13 +05:30
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
2f73e1385b
SWDEV-525933 - add constexpr operators for fp16/bf16 ( #199 )
2025-05-09 09:53:58 +01:00
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
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
12febe6782
SWDEV-514560 - add fp6 header implementation ( #54 )
...
Co-authored-by: rahul manocha <rmanocha_amdeng>
2025-05-01 15:17:38 +01:00
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
99142c3dd9
SWDEV-526612 - Add missing copyright notices ( #201 )
2025-04-18 20:54:27 +05:30
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
588cf0fc69
SWDEV-520627 - include warp functions header for warpSize ( #177 )
...
Change-Id: Id3fff8f2722d521071ef0ff71b09fc365ef6fa82
2025-04-15 14:40:27 +05:30
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
6d10577761
SWDEV-521083 - Fix atomicMin/Max issues ( #151 )
...
Fix atomicMin/Max(), atomicMin/Max_system() issue on
float types.
2025-04-10 12:30:55 -04:00
628777b73d
SWDEV-461087 - fp4/fp6/fp8 ocp headers ( #41 )
...
This now has host conversions too, which is directly from Christopher's
work on fcbx.
Signed-off-by: Christopher M. Riedl
* add const to func parameter
* do not depend on builtins, use gfx950 detection
2025-04-10 17:22:15 +01:00
5c030840d6
SWDEV-520627 - include sync_warp header instead of warp function header ( #18 )
...
Change-Id: Ic3f54b0f5bfee8565a8bbb6218fb0ccdb900c9ea
2025-04-10 21:50:25 +05:30
5214d1ca07
SWDEV-525969 - add gfx950 to use ocp type for fp8 ( #157 )
...
* do not use __gfx94plus_clr__ macro in fp8 header
2025-04-09 16:21:39 +01:00
8c6b90996e
SWDEV-523281 - [clr] Implementation of hipLaunchKernelExC and hipDrvLaunchKernelEx API with support for cooperative launch ( #92 )
2025-04-03 20:10:05 +09:00
376f23b86a
SWDEV-516595 - Add __shfl functions with __hip_bfloat16 datatype ( #42 )
...
Also removes asserts in cooperative groups shfl functions since
__hip_bfloat16 shfl is present now
Change-Id: I57578b6e68dccc10c2ddcd194e9cc18bc7732ce1
2025-03-25 15:38:01 +05:30
340d6bb69f
SWDEV-420237 - Add __reduce_add_sync()
...
Change-Id: Ic8e4fab6b7aeb879d40b2c1419b30d1355a2bbdc
2025-03-12 03:20:49 -04:00
d91e1f19d0
SWDEV-512613 - Improve device atomics functions
...
Also part of SWDEV-510994.
1. Fix atomicMin/Max_system() for float and double.
2. Remove logics of gfx941 which isn't supported.
Change-Id: Iacfdc1bc13e8da2f5df8751bb315b37d33cea667
2025-03-06 10:05:59 -05:00
Belton-Schure, Aidan
Sicarov, Dragoslav
Sang, Tao
Assiouras, Ioannis
Brzak, Branislav
Chaudhary, Jatin Jaikishan
Dittakavi, Satyanvesh
Xie, Jiabao(Jimbo)
Patel, Jaydeepkumar
Betigeri, Sourabh
Kandula, Venkateshwar reddy
Jayaprakash, Karthik
Hernandez, Gerardo
Arsenault, Matthew
Milanov, Aleksandar