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

1693 Commits

Author SHA1 Message Date
Rahul Manocha 436271e407 SWDEV-468039 - FP8 host only conversion support on mi200 
Change-Id: I0891f42d1b7c0d94d099fe26df5db3eff64ba564


[ROCm/clr commit: 39bbc0341d]
 2024-08-07 20:51:00 -04:00
Jaydeep Patel c51153f759 SWDEV-470886 - Add maybe_undef attribute for shfl device function due to not all lanes of wave define var and compiler needs to know about this. 
Change-Id: I3a683887e033305ac55362f356838b491a6d50f2


[ROCm/clr commit: 6344ddb2f3]
 2024-08-05 00:53:13 -04:00
Sourabh Betigeri f518648183 SWDEV-365151 - Fix the fns32 to do 32 bit computation and adds a wrapper to ease porting from CUDA 
Change-Id: I0b5a9ca11c98f8c1c40cfba7f4e057bfda2d756e


[ROCm/clr commit: 7298b80112]
 2024-07-26 11:20:14 -04:00
Sourabh Betigeri 5b0cc86295 SWDEV-475394 - Fix for the return type to be in-line with CUDA 
Change-Id: I7c833571d47b4e86a86e4a0095b61947d16ecab6


[ROCm/clr commit: 4fbd7abbb2]
 2024-07-25 16:33:33 -04:00
Rahul Manocha de67a2a1dc SWDEV-468039 - FP8 OCP headers 
Change-Id: Iecd32c5a0357781da07395d32f894415954b7b22


[ROCm/clr commit: 353f15afa6]
 2024-07-25 12:42:23 -04:00
Jatin Chaudhary 0193d66679 SWDEV-466747 - add shfl functions in bfloat16 
Change-Id: Ide7d7e1d449783cced8867abf43ff45f5bce113a


[ROCm/clr commit: e43176bde9]
 2024-07-23 10:51:02 -04:00
Jatin Chaudhary 6240b203dc SWDEV-466747 - optimize conversions for bfloat16 operations 
Since we made the members public, we can optimize some operations which
do not require redundant conversions to half_raw types.

Change-Id: I31555ef18e695d8e24b89f0418187fa4e932a38a


[ROCm/clr commit: 6a655a77e7]
 2024-07-17 18:37:25 -04:00
Julia Jiang 3c7ae28776 SWDEV-472710 - Adding gitattributes and remove trailing spaces 
Change-Id: Ic8ad2071745f0ffe6a2e120bfebb6d90bf270f87


[ROCm/clr commit: dd30e0e893]
 2024-07-15 12:39:56 -04:00
Marko Arandjelovic 6159b0eba0 SWDEV-472345 - Fix coalesced group size 
In case when the tile size is greater than the number of active threads,
the coalesced group size should be equal to the number of active threads.

Change-Id: I1d41322f2428a07862a590cb5d34b01243383b7c


[ROCm/clr commit: 152f343124]
 2024-07-12 04:29:53 -04:00
sdashmiz ba07492b20 SWDEV-421021 - Add APIs cuMemcpyNodeGet/Set params 
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I24bc0da56aad34c9d5876a3d83b59515f11dc3ea


[ROCm/clr commit: 57e79802cd]
 2024-07-10 12:53:01 -04:00
Jatin Chaudhary 58d27f6f47 SWDEV-460834 - add unsafe atomic add for fp16 and bf16 
Change-Id: I6de5c2c425c9f8ac7f6c4e5c83c8b8b7ac8fe4cb


[ROCm/clr commit: ecd812b2d8]
 2024-07-02 16:45:17 -04:00
Jatin Chaudhary 138caca238 SWDEV-450109 - add missing coop group operations 
any, all, ballot, match_any, match_all for coalesced_group and
tiled_group. add binary_partition.

Change-Id: I8ed0e654f36d788653352a0617beb5b8d44749bc


[ROCm/clr commit: 36d58c6518]
 2024-06-20 10:42:06 -04:00
Ioannis Assiouras 75104df3b2 SWDEV-464648 - code and comment cleanups 
Change-Id: I5ba3f1bff500b3cd5903c2f441017735e688f83f


[ROCm/clr commit: 8f42ad6aa3]
 2024-06-07 22:38:09 +01:00
Fabian Ritter 10b3a313d2 Add __device__ specifier to HIP's warpSize constexpr variable 
The warpSize variable is set to the value of the __AMDGCN_WAVEFRONT_SIZE macro,
which is a meaningless default in host code.
The resolution for SWDEV-449015 will introduce diagnostics for uses of this
macro in host code, which includes the current definition of the warpSize
variable. With the __device__ specifier, the definition of the warpSize
variable will not cause these diagnostics.

