AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
13,660 Commits 3 Branches 0 Tags
1aae121e864b7c19be35ccb4bfb393fddbeff0f7
Commit Graph

1785 Commits

Author SHA1 Message Date
Sang, Tao 1aae121e86 SWDEV-533964 - use __builtin_readcyclecounter() only in clock64 (#368) 
[ROCm/clr commit: 5fc6208b19]
 2025-07-22 10:41:23 +05:30
Jiang, Julia 005672c314 SWDEV-523483 - Update atomics add-operation in CLR (#751) 
[ROCm/clr commit: 9cd8757717]
 2025-07-19 06:48:54 +05:30
Dittakavi, Satyanvesh b6d69d162e SWDEV-542832 - Refactor vector types alignment (#713) 
[ROCm/clr commit: 918b0bdb7a]
 2025-07-15 21:09:31 +05:30
Belton-Schure, Aidan b735783ae7 SWDEV-536813 - Make alignment 16 for special cases (#609) 
[ROCm/clr commit: 5e77c58bd0]
 2025-06-27 14:20:16 +05:30
Sicarov, Dragoslav dbfd869936 SWDEV-533232 Add num_threads API in cooperative_groups (#437) 
* Add num_threads API in cooperative_groups
* Reimplementsize API as its alias.
* Update changelog

[ROCm/clr commit: a4631b5700]
 2025-06-26 15:05:17 +05:30
Belton-Schure, Aidan 90508da214 Fix rocWMMA specialization (#532) 
Co-authored-by: Aidan Belton <abeltons@amd.com>

[ROCm/clr commit: 49690b8b38]
 2025-06-23 18:22:37 +05:30
Belton-Schure, Aidan 56bee0426a 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

[ROCm/clr commit: 3faece02dc]
 2025-06-23 18:21:51 +05:30
Sang, Tao 0faa7d7c95 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.

[ROCm/clr commit: 428fce9fd6]
 2025-06-19 19:12:59 +05:30
Belton-Schure, Aidan 615b904729 SWDEV-535738 - Properly define size_t for windows (#481) 
[ROCm/clr commit: adda2602e1]
 2025-06-18 19:32:11 +05:30
Assiouras, Ioannis 079b490e64 SWDEV-537656 - Added c++11 include guard for constexpr HIP_vector_type operator== (#517) 
[ROCm/clr commit: e88d9fabe0]
 2025-06-13 12:05:41 +05:30
Sang, Tao cca39a7e19 SWDEV-519446 - Part 1: update normal atomic function file header (#345) 
Remove the scope of __has_builtin(__hip_atomic_compare_exchange_strong) == false

[ROCm/clr commit: 67f6409eb1]
 2025-06-12 09:16:51 +05:30
Brzak, Branislav cefb51e63e SWDEV-537139 - Add missing typename qualifier to permuteType (#505) 
[ROCm/clr commit: b7778ecd65]
 2025-06-11 08:03:10 +05:30
Sang, Tao 06616a4837 SWDEV-511020 - fix bugs in surface function header (#449) 
[ROCm/clr commit: 10ea744996]
 2025-06-11 08:00:42 +05:30
Chaudhary, Jatin Jaikishan c914642086 SWDEV-529899 - [ABI Break] remove unnamed namespace from fp16 operations (#244) 
Change-Id: I4feb9089b4cee145300dc78fc43a4f42473f5b1c

[ROCm/clr commit: 543925e048]
 2025-06-06 21:22:37 +05:30
Assiouras, Ioannis af6016fe89 SWDEV-536048 - Add explicit cast to int when warpSize is used with == operator (#473) 
[ROCm/clr commit: a01134a1ab]
 2025-06-06 12:22:31 +05:30
Dittakavi, Satyanvesh 2001b250b5 SWDEV-535738 - Use the type traits implementation from __hip_internal namespace in hip headers (#471) 
[ROCm/clr commit: 1bce89ca51]
 2025-06-05 00:14:56 +05:30
Belton-Schure, Aidan 415ba0773b Add HIP_vector_base constructor (#456) 
Change-Id: I57571f6bf9876d3a9ff312b952858527ec6abc7b

[ROCm/clr commit: c2f6c98006]
 2025-06-03 12:23:38 +05:30
Xie, Jiabao(Jimbo) 1539c16ee8 SWDEV-535184 - Fix to switch HIP_RUNTIME_API_TABLE_STEP_VERSION 9 and 10 (#448) 
[ROCm/clr commit: 54a2e943d7]
 2025-06-02 21:24:40 +05:30
Patel, Jaydeepkumar cb73b09ce1 SWDEV-535738 - Use __hip_internal instead of std namespace for device code. (#455) 
[ROCm/clr commit: a018d4278a]
 2025-06-02 20:09:17 +05:30
Belton-Schure, Aidan 0fcf5ef2c5 Fix vec operator[] (#419) 
Change-Id: I1b6767b8b7f7dbcf9998662dd9f8ae3a96e1deca

[ROCm/clr commit: 9533ceb8aa]
 2025-05-29 13:09:34 +05:30
Betigeri, Sourabh 328aeb758c SWDEV-508963 - Change hipMemcpy subset of methods to use const qualifier for src (#139) (#418) 
Change-Id: Ia5b68bacff634d3d9e5757f515bac50887895bc3

[ROCm/clr commit: c0c3aa3637]
 2025-05-29 13:07:51 +05:30
Dittakavi, Satyanvesh dd3eaa86ff 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

[ROCm/clr commit: 54503f0d67]
 2025-05-29 12:52:10 +05:30
Belton-Schure, Aidan 8074e10ce6 Add operator[] to vec class (#402) 
Change-Id: I452cd5d1b258ccbc354a2f424d062f0fa91100ea

[ROCm/clr commit: 0f46bd5131]
 2025-05-28 13:38:38 +05:30
Brzak, Branislav d17388cd6b SWDEV-518962 - Match hipCtxGetApiVersion signature with Cuda (#327) 
[ROCm/clr commit: ee2d4a0358]
 2025-05-26 20:28:22 +05:30
Kandula, Venkateshwar reddy 10d0ae9638 [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>

[ROCm/clr commit: ad44dd527f]
 2025-05-22 15:09:31 +05:30
Dittakavi, Satyanvesh 26a927621b 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

[ROCm/clr commit: 946ea72c94]
 2025-05-22 13:28:31 +05:30
Dittakavi, Satyanvesh 4f6024c376 SWDEV-482125 - HIPRTC implementation of standard headers moved from std to __hip_internal namespace (#376) 
Change-Id: Ia359ed45e0916ad397401113a69a6b26a590d807

[ROCm/clr commit: b27e2abc41]
 2025-05-22 13:27:19 +05:30
Belton-Schure, Aidan 720660d552 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

[ROCm/clr commit: 3cd3b3ffc5]
 2025-05-21 21:30:55 +05:30
Assiouras, Ioannis 3ea732b38a SWDEV-508965 - [6.4 Preview] Remove HIP_MEMSET_NODE_PARAMS struct (#121) (#364) 
[ROCm/clr commit: 968f5599a8]
 2025-05-21 21:16:42 +05:30
Sang, Tao 12b43826dd SWDEV-519346 - Fix __clock64() compiling issue in SPIRV (#207) 
[ROCm/clr commit: 78f3caba0b]
 2025-05-20 16:57:17 -04:00
Jayaprakash, Karthik a17e22a78c SWDEV-529929 - hipMemGetHandleForAddressRange implementation. (#245) 
[ROCm/clr commit: 12131de4a9]
 2025-05-20 15:56:04 -04:00
Belton-Schure, Aidan c04355bbe3 Add __syncwarp operation (#160) 
Change-Id: I6a3783beafdbb9f11a3b37333f4ff3f5be27ea54

[ROCm/clr commit: c50610b44d]
 2025-05-15 14:20:13 +05:30
Hernandez, Gerardo f6e2ce709e 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

[ROCm/clr commit: 5606debd8e]
 2025-05-13 16:35:58 +01:00
Brzak, Branislav cc6bfe71fc SWDEV-528683 - Hardcode valid wavefront compile time options (#306) 
[ROCm/clr commit: f9199ac205]
 2025-05-12 19:29:39 +02:00
Chaudhary, Jatin Jaikishan a8630e866d SWDEV-525933 - add constexpr operators for fp16/bf16 (#199) 
[ROCm/clr commit: 2f73e1385b]
 2025-05-09 09:53:58 +01:00
Chaudhary, Jatin Jaikishan b5f67d4804 SWDEV-529854 - __hmax/__hmin should handle nan's (#246) 
[ROCm/clr commit: a71c6eb1a0]
 2025-05-06 09:42:15 +01:00
Chaudhary, Jatin Jaikishan a12739ecd9 SWDEV-529927 - add missing operations for fp16/bf16 (#238) 
[ROCm/clr commit: b1ebf33850]
 2025-05-06 09:41:21 +01:00
Arsenault, Matthew 13d8f9adae 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.

[ROCm/clr commit: 1db9a7d48b]
 2025-05-05 19:35:12 +02:00
Chaudhary, Jatin Jaikishan f619372ae6 SWDEV-514560 - add fp6 header implementation (#54) 
Co-authored-by: rahul manocha <rmanocha_amdeng>

[ROCm/clr commit: 12febe6782]
 2025-05-01 15:17:38 +01:00
Hernandez, Gerardo ba5a9a5395 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.

[ROCm/clr commit: 1a8d766836]
 2025-04-22 19:47:58 +05:30
Brzak, Branislav 5df4734f71 SWDEV-526612 - Add missing copyright notices (#201) 
[ROCm/clr commit: 99142c3dd9]
 2025-04-18 20:54:27 +05:30
Chaudhary, Jatin Jaikishan 7257b705ce SWDEV-512924 - add fp4 API (#52) 
* Remove C-style include guard

* clean up issues in the PR


[ROCm/clr commit: 5d638d831c]
 2025-04-15 17:53:50 +01:00
Chaudhary, Jatin Jaikishan 7c3dcb707e SWDEV-520627 - include warp functions header for warpSize (#177) 
Change-Id: Id3fff8f2722d521071ef0ff71b09fc365ef6fa82

[ROCm/clr commit: 588cf0fc69]
 2025-04-15 14:40:27 +05:30
Milanov, Aleksandar c83df8a653 SWDEV-526208 - Fix miscalculation of coalesced tiled partition mask (#162) 
[ROCm/clr commit: c4fa3ef927]
 2025-04-11 19:40:26 +02:00
Hernandez, Gerardo 9264e97cbb 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

[ROCm/clr commit: 66496258b4]
 2025-04-11 17:00:59 +05:30
Sang, Tao 929209b988 SWDEV-521083 - Fix atomicMin/Max issues (#151) 
Fix atomicMin/Max(), atomicMin/Max_system() issue on
float types.

[ROCm/clr commit: 6d10577761]
 2025-04-10 12:30:55 -04:00
Chaudhary, Jatin Jaikishan 665e88008b 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

[ROCm/clr commit: 628777b73d]
 2025-04-10 17:22:15 +01:00
Chaudhary, Jatin Jaikishan 07c7d3e860 SWDEV-520627 - include sync_warp header instead of warp function header (#18) 
Change-Id: Ic3f54b0f5bfee8565a8bbb6218fb0ccdb900c9ea

[ROCm/clr commit: 5c030840d6]
 2025-04-10 21:50:25 +05:30
Chaudhary, Jatin Jaikishan b0d9ea7454 SWDEV-525969 - add gfx950 to use ocp type for fp8 (#157) 
* do not use __gfx94plus_clr__ macro in fp8 header


[ROCm/clr commit: 5214d1ca07]
 2025-04-09 16:21:39 +01:00
Betigeri, Sourabh 487ede31a9 SWDEV-523281 - [clr] Implementation of hipLaunchKernelExC and hipDrvLaunchKernelEx API with support for cooperative launch (#92) 
[ROCm/clr commit: 8c6b90996e]
 2025-04-03 20:10:05 +09:00