ccb09057d8
SWDEV-490198: _sync() will be enabled by default in 6.4
...
Change-Id: Id029424a9c0f6b144a7aa0e96fe8acc4a138ec51
[ROCm/clr commit: 556390f9c7
2024-10-25 09:54:40 -04:00
b8a5d8396d
SWDEV-468039,SWDEV-482579 - Enable FP8 SW Conversions on pre gfx940 archs
...
1) SW Conversions for ocp and fnuz are enabled on pre mi300 archs
2) for mi300 only fnuz is enabled
3) for gfx1200 only ocp is enabled
Change-Id: I90373752a2d15eff20d5deec874ed396ba4e1788
[ROCm/clr commit: e729f08704
2024-10-17 11:49:22 -04:00
da59165313
SWDEV-486137 - match behavior of int variants of hadd/uhadd/rhadd/urhadd
...
Match cases and handle cases where it can overflow.
Change-Id: I3d6f802686af230a622ef9891a844135ad3d1ae5
[ROCm/clr commit: b977101893
2024-10-09 13:47:33 -04:00
98539759f1
SWDEV-483436 - User spt stream as def with -fgpu-default-stream=per-thread for hipMemsetAsync.
...
Change-Id: Ia85c2b4c40fc9250754d3b64fb9fd1c615362572
[ROCm/clr commit: d6193a2f23
2024-09-29 01:42:33 -04:00
43a1d720d5
Fix HIP API trace versioning
...
Change-Id: I33f2be4668c96e2225d4ca9a253e61ec2dc65102
[ROCm/clr commit: 07c9c7fe56
2024-09-25 10:32:14 -04:00
29e9bed35d
SWDEV481762 - Updated definition of 'DEPRECATED' in header file
...
Change-Id: I88986b8e1815f3d816595f3eb2da8a6c1c1c2993
Jenifer helped make a combined PSDB build, together with the change in hip repos
https://gerrit-git.amd.com/c/compute/ec/hip/+/1114046
Combined PSDB verification passed.
http://rocm-ci.amd.com/job/compute-psdb-staging-hip/17293/
[ROCm/clr commit: bb03ef11a3
2024-09-05 15:41:04 -04:00
11632a954a
SWDEV-477218 - Implement hipDeviceGetTexture1DLinearMaxWidth
...
Change-Id: I8103f710abeb869f5f84be61c57a30b24356def6
[ROCm/clr commit: 8be00b6602
2024-09-05 15:09:38 -04:00
a00f071579
SWDEV-470372 - Added hipExtHostAlloc API
...
This change adds a new HIP API `hipExtHostAlloc` which preserves
the functionality of `hipHostMalloc`.
Change-Id: I13504c6fc13465ddd7aed329795bb4f2fef1baff
[ROCm/clr commit: 2c84211b58
2024-08-27 08:26:03 -04:00
5e3fe9bd1f
SWDEV-480489 - fix unsafeAtomicAdd
...
Integration into pytorch pointed out some issues, value narrowing, to
fix this we are now using unions. Also removed check for -munsafe*
compiler flag. The check is now just on builtin detection.
Change-Id: I49364503fa429bd862952f9b29879072afa6d553
[ROCm/clr commit: bb52d9ed62
2024-08-27 06:29:11 -04:00
aacf75f480
SWDEV-480725: missing __ockl_wfall __ockl_wfany in amd_hip_bf16.h
...
Change-Id: Iff4aeec411bfeaf4cc187c515e2da3d5898f89cb
[ROCm/clr commit: 6df2da65cd
2024-08-25 22:49:14 -04:00
6907974f90
SWDEV-473942 - SWDEV-431367 - Correct atomicMax(_system) and atomicMin(_system)
...
- Fixes -0.0 and +0.0 comparison. For atomicMax if the value on
address is -0.0 and on val is +0.0, gfx90a's unsafe atomics will swap
them. This behavior should be consistent with cas loop as well.
- _system variants of atomicMax and atomicMin are resulting in
incorrect output. Updated these to use the similar implementation as
atomicMax and atomicMin.
Change-Id: I20df36ee29ae0434a6b564f2ba71193fe41cfa59
[ROCm/clr commit: d69cc35750
2024-08-13 10:38:50 -04:00
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
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
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
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
de67a2a1dc
SWDEV-468039 - FP8 OCP headers
...
Change-Id: Iecd32c5a0357781da07395d32f894415954b7b22
[ROCm/clr commit: 353f15afa6
2024-07-25 12:42:23 -04:00
0193d66679
SWDEV-466747 - add shfl functions in bfloat16
...
Change-Id: Ide7d7e1d449783cced8867abf43ff45f5bce113a
[ROCm/clr commit: e43176bde9
2024-07-23 10:51:02 -04:00
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
3c7ae28776
SWDEV-472710 - Adding gitattributes and remove trailing spaces
...
Change-Id: Ic8ad2071745f0ffe6a2e120bfebb6d90bf270f87
[ROCm/clr commit: dd30e0e893
2024-07-15 12:39:56 -04:00
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
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
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
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
75104df3b2
SWDEV-464648 - code and comment cleanups
...
Change-Id: I5ba3f1bff500b3cd5903c2f441017735e688f83f
[ROCm/clr commit: 8f42ad6aa3
2024-06-07 22:38:09 +01:00
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
a7ef9d9d39
SWDEV-301667 - Add enum for TASK command
...
Change-Id: I5d2712d913bdbc42bc6eaccf9eb2f8cb9976b22e
[ROCm/clr commit: f18a3405d7
2024-05-28 14:18:46 -04:00
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
a81f82931b
SWDEV-449620 - Added hipExternalMemoryGetMappedMipmappedArray in the dispatch table
...
Change-Id: I79952f131dea1949296685ee92cae05a4f9aec30
[ROCm/clr commit: 7b46701142
2024-05-28 06:31:10 +00:00
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
9818ad44a2
SWDEV-453739 - Added definition for undeclared methods
...
Change-Id: I3f10e0a3fb4bbaf9d873a6a988847f05ac43a312
[ROCm/clr commit: 45dccd6eab
2024-05-28 06:28:17 +00:00
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
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
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
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
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
5e42260071
SWDEV-379007 initial implementation of fp8 header
...
Change-Id: Id9a5a85641882961e4d860a815217c641e6f3387
[ROCm/clr commit: ca07f59fb1
2024-04-16 05:37:59 -04:00
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
a17f9d2d75
SWDEV-448288 - Remove references to deprecated llvm references.
...
Change-Id: I7b58177c41dc0c6c59813977cb90e65a6cb3be72
[ROCm/clr commit: 5b28e386f6
2024-03-26 10:23:18 -04:00
15f01694f4
SWDEV-444021 - Implement hipGetFuncBySymbol
...
Change-Id: I7ef13d02c5b5c6ce2386ccb92b5602d005b35988
[ROCm/clr commit: 684fd60c8f
2024-03-14 02:56:30 -04:00
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
41ec66ff67
SWDEV-448586 - Added implementation for new API hipStreamBeginCaptureToGraph
...
Change-Id: I1ce802102cef2b66c92d3375f769983841de793f
[ROCm/clr commit: 4feb1f9337
2024-03-07 05:24:49 +00:00
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
609893e98f
SWDEV-445960 - Fix __umul64hi __mul64hi on windows
...
Change-Id: I4857d269bc80fcc0e6b77df941f24ff1964a4d42
[ROCm/clr commit: 74cae705ae
2024-03-03 12:42:42 -05:00
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
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
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
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
7bccc07147
SWDEV-421025 - hipGraphInstantiateWithParams API changes
...
Change-Id: Ib07d4dd1698220b68ed27f91d58d3bd315a8804c
[ROCm/clr commit: 1a3901fa49
2024-02-05 05:08:11 +00:00
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
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
Sameer Sahasrabuddhe
Rahul Manocha
Jatin Chaudhary
Jaydeep Patel
Jonathan R. Madsen
Julia Jiang
victzhan
Ioannis Assiouras
Satyanvesh Dittakavi
Sourabh Betigeri
Marko Arandjelovic
sdashmiz
Fabian Ritter
Saleel Kudchadker
Anusha GodavarthySurya
kjayapra-amd
Rahul Manocha
jiabaxie