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

1693 Commits

Author SHA1 Message Date
Rahul Manocha 39bbc0341d SWDEV-468039 - FP8 host only conversion support on mi200 
Change-Id: I0891f42d1b7c0d94d099fe26df5db3eff64ba564
 2024-08-07 20:51:00 -04:00
Jaydeep Patel 6344ddb2f3 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
 2024-08-05 00:53:13 -04:00
Sourabh Betigeri 7298b80112 SWDEV-365151 - Fix the fns32 to do 32 bit computation and adds a wrapper to ease porting from CUDA 
Change-Id: I0b5a9ca11c98f8c1c40cfba7f4e057bfda2d756e
 2024-07-26 11:20:14 -04:00
Sourabh Betigeri 4fbd7abbb2 SWDEV-475394 - Fix for the return type to be in-line with CUDA 
Change-Id: I7c833571d47b4e86a86e4a0095b61947d16ecab6
 2024-07-25 16:33:33 -04:00
Rahul Manocha 353f15afa6 SWDEV-468039 - FP8 OCP headers 
Change-Id: Iecd32c5a0357781da07395d32f894415954b7b22
 2024-07-25 12:42:23 -04:00
Jatin Chaudhary e43176bde9 SWDEV-466747 - add shfl functions in bfloat16 
Change-Id: Ide7d7e1d449783cced8867abf43ff45f5bce113a
 2024-07-23 10:51:02 -04:00
Jatin Chaudhary 6a655a77e7 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
 2024-07-17 18:37:25 -04:00
Julia Jiang dd30e0e893 SWDEV-472710 - Adding gitattributes and remove trailing spaces 
Change-Id: Ic8ad2071745f0ffe6a2e120bfebb6d90bf270f87
 2024-07-15 12:39:56 -04:00
Marko Arandjelovic 152f343124 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
 2024-07-12 04:29:53 -04:00
sdashmiz 57e79802cd SWDEV-421021 - Add APIs cuMemcpyNodeGet/Set params 
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I24bc0da56aad34c9d5876a3d83b59515f11dc3ea
 2024-07-10 12:53:01 -04:00
Jatin Chaudhary ecd812b2d8 SWDEV-460834 - add unsafe atomic add for fp16 and bf16 
Change-Id: I6de5c2c425c9f8ac7f6c4e5c83c8b8b7ac8fe4cb
 2024-07-02 16:45:17 -04:00
Jatin Chaudhary 36d58c6518 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
 2024-06-20 10:42:06 -04:00
Ioannis Assiouras 8f42ad6aa3 SWDEV-464648 - code and comment cleanups 
Change-Id: I5ba3f1bff500b3cd5903c2f441017735e688f83f
 2024-06-07 22:38:09 +01:00
Fabian Ritter 819e537dc5 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
 2024-06-04 09:32:52 -04:00
Saleel Kudchadker f18a3405d7 SWDEV-301667 - Add enum for TASK command 
Change-Id: I5d2712d913bdbc42bc6eaccf9eb2f8cb9976b22e
 2024-05-28 14:18:46 -04:00
Jatin Chaudhary 2b290953b0 SWDEV-462027 - bfloat162 struct members are public 
bfloat16 member is protected, bfloat162 members are public

Change-Id: Ice090401553b9526cb34d003b3996b85d538bc4e
 2024-05-28 06:31:10 +00:00
Ioannis Assiouras 7b46701142 SWDEV-449620 - Added hipExternalMemoryGetMappedMipmappedArray in the dispatch table 
Change-Id: I79952f131dea1949296685ee92cae05a4f9aec30
 2024-05-28 06:31:10 +00:00
shadi e705e5e0d9 SWDEV-421027 - Add more Graph APIs 
Signed-off-by: shadi <shadi.dashmiz@amd.com>
Change-Id: I0a1fc284e48317a49ca88d4ed4e3a10e752efd58
 2024-05-28 06:28:17 +00:00
Anusha GodavarthySurya 45dccd6eab SWDEV-453739 - Added definition for undeclared methods 
Change-Id: I3f10e0a3fb4bbaf9d873a6a988847f05ac43a312
 2024-05-28 06:28:17 +00:00
shadi f2b01782ac SWDEV-420016 - Add more driver side graph APIs 
Signed-off-by: shadi <shadi.dashmiz@amd.com>
Change-Id: Iff3ee7dcbcd24836f227fdc9bd5ff4b554ac914f
 2024-04-25 12:50:43 -04:00
Ioannis Assiouras 2841aab017 SWDEV-451099 - Added include for __half type definitions for non-HIP code on windows 
Change-Id: Id80cef5a36db8707276de052cbaf73b6826d222f
 2024-04-24 15:31:31 -04:00
Jatin Chaudhary d7b0d78fad 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
 2024-04-17 05:31:21 -04:00
Sameer Sahasrabuddhe 03562a2547 SWDEV-454959 - ignore the upper half of the mask in wave32 mode 
Change-Id: If027dd8cbe5cbe142fff353cb72c16f08e9aea8e
 2024-04-17 10:12:57 +05:30
Jatin Chaudhary 49349f168c 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
 2024-04-16 18:35:08 -04:00
Jatin Chaudhary ca07f59fb1 SWDEV-379007 initial implementation of fp8 header 
Change-Id: Id9a5a85641882961e4d860a815217c641e6f3387
 2024-04-16 05:37:59 -04:00
Jatin Chaudhary 481912a1fd SWDEV-379007 - add __hip_bfloat16_raw types 
This also brings bfloat16 implementation closer to CUDA's.

Change-Id: I23f381141faacd6537923ae9b88ada4d661db496
 2024-04-09 05:32:13 -04:00