This change does not stop the variable from being used in host code since clang
intentionally does not diagnose uses of __device__ constexpr variables in host
code.

Change-Id: I0317217affe94fdf2dfd9ad0f134e68f5173245f


[ROCm/clr commit: 819e537dc5]
 2024-06-04 09:32:52 -04:00
Saleel Kudchadker a7ef9d9d39 SWDEV-301667 - Add enum for TASK command 
Change-Id: I5d2712d913bdbc42bc6eaccf9eb2f8cb9976b22e


[ROCm/clr commit: f18a3405d7]
 2024-05-28 14:18:46 -04:00
Jatin Chaudhary 026346a6c7 SWDEV-462027 - bfloat162 struct members are public 
bfloat16 member is protected, bfloat162 members are public

Change-Id: Ice090401553b9526cb34d003b3996b85d538bc4e


[ROCm/clr commit: 2b290953b0]
 2024-05-28 06:31:10 +00:00
Ioannis Assiouras a81f82931b SWDEV-449620 - Added hipExternalMemoryGetMappedMipmappedArray in the dispatch table 
Change-Id: I79952f131dea1949296685ee92cae05a4f9aec30


[ROCm/clr commit: 7b46701142]
 2024-05-28 06:31:10 +00:00
shadi 19ce99d104 SWDEV-421027 - Add more Graph APIs 
Signed-off-by: shadi <shadi.dashmiz@amd.com>
Change-Id: I0a1fc284e48317a49ca88d4ed4e3a10e752efd58


[ROCm/clr commit: e705e5e0d9]
 2024-05-28 06:28:17 +00:00
Anusha GodavarthySurya 9818ad44a2 SWDEV-453739 - Added definition for undeclared methods 
Change-Id: I3f10e0a3fb4bbaf9d873a6a988847f05ac43a312


[ROCm/clr commit: 45dccd6eab]
 2024-05-28 06:28:17 +00:00
shadi 15b8dc404c SWDEV-420016 - Add more driver side graph APIs 
Signed-off-by: shadi <shadi.dashmiz@amd.com>
Change-Id: Iff3ee7dcbcd24836f227fdc9bd5ff4b554ac914f


[ROCm/clr commit: f2b01782ac]
 2024-04-25 12:50:43 -04:00
Ioannis Assiouras b052e90ea8 SWDEV-451099 - Added include for __half type definitions for non-HIP code on windows 
Change-Id: Id80cef5a36db8707276de052cbaf73b6826d222f


[ROCm/clr commit: 2841aab017]
 2024-04-24 15:31:31 -04:00
Jatin Chaudhary 5a7b0c50ad SWDEV-379007 - fix bool check for fp8_fnuz 
for fnuz nums zero val is 0x00, -0 i.e. 0x80 would be a NaN

Change-Id: Ibdc4fb4b9fb307b5952434f08d45a8ddd6262db8


[ROCm/clr commit: d7b0d78fad]
 2024-04-17 05:31:21 -04:00
Sameer Sahasrabuddhe c44b565943 SWDEV-454959 - ignore the upper half of the mask in wave32 mode 
Change-Id: If027dd8cbe5cbe142fff353cb72c16f08e9aea8e


[ROCm/clr commit: 03562a2547]
 2024-04-17 10:12:57 +05:30
Jatin Chaudhary 80f02a1534 SWDEV-379007 - use avx instruction for bf16 cvt 
AMD CPUs have had avx512_bf16 support for quite some time now (from
consumer Ryzen 7000 series to enterprise grade CPUs). This
patch should allow users to use the hardware bf16 unit when running the
__host__ variants of the function. This can be enabled via `hipcc ...
-mavx512vl -mavx512bf16`.

Change-Id: I67c377afc95ddfe8d45a048dce078a247d4a1878


[ROCm/clr commit: 49349f168c]
 2024-04-16 18:35:08 -04:00
Jatin Chaudhary 5e42260071 SWDEV-379007 initial implementation of fp8 header 
Change-Id: Id9a5a85641882961e4d860a815217c641e6f3387


[ROCm/clr commit: ca07f59fb1]
 2024-04-16 05:37:59 -04:00
Jatin Chaudhary 565f915ebe SWDEV-379007 - add __hip_bfloat16_raw types 
This also brings bfloat16 implementation closer to CUDA's.

Change-Id: I23f381141faacd6537923ae9b88ada4d661db496


