Skip to content

SWDEV-516877 - Fix unsafe uses of unsafe for SPIRV #141

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 46 commits into
base: amd-staging
Choose a base branch
from
Open

SWDEV-516877 - Fix unsafe uses of unsafe for SPIRV #141

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
46 commits
Select commit Hold shift + click to select a range
8c8c00f
Promote till commit '667f048d187af06e1ed8aa84a387da68a1ba0e76'
mangupta Apr 21, 2023
5cba302
Promote till commit '636983de499beb115e88e72f84f6d1e32f2c6b59'
mangupta May 2, 2023
53bdffa
SWDEV-1 - - Switch to new patch version
mangupta May 2, 2023
2079243
SWDEV-403384 - Change OpenCL version number from 3581 to 3582
gargrahul May 31, 2023
07a6392
Promote till commit '15bafc9a935cfbdf3ea4878ee8ef8ec2db521e96'
mangupta Jun 7, 2023
c080c14
SWDEV-1 - Bump OpenCL build number to 3583
gargrahul Jun 15, 2023
cc9c187
Promote till commit '90f53925a3c39d0855a903ccc97f2ba5d5258ba6'
mangupta Jul 12, 2023
276d659
SWDEV-2 - Change OpenCL version number from 3583 to 3588
mangupta Jul 12, 2023
a835d18
Promote till commit '7912f3af89e1afdb856da24f5ebc9f60283951e0'
mangupta Jul 19, 2023
f5021ed
SWDEV-2 - Change OpenCL version number from 3588 to 3590
mangupta Jul 19, 2023
1a52c4f
Promote till commit d18f4b1ad406b6ab6d28e44437adf7928c42148e
rakesroy Aug 24, 2023
cefafdb
SWDEV-2 - Change OpenCL version number from 3590 to 3594
rakesroy Aug 24, 2023
a09af41
Promote till commit 9bd592e2ff3c99cdf26dbd4d1b1e1b181f605609
rakesroy Sep 5, 2023
fe4cce8
SWDEV-2 - Change OpenCL version number from 3594 to 3596
rakesroy Sep 5, 2023
315c7d0
Promote till commit 'fbea58ba117c96f80f6d5a2911bb941f7b1f7d23'
rakesroy Sep 29, 2023
dd55732
SWDEV-2 - Change OpenCL version number from 3596 to 3600
rakesroy Sep 29, 2023
c87c347
Promote till commit 2535fac29634a06c7da5cc570e4645cf8481fd8d
rakesroy Oct 16, 2023
a505a6a
SWDEV-2 - Change OpenCL version number from 3600 to 3602
rakesroy Oct 16, 2023
2192e20
Promote till commit 'd6d235a111247bbae1be71eb7ed2feee4207bd0a'
mangupta Dec 21, 2023
08aa6bd
SWDEV-2 - Change OpenCL version number from 3602 to 3610
mangupta Dec 21, 2023
e67b18e
Promote till commit '15b8cf911aaff1a65618d39f10fcf53fd7f9f7dd'
mangupta Jan 3, 2024
3a954af
SWDEV-2 - Change OpenCL version number from 3610 to 3612
mangupta Jan 3, 2024
aded7cd
Promote till commit 'dec1158d04d5d20bb91ada1a46b495312f0b831b'
rakesroy Jan 18, 2024
8b6f27a
SWDEV-2 - Change OpenCL version number from 3612 to 3614
rakesroy Jan 18, 2024
bfb1ccc
SWDEV-444098 - Porting the fix to mainline for removal of rocm-ocl-ic…
jujiang-del Feb 8, 2024
67f5fae
Revert "SWDEV-444098 - Porting the fix to mainline for removal of roc…
jujiang-del Mar 18, 2024
6b005a9
Promote till commit '47e3ed545d134bf8c6c1039ebdef4c2ac877fdd7'
rakesroy Apr 4, 2024
1d20b98
SWDEV-2 - Change OpenCL version number from 3614 to 3620
rakesroy Apr 4, 2024
a64b823
Promote till commit '863c56262e5ea5955cdcbce5423302521ee1e44a'
rakesroy Apr 30, 2024
6e2222b
SWDEV-2 - Change OpenCL version number from 3620 to 3622
rakesroy Apr 30, 2024
b76b53d
Promote till commit '3a5cbb91b930404b0c26967d539a5a446615c0c9'
rakesroy Jun 25, 2024
e577c9a
SWDEV-2 - Change OpenCL version number from 3622 to 3630
rakesroy Jun 25, 2024
30eaaab
Promote till commit '7d3c0c5e1064b25e025bc514952af9a2cc031e7e'
rakesroy Aug 6, 2024
de91dc5
SWDEV-2 - Change OpenCL version number from 3630 to 3632
rakesroy Aug 6, 2024
6172922
Promote till commit '9b33db9b24acf880f03a57de0a4756c31ecb70d7'
rakesroy Aug 20, 2024
1a9c8df
SWDEV-2 - Change OpenCL version number from 3632 to 3634
rakesroy Aug 20, 2024
53431b0
Promote till commit 'efce2f77c4b5a1b9a2c16f33b7b9062faf09b007'
satyanveshd Dec 1, 2024
220465d
SWDEV-2 - Change OpenCL version number from 3634 to 3642
satyanveshd Dec 1, 2024
8a1b617
Promote till commit 'e3b87544482f43760e0bf1c49e628039199c4bdf'
satyanveshd Dec 6, 2024
03366da
SWDEV-2 - Change OpenCL version number from 3642 to 3644
satyanveshd Dec 6, 2024
9cbb62c
Promote till commit '6c755a411648c892d917d2415c5e646ab1e37fe3'
satyanveshd Dec 13, 2024
9c3f9f8
SWDEV-2 - Change OpenCL version number from 3644 to 3646
satyanveshd Dec 13, 2024
4a9154c
Promote till commit 'f1c05e902656c15d7a6c229e0b400a1656dbcb5f'
satyanveshd Dec 20, 2024
0c88197
SWDEV-2 - Change OpenCL version number from 3646 to 3648
satyanveshd Dec 20, 2024
499cee3
SPIRV cannot robustly deal with some builtins yet, so bypass them for…
AlexVlx Feb 25, 2025
a0c9655
Remove noise.
AlexVlx Feb 26, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion hipamd/include/hip/amd_detail/amd_hip_bf16.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1835,7 +1835,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 h2trunc(const __hip_bfloat162 h) {
*/
__BF16_DEVICE_STATIC__ __hip_bfloat162 unsafeAtomicAdd(__hip_bfloat162* address,
__hip_bfloat162 value) {
#if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2bf16)
#if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2bf16) && !defined(__SPIRV__)
typedef short __attribute__((ext_vector_type(2))) vec_short2;
static_assert(sizeof(vec_short2) == sizeof(__hip_bfloat162_raw));
union {
Expand Down
2 changes: 1 addition & 1 deletion hipamd/include/hip/amd_detail/amd_hip_fp16.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1530,7 +1530,7 @@ THE SOFTWARE.
// Atomic
#if defined(__clang__) && defined(__HIP__)
inline __device__ __half2 unsafeAtomicAdd(__half2* address, __half2 value) {
#if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2f16)
#if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2f16) && !defined(__SPIRV__)
// The api expects an ext_vector_type of half
typedef _Float16 __attribute__((ext_vector_type(2))) vec_fp162;
static_assert(sizeof(vec_fp162) == sizeof(__half2_raw));
Expand Down