* 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]
* 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]
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]
* 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]
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]
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]
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]
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]
Also removes asserts in cooperative groups shfl functions since
__hip_bfloat16 shfl is present now
Change-Id: I57578b6e68dccc10c2ddcd194e9cc18bc7732ce1
[ROCm/clr commit: 376f23b86a]
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
[ROCm/clr commit: d91e1f19d0]
1) Add Linker APIs to runtime to support SPIRV linking
2) Migrate Internal implementations to runtime and share with rtc
3) Add Support to bundled and unbundled SPIRV Code object linking.
Change-Id: Ic1fd4431f842a208a2468e8aec54a65b5fa6b0e3
[ROCm/clr commit: 5930f047bb]
- Use correct header in device_library_decl
- use std:: instead of __hip_internal:: for host compilation
- hide device specific stuff behind __clang__ and __HIP__ check
Change-Id: I2f3647e00555ed0e79f9954a459c41394c3cd49b
[ROCm/clr commit: c3f49c8788]
thread_rank() gives thread index in a block. Limit the range to the
current warp size.
Change-Id: Ib5c9831236096485cf99ba7ab0b911a3b10de31c
[ROCm/clr commit: bd7d40a4d8]
This removes almost all uses of the deprecated
__AMDGCN_WAVEFRONT_SIZE macro, which is unavailable
when targeting SPIR-V, and adds a SPIR-V compatible
formulation of warpSize (which should end up as the
sole definition of warpSize once we remove support
for treating it as a compile time constant). It
is incomplete in that the cooperative_groups
implementation will need additional surgery.
Squashed commit of the following:
commit 6840826c3fec8516857dc4f2092d84358550f588
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 23:36:32 2024 +0000
Add deprecation warning for constexpr uses of `warpSize`.
commit a72307a7353034c2de53fd164e016967945fd0d1
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 23:12:14 2024 +0000
Prepare HIP RT for SPIR-V.
commit 5e40dd746ac4f8c93b521ef048ff9d494905ba95
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 22:46:05 2024 +0000
Revert stale change.
commit 231fe91c53dba4cabd832fc84eaa6ddb402271a0
Merge: a48905ec9 12dc02b4f
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Fri Dec 6 22:37:24 2024 +0000
Merge branch 'amd-staging' of https://github.com/ROCm/clr into amd-staging
commit a48905ec9cfe0e017cc64943195be82b530117d7
Author: Alex Voicu <alexandru.voicu@amd.com>
Date: Tue Sep 17 03:14:56 2024 +0100
Add scaffolding for SPIR-V support.
Change-Id: I2e84bbe90df58a5f9a8709b619905f04fa5b96dc
[ROCm/clr commit: dd4378611a]
Betigeri, Sourabh
Dittakavi, Satyanvesh
Belton-Schure, Aidan
Brzak, Branislav
Kandula, Venkateshwar reddy
Assiouras, Ioannis
Sang, Tao
Jayaprakash, Karthik
Hernandez, Gerardo
Chaudhary, Jatin Jaikishan
Arsenault, Matthew
Milanov, Aleksandar
Rahul Manocha
Jimbo Xie
Steven Chung
Konstantin Zhuravlyov
Jaydeep Patel
zichguan-amd