[ROCm/clr commit: 481912a1fd]
 2024-04-09 05:32:13 -04:00
kjayapra-amd a17f9d2d75 SWDEV-448288 - Remove references to deprecated llvm references. 
Change-Id: I7b58177c41dc0c6c59813977cb90e65a6cb3be72


[ROCm/clr commit: 5b28e386f6]
 2024-03-26 10:23:18 -04:00
Satyanvesh Dittakavi 15f01694f4 SWDEV-444021 - Implement hipGetFuncBySymbol 
Change-Id: I7ef13d02c5b5c6ce2386ccb92b5602d005b35988


[ROCm/clr commit: 684fd60c8f]
 2024-03-14 02:56:30 -04:00
Sameer Sahasrabuddhe dcd3673e66 SWDEV-295298: hide warp sync builtins in ROCm 6.2 (absent in 6.1) 
Change-Id: I6d59e4f3a133dce5ea0521fc9d09f93f9fd0f376


[ROCm/clr commit: 0421a35f71]
 2024-03-11 02:44:20 -04:00
Anusha GodavarthySurya 41ec66ff67 SWDEV-448586 - Added implementation for new API hipStreamBeginCaptureToGraph 
Change-Id: I1ce802102cef2b66c92d3375f769983841de793f


[ROCm/clr commit: 4feb1f9337]
 2024-03-07 05:24:49 +00:00
Rahul Manocha ee52a05009 SWDEV-448281 - Fix for CG Tile Partition Dynamic Size 
Change-Id: Ibdccec45470caf55a8c00194e6da0dde73e4b180


[ROCm/clr commit: e5e7687fd8]
 2024-03-04 13:00:24 -05:00
Ioannis Assiouras 609893e98f SWDEV-445960 - Fix __umul64hi __mul64hi on windows 
Change-Id: I4857d269bc80fcc0e6b77df941f24ff1964a4d42


[ROCm/clr commit: 74cae705ae]
 2024-03-03 12:42:42 -05:00
Jonathan R. Madsen 1f91fcf42e SWDEV-443227 - Update hip_api_trace.cpp 
- enforce incrementing the table versioning number when a table size changes outside of ifdef for ROCPROFILER_REGISTER
- add new HIP_ENFORCE_ABI entries
- update the HipDispatchTable size and bump HIP_RUNTIME_API_TABLE_STEP_VERSION to 1
- re-enable rocprofiler-register

Change-Id: Ie0cc1d8491c5640056e5dd393ea243e4dce4e8a9


[ROCm/clr commit: d84c5ae3af]
 2024-02-27 13:29:33 -05:00
Jaydeep Patel 97ad1412ea SWDEV-445539, SWDEV-445549 - Restore kernel launch prof data. 
Issue is due to INIT_hipLaunchKernel_CB_ARGS_DATA turns out to be empty when we generate new hip_prof_str.h with since
extra space introduced by https://gerrit-git.amd.com/c/compute/ec/clr/+/846275/77/hipamd/src/hip_module.cpp#667

Change-Id: Id251956c1663483cc7ab11a3722769cb5b0d3664


[ROCm/clr commit: d1b75659f2]
 2024-02-14 00:05:41 -05:00
jiabaxie b251260e2a SWDEV-437702 - implement hipGetProcAddress 
This should be used in place of dlsym or GetProcAddress (linux and windows respectively)

Change-Id: I5501b538e03892e8e5a2282678d848fcaf21d911


[ROCm/clr commit: 0479cdb3dd]
 2024-02-08 11:26:02 -05:00
