* 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
* 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>
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
* 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
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.
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.
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
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
Also removes asserts in cooperative groups shfl functions since
__hip_bfloat16 shfl is present now
Change-Id: I57578b6e68dccc10c2ddcd194e9cc18bc7732ce1
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
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
- 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
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 78f62d323
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
Xie, Jiabao(Jimbo)
Patel, Jaydeepkumar
Belton-Schure, Aidan
Betigeri, Sourabh
Dittakavi, Satyanvesh
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