kjayapra-amd 5b28e386f6 SWDEV-448288 - Remove references to deprecated llvm references. 
Change-Id: I7b58177c41dc0c6c59813977cb90e65a6cb3be72
 2024-03-26 10:23:18 -04:00
Satyanvesh Dittakavi 684fd60c8f SWDEV-444021 - Implement hipGetFuncBySymbol 
Change-Id: I7ef13d02c5b5c6ce2386ccb92b5602d005b35988
 2024-03-14 02:56:30 -04:00
Sameer Sahasrabuddhe 0421a35f71 SWDEV-295298: hide warp sync builtins in ROCm 6.2 (absent in 6.1) 
Change-Id: I6d59e4f3a133dce5ea0521fc9d09f93f9fd0f376
 2024-03-11 02:44:20 -04:00
Anusha GodavarthySurya 4feb1f9337 SWDEV-448586 - Added implementation for new API hipStreamBeginCaptureToGraph 
Change-Id: I1ce802102cef2b66c92d3375f769983841de793f
 2024-03-07 05:24:49 +00:00
Rahul Manocha e5e7687fd8 SWDEV-448281 - Fix for CG Tile Partition Dynamic Size 
Change-Id: Ibdccec45470caf55a8c00194e6da0dde73e4b180
 2024-03-04 13:00:24 -05:00
Ioannis Assiouras 74cae705ae SWDEV-445960 - Fix __umul64hi __mul64hi on windows 
Change-Id: I4857d269bc80fcc0e6b77df941f24ff1964a4d42
 2024-03-03 12:42:42 -05:00
Jonathan R. Madsen d84c5ae3af 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
 2024-02-27 13:29:33 -05:00
Jaydeep Patel d1b75659f2 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
 2024-02-14 00:05:41 -05:00
jiabaxie 0479cdb3dd SWDEV-437702 - implement hipGetProcAddress 
This should be used in place of dlsym or GetProcAddress (linux and windows respectively)

Change-Id: I5501b538e03892e8e5a2282678d848fcaf21d911
 2024-02-08 11:26:02 -05:00
Anusha GodavarthySurya 5303cd2740 SWDEV-439637 - Fix undefined symbol error hipTexRefGetArray & hipTexRefGetBorderColor (github id- #3382) 
Change-Id: Iae1956779e84759f23011151b60a01cef0d0ef32
 2024-02-06 03:04:07 -05:00
Rahul Manocha 1a3901fa49 SWDEV-421025 - hipGraphInstantiateWithParams API changes 
Change-Id: Ib07d4dd1698220b68ed27f91d58d3bd315a8804c
 2024-02-05 05:08:11 +00:00
Rahul Manocha ceffb24416 SWDEV-421025 - Hip Launch Attribute ID/Value datatype for rocprof 
Change-Id: Icfcd1ecf2f2956ba015820ba250f06b438169a6a
 2024-02-05 05:08:11 +00:00
Sameer Sahasrabuddhe d80168eb87 SWDEV-295298: hide warp sync builtins with a macro in ROCm 6.1 
Change-Id: Ie2efd233c0bcf8ad0e06223ee854fe4bd1060443
 2024-02-05 05:08:11 +00:00
Anusha GodavarthySurya 242132caa7 SWDEV-439637 - Fix undefined symbol error(github id- #3382) 
Change-Id: Ic7812e21618f48c737ce2c51b22a153099cc5d66
 2024-02-05 05:08:11 +00:00
Jatin Chaudhary a01257f000 SWDEV-431367 - fix float compare for atomicMax/Min where -0.0 < +0.0 
Change-Id: Ief99aac3775b0a0c6372215e5a0a9ecbbdca4ca5
 2024-02-05 05:08:11 +00:00
Satyanvesh Dittakavi 0f9f8b52da SWDEV-435750 - Make threadIdx, blockIdx compatible with CUDA 
Change-Id: If0bd2c2c72e8418e30b5a202d9c10da334eff75e
 2024-01-18 02:16:15 -05:00
Jaydeep Patel fb2c100d8c SWDEV-398109 - Enable smid support for CU mode for gfx10+. 
Change-Id: Ief1e7c952c3de7e4a547bbdcc9817f51aa4d5d36
 2024-01-16 20:13:23 -05:00
Jonathan R. Madsen 0f1fdc7794 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
 2024-01-13 01:02:59 -06:00
sdashmiz d23835ffbe SWDEV-421027 - Add hipGraphAddNode 
Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: Ic8cf293ff483ee2547b52d2975062bcb9a6f5d17
 2024-01-12 11:36:30 -05:00
Sameer Sahasrabuddhe c5ab5680b4 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
 2024-01-11 22:31:27 -05:00
jujiang 2a2fea6c07 SWDEV-410522 - Fix a doxygen warning in a header in AMD platform 
Change-Id: I29cd222af453c42a5d472081747bacf0c3f8d17c
 2024-01-11 13:40:25 -05:00
sdashmiz 087767d9ac SWDEV-403382 - Add missing graph APIs 
- add external semaphor signal and wait node

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: If205d48bc4fb60ce2966ca6ed0419e7700755bdd
 2024-01-09 14:19:50 -05:00
Jaydeep Patel c3fc607e7d SWDEV-437440 - Specify which defination to use at compile time. 
Change-Id: Ibeee914f293d26fa68ab793cac8b662aa34e8df3
 2023-12-21 19:59:39 -05:00
Rahul Garg 92aa9d0fba SWDEV-388256 - Add support of hipExtGetLastError 
Change-Id: Ie119d5ea611019996b8a80f7b820ca6a160496a6
 2023-12-18 22:03:22 -05:00