Anusha GodavarthySurya 78aee3d79f SWDEV-439637 - Fix undefined symbol error hipTexRefGetArray & hipTexRefGetBorderColor (github id- #3382) 
Change-Id: Iae1956779e84759f23011151b60a01cef0d0ef32


[ROCm/clr commit: 5303cd2740]
 2024-02-06 03:04:07 -05:00
Rahul Manocha 7bccc07147 SWDEV-421025 - hipGraphInstantiateWithParams API changes 
Change-Id: Ib07d4dd1698220b68ed27f91d58d3bd315a8804c


[ROCm/clr commit: 1a3901fa49]
 2024-02-05 05:08:11 +00:00
Rahul Manocha 0aadfd0a6f SWDEV-421025 - Hip Launch Attribute ID/Value datatype for rocprof 
Change-Id: Icfcd1ecf2f2956ba015820ba250f06b438169a6a


[ROCm/clr commit: ceffb24416]
 2024-02-05 05:08:11 +00:00
Sameer Sahasrabuddhe ebcecf1f4f SWDEV-295298: hide warp sync builtins with a macro in ROCm 6.1 
Change-Id: Ie2efd233c0bcf8ad0e06223ee854fe4bd1060443


[ROCm/clr commit: d80168eb87]
 2024-02-05 05:08:11 +00:00
Anusha GodavarthySurya 401a2635ba SWDEV-439637 - Fix undefined symbol error(github id- #3382) 
Change-Id: Ic7812e21618f48c737ce2c51b22a153099cc5d66


[ROCm/clr commit: 242132caa7]
 2024-02-05 05:08:11 +00:00
Jatin Chaudhary f6fe357508 SWDEV-431367 - fix float compare for atomicMax/Min where -0.0 < +0.0 
Change-Id: Ief99aac3775b0a0c6372215e5a0a9ecbbdca4ca5


[ROCm/clr commit: a01257f000]
 2024-02-05 05:08:11 +00:00
Satyanvesh Dittakavi 62fbe82c51 SWDEV-435750 - Make threadIdx, blockIdx compatible with CUDA 
Change-Id: If0bd2c2c72e8418e30b5a202d9c10da334eff75e


[ROCm/clr commit: 0f9f8b52da]
 2024-01-18 02:16:15 -05:00
Jaydeep Patel baaad8b45a SWDEV-398109 - Enable smid support for CU mode for gfx10+. 
Change-Id: Ief1e7c952c3de7e4a547bbdcc9817f51aa4d5d36


[ROCm/clr commit: fb2c100d8c]
 2024-01-16 20:13:23 -05:00
Jonathan R. Madsen 3c5da3e666 Add rocprofiler-register support 
- Add API table versioning defines to hip_api_trace.hpp
- Add rocprofiler-register implementation in hip_api_trace.cpp
  - created ToolInit implementation
  - moved some local functions into anonymous namespace
  - renamed some function to be overloads
  - made most of the initialization code templated (reduced code duplication)
  - enforce ABI
- Updated hipamd/src/CMakeLists.txt
  - find rocprofiler-register package (enabled by default but not required)
- Address review comments
  - simplify size calculation for dispatch table
  - remove setting .size field in GetDispatchTableImpl
- Fix windows build
- Do not enforce ABI on Windows

Change-Id: I08766e8a083528a1236996274bf4522e0e8a1b9f


[ROCm/clr commit: 0f1fdc7794]
 2024-01-13 01:02:59 -06:00
sdashmiz 29eba6df09 SWDEV-421027 - Add hipGraphAddNode 
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: Ic8cf293ff483ee2547b52d2975062bcb9a6f5d17


[ROCm/clr commit: d23835ffbe]
 2024-01-12 11:36:30 -05:00
Sameer Sahasrabuddhe 2dba215643 SWDEV-295298: introduce warp sync functions 
The following builtins from the CUDA spec are implemented:

  - __all_sync, __any_sync, __ballot_sync and __activemask
  - __match_any_sync and __match_all_sync
  - __shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync

The following builtins are NOT implemented, pending support in the compiler:

  - __reduce_add_sync, __reduce_min_sync, __reduce_max_sync
  - __reduce_and_sync, __reduce_or_sync, __reduce_xor_sync

Change-Id: I07dedbbfe5449f4b5c9b040bed59f5603ccec8c3


[ROCm/clr commit: c5ab5680b4]
 2024-01-11 22:31:27 -05:00
jujiang 898e647f84 SWDEV-410522 - Fix a doxygen warning in a header in AMD platform 
Change-Id: I29cd222af453c42a5d472081747bacf0c3f8d17c


[ROCm/clr commit: 2a2fea6c07]
 2024-01-11 13:40:25 -05:00
sdashmiz 57c0674538 SWDEV-403382 - Add missing graph APIs 
- add external semaphor signal and wait node

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: If205d48bc4fb60ce2966ca6ed0419e7700755bdd


[ROCm/clr commit: 087767d9ac]
 2024-01-09 14:19:50 -05:00
Jaydeep Patel 5f13c16039 SWDEV-437440 - Specify which defination to use at compile time. 
Change-Id: Ibeee914f293d26fa68ab793cac8b662aa34e8df3


[ROCm/clr commit: c3fc607e7d]
 2023-12-21 19:59:39 -05:00
Rahul Garg 0051334261 SWDEV-388256 - Add support of hipExtGetLastError 
Change-Id: Ie119d5ea611019996b8a80f7b820ca6a160496a6


[ROCm/clr commit: 92aa9d0fba]
 2023-12-18 22:03:22 -05:00