Skip to content

Commit b13561c

Browse files
hdelansteffenlarsen
hdelan
and
steffenlarsen
authored
[SYCL][HIP] Add gfx9+ hip atomics (#8170)
Adding gfx90a, gfx940 atomic builtins. --------- Co-authored-by: Steffen Larsen <[email protected]>
1 parent 16590f9 commit b13561c

14 files changed

+108
-34
lines changed

‎libclc/CMakeLists.txt

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -336,6 +336,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
336336
# Disables NVVM reflection to defer to after linking
337337
set( flags "SHELL:-Xclang -target-feature" "SHELL:-Xclang +ptx72"
338338
"SHELL:-march=sm_86" "SHELL:-mllvm --nvvm-reflect-enable=false")
339+
elseif( ${ARCH} STREQUAL amdgcn )
340+
# AMDGCN needs libclc to be compiled to high bc version since all atomic
341+
# clang builtins need to be accessible
342+
set( flags "SHELL:-mcpu=gfx940")
339343
else()
340344
set ( flags )
341345
endif()

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_add.cl

Lines changed: 36 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -10,13 +10,34 @@
1010
#include <spirv/spirv.h>
1111
#include <spirv/spirv_types.h>
1212

13+
extern constant int __oclc_ISA_version;
14+
1315
AMDGPU_ATOMIC(_Z18__spirv_AtomicIAdd, int, i, __hip_atomic_fetch_add)
1416
AMDGPU_ATOMIC(_Z18__spirv_AtomicIAdd, unsigned int, j, __hip_atomic_fetch_add)
1517
AMDGPU_ATOMIC(_Z18__spirv_AtomicIAdd, long, l, __hip_atomic_fetch_add)
1618
AMDGPU_ATOMIC(_Z18__spirv_AtomicIAdd, unsigned long, m, __hip_atomic_fetch_add)
17-
AMDGPU_ATOMIC(_Z21__spirv_AtomicFAddEXT, float, f, __hip_atomic_fetch_add)
1819

19-
#define AMDGPU_ATOMIC_FP64_ADD_IMPL(AS, AS_MANGLED, SUB1, SUB2) \
20+
#define AMDGPU_ATOMIC_FP32_ADD_IMPL(AS, AS_MANGLED, SUB1, CHECK, NEW_BUILTIN) \
21+
_CLC_DEF float \
22+
_Z21__spirv_AtomicFAddEXTP##AS_MANGLED##fN5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagEf( \
23+
volatile AS float *p, enum Scope scope, \
24+
enum MemorySemanticsMask semantics, float val) { \
25+
if (CHECK) \
26+
return NEW_BUILTIN(p, val); \
27+
int atomic_scope = 0, memory_order = 0; \
28+
GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \
29+
return __hip_atomic_fetch_add(p, val, memory_order, atomic_scope); \
30+
}
31+
32+
AMDGPU_ATOMIC_FP32_ADD_IMPL(global, U3AS1, 1, AMDGPU_ARCH_BETWEEN(9010, 10000),
33+
__builtin_amdgcn_global_atomic_fadd_f32)
34+
AMDGPU_ATOMIC_FP32_ADD_IMPL(local, U3AS3, 1, AMDGPU_ARCH_GEQ(8000),
35+
__builtin_amdgcn_ds_atomic_fadd_f32)
36+
AMDGPU_ATOMIC_FP32_ADD_IMPL(, , 0, AMDGPU_ARCH_BETWEEN(9400, 10000),
37+
__builtin_amdgcn_flat_atomic_fadd_f32)
38+
39+
#define AMDGPU_ATOMIC_FP64_ADD_IMPL(AS, AS_MANGLED, SUB1, SUB2, CHECK, \
40+
NEW_BUILTIN) \
2041
_CLC_DEF long \
2142
_Z29__spirv_AtomicCompareExchangeP##AS_MANGLED##lN5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagES##SUB2##_ll( \
2243
volatile AS long *, enum Scope, enum MemorySemanticsMask, \
@@ -28,6 +49,8 @@ AMDGPU_ATOMIC(_Z21__spirv_AtomicFAddEXT, float, f, __hip_atomic_fetch_add)
2849
_Z21__spirv_AtomicFAddEXTP##AS_MANGLED##dN5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagEd( \
2950
volatile AS double *p, enum Scope scope, \
3051
enum MemorySemanticsMask semantics, double val) { \
52+
if (CHECK) \
53+
return NEW_BUILTIN(p, val); \
3154
int atomic_scope = 0, memory_order = 0; \
3255
volatile AS long *int_pointer = (volatile AS long *)p; \
3356
long old_int_val = 0, new_int_val = 0; \
@@ -46,12 +69,20 @@ AMDGPU_ATOMIC(_Z21__spirv_AtomicFAddEXT, float, f, __hip_atomic_fetch_add)
4669
}
4770

4871
#ifdef cl_khr_int64_base_atomics
49-
AMDGPU_ATOMIC_FP64_ADD_IMPL(global, U3AS1, 1, 5)
50-
AMDGPU_ATOMIC_FP64_ADD_IMPL(local, U3AS3, 1, 5)
51-
AMDGPU_ATOMIC_FP64_ADD_IMPL(, , 0, 4)
72+
AMDGPU_ATOMIC_FP64_ADD_IMPL(global, U3AS1, 1, 5,
73+
AMDGPU_ARCH_BETWEEN(9010, 10000),
74+
__builtin_amdgcn_global_atomic_fadd_f64)
75+
AMDGPU_ATOMIC_FP64_ADD_IMPL(local, U3AS3, 1, 5,
76+
AMDGPU_ARCH_BETWEEN(9010, 10000),
77+
__builtin_amdgcn_ds_atomic_fadd_f64)
78+
AMDGPU_ATOMIC_FP64_ADD_IMPL(, , 0, 4, AMDGPU_ARCH_BETWEEN(9400, 10000),
79+
__builtin_amdgcn_flat_atomic_fadd_f64)
5280
#endif
5381

5482
#undef AMDGPU_ATOMIC
5583
#undef AMDGPU_ATOMIC_IMPL
84+
#undef AMDGPU_ATOMIC_FP32_ADD_IMPL
5685
#undef AMDGPU_ATOMIC_FP64_ADD_IMPL
86+
#undef AMDGPU_ARCH_GEQ
87+
#undef AMDGPU_ARCH_BETWEEN
5788
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_and.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,4 +17,6 @@ AMDGPU_ATOMIC(_Z17__spirv_AtomicAnd, unsigned long, m, __hip_atomic_fetch_and)
1717

1818
#undef AMDGPU_ATOMIC
1919
#undef AMDGPU_ATOMIC_IMPL
20+
#undef AMDGPU_ARCH_GEQ
21+
#undef AMDGPU_ARCH_BETWEEN
2022
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_cmpxchg.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -48,4 +48,6 @@ AMDGPU_ATOMIC_CMPXCHG(float, f)
4848
#undef AMDGPU_ATOMIC_IMPL
4949
#undef AMDGPU_ATOMIC_CPMXCHG
5050
#undef AMDGPU_ATOMIC_CPMXCHG_IMPL
51+
#undef AMDGPU_ARCH_GEQ
52+
#undef AMDGPU_ARCH_BETWEEN
5153
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_helpers.h

Lines changed: 11 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,10 @@
99
#include <spirv/spirv.h>
1010
#include <spirv/spirv_types.h>
1111

12+
#define AMDGPU_ARCH_GEQ(LOWER) __oclc_ISA_version >= LOWER
13+
#define AMDGPU_ARCH_BETWEEN(LOWER, UPPER) \
14+
__oclc_ISA_version >= LOWER &&__oclc_ISA_version < UPPER
15+
1216
#define GET_ATOMIC_SCOPE_AND_ORDER(IN_SCOPE, OUT_SCOPE, IN_SEMANTICS, \
1317
OUT_ORDER) \
1418
{ \
@@ -52,20 +56,18 @@
5256
} \
5357
}
5458

55-
#define AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, AS, AS_MANGLED, \
59+
#define AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, AS, AS_MANGLED, \
5660
SUB1, BUILTIN) \
57-
_CLC_DEF TYPE \
61+
_CLC_DEF TYPE \
5862
FUNC_NAME##P##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \
59-
volatile AS TYPE *p, enum Scope scope, \
60-
enum MemorySemanticsMask semantics, TYPE val) { \
61-
int atomic_scope = 0, memory_order = 0; \
62-
GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \
63-
TYPE ret = BUILTIN(p, val, memory_order, atomic_scope); \
64-
return *(TYPE *)&ret; \
63+
volatile AS TYPE *p, enum Scope scope, \
64+
enum MemorySemanticsMask semantics, TYPE val) { \
65+
int atomic_scope = 0, memory_order = 0; \
66+
GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \
67+
return BUILTIN(p, val, memory_order, atomic_scope); \
6568
}
6669

6770
#define AMDGPU_ATOMIC(FUNC_NAME, TYPE, TYPE_MANGLED, BUILTIN) \
6871
AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, global, U3AS1, 1, BUILTIN) \
6972
AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, local, U3AS3, 1, BUILTIN) \
7073
AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, , , 0, BUILTIN)
71-

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_load.cl

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -17,8 +17,7 @@
1717
enum MemorySemanticsMask semantics) { \
1818
int atomic_scope = 0, memory_order = 0; \
1919
GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \
20-
TYPE res = __hip_atomic_load(p, memory_order, atomic_scope); \
21-
return *(TYPE *)&res; \
20+
return __hip_atomic_load(p, memory_order, atomic_scope); \
2221
}
2322

2423
#define AMDGPU_ATOMIC_LOAD(TYPE, TYPE_MANGLED) \
@@ -38,4 +37,6 @@ AMDGPU_ATOMIC_LOAD(float, Kf)
3837
#undef AMDGPU_ATOMIC_IMPL
3938
#undef AMDGPU_ATOMIC_LOAD
4039
#undef AMDGPU_ATOMIC_LOAD_IMPL
40+
#undef AMDGPU_ARCH_GEQ
41+
#undef AMDGPU_ARCH_BETWEEN
4142
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_max.cl

Lines changed: 17 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -11,22 +11,33 @@
1111
#include <spirv/spirv.h>
1212
#include <spirv/spirv_types.h>
1313

14+
extern constant int __oclc_ISA_version;
15+
1416
AMDGPU_ATOMIC(_Z18__spirv_AtomicSMax, int, i, __hip_atomic_fetch_max)
1517
AMDGPU_ATOMIC(_Z18__spirv_AtomicUMax, unsigned int, j, __hip_atomic_fetch_max)
1618
AMDGPU_ATOMIC(_Z18__spirv_AtomicSMax, long, l, __hip_atomic_fetch_max)
1719
AMDGPU_ATOMIC(_Z18__spirv_AtomicUMax, unsigned long, m, __hip_atomic_fetch_max)
1820

19-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, float, f, int, i, global, U3AS1, 1, 5_ii)
20-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, float, f, int, i, local, U3AS3, 1, 5_ii)
21-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, float, f, int, i, , , 0, 4_ii)
21+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, float, f, int, i, global, U3AS1, 1, 5_ii,
22+
false, )
23+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, float, f, int, i, local, U3AS3, 1, 5_ii,
24+
false, )
25+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, float, f, int, i, , , 0, 4_ii, false, )
2226

2327
#ifdef cl_khr_int64_base_atomics
24-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, global, U3AS1, 1, 5_ll)
25-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, local, U3AS3, 1, 5_ll)
26-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, , , 0, 4_ll)
28+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, global, U3AS1, 1, 5_ll,
29+
AMDGPU_ARCH_BETWEEN(9010, 10000),
30+
__builtin_amdgcn_global_atomic_fmax_f64)
31+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, local, U3AS3, 1, 5_ll,
32+
false, )
33+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, , , 0, 4_ll,
34+
AMDGPU_ARCH_BETWEEN(9010, 10000),
35+
__builtin_amdgcn_flat_atomic_fmax_f64)
2736
#endif
2837

2938
#undef AMDGPU_ATOMIC
3039
#undef AMDGPU_ATOMIC_IMPL
40+
#undef AMDGPU_ARCH_GEQ
41+
#undef AMDGPU_ARCH_BETWEEN
3142
#undef AMDGPU_ATOMIC_FP_MINMAX_IMPL
3243
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_min.cl

Lines changed: 17 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -11,22 +11,33 @@
1111
#include <spirv/spirv.h>
1212
#include <spirv/spirv_types.h>
1313

14+
extern constant int __oclc_ISA_version;
15+
1416
AMDGPU_ATOMIC(_Z18__spirv_AtomicSMin, int, i, __hip_atomic_fetch_min)
1517
AMDGPU_ATOMIC(_Z18__spirv_AtomicUMin, unsigned int, j, __hip_atomic_fetch_min)
1618
AMDGPU_ATOMIC(_Z18__spirv_AtomicSMin, long, l, __hip_atomic_fetch_min)
1719
AMDGPU_ATOMIC(_Z18__spirv_AtomicUMin, unsigned long, m, __hip_atomic_fetch_min)
1820

19-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, float, f, int, i, global, U3AS1, 1, 5_ii)
20-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, float, f, int, i, local, U3AS3, 1, 5_ii)
21-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, float, f, int, i, , , 0, 4_ii)
21+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, float, f, int, i, global, U3AS1, 1, 5_ii,
22+
false, )
23+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, float, f, int, i, local, U3AS3, 1, 5_ii,
24+
false, )
25+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, float, f, int, i, , , 0, 4_ii, false, )
2226

2327
#ifdef cl_khr_int64_base_atomics
24-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, global, U3AS1, 1, 5_ll)
25-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, local, U3AS3, 1, 5_ll)
26-
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, , , 0, 4_ll)
28+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, global, U3AS1, 1, 5_ll,
29+
AMDGPU_ARCH_BETWEEN(9010, 10000),
30+
__builtin_amdgcn_global_atomic_fmin_f64)
31+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, local, U3AS3, 1, 5_ll,
32+
false, )
33+
AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, , , 0, 4_ll,
34+
AMDGPU_ARCH_BETWEEN(9010, 10000),
35+
__builtin_amdgcn_flat_atomic_fmin_f64)
2736
#endif
2837

2938
#undef AMDGPU_ATOMIC
3039
#undef AMDGPU_ATOMIC_IMPL
40+
#undef AMDGPU_ARCH_GEQ
41+
#undef AMDGPU_ARCH_BETWEEN
3142
#undef AMDGPU_ATOMIC_FP_MINMAX_IMPL
3243
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_minmax.h

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -10,9 +10,9 @@
1010
#include <spirv/spirv.h>
1111
#include <spirv/spirv_types.h>
1212

13-
#define AMDGPU_ATOMIC_FP_MINMAX_IMPL(OPNAME, OP, TYPE, TYPE_MANGLED, \
14-
STORAGE_TYPE, STORAGE_TYPE_MANGLED, AS, \
15-
AS_MANGLED, SUB1, SUB2) \
13+
#define AMDGPU_ATOMIC_FP_MINMAX_IMPL( \
14+
OPNAME, OP, TYPE, TYPE_MANGLED, STORAGE_TYPE, STORAGE_TYPE_MANGLED, AS, \
15+
AS_MANGLED, SUB1, SUB2, CHECK, NEW_BUILTIN) \
1616
_CLC_DEF STORAGE_TYPE \
1717
_Z29__spirv_AtomicCompareExchangeP##AS_MANGLED##STORAGE_TYPE_MANGLED##N5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagES##SUB2( \
1818
volatile AS STORAGE_TYPE *, enum Scope, enum MemorySemanticsMask, \
@@ -26,6 +26,8 @@
2626
_Z21__spirv_AtomicF##OPNAME##EXTP##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \
2727
volatile AS TYPE *p, enum Scope scope, \
2828
enum MemorySemanticsMask semantics, TYPE val) { \
29+
if (CHECK) \
30+
return NEW_BUILTIN(p, val); \
2931
int atomic_scope = 0, memory_order = 0; \
3032
volatile AS STORAGE_TYPE *int_pointer = (volatile AS STORAGE_TYPE *)p; \
3133
STORAGE_TYPE old_int_val = 0, new_int_val = 0; \
@@ -45,4 +47,3 @@
4547
\
4648
return old_val; \
4749
}
48-

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_or.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,4 +17,6 @@ AMDGPU_ATOMIC(_Z16__spirv_AtomicOr, unsigned long, m, __hip_atomic_fetch_or)
1717

1818
#undef AMDGPU_ATOMIC
1919
#undef AMDGPU_ATOMIC_IMPL
20+
#undef AMDGPU_ARCH_GEQ
21+
#undef AMDGPU_ARCH_BETWEEN
2022
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_store.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,4 +38,6 @@ AMDGPU_ATOMIC_STORE(float, f)
3838
#undef AMDGPU_ATOMIC_IMPL
3939
#undef AMDGPU_ATOMIC_STORE
4040
#undef AMDGPU_ATOMIC_STORE_IMPL
41+
#undef AMDGPU_ARCH_GEQ
42+
#undef AMDGPU_ARCH_BETWEEN
4143
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_sub.cl

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,7 @@
1818
enum MemorySemanticsMask semantics, TYPE val) { \
1919
int atomic_scope = 0, memory_order = 0; \
2020
GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order) \
21-
TYPE ret = BUILTIN(p, val, memory_order); \
22-
return *(TYPE *)&ret; \
21+
return BUILTIN(p, val, memory_order); \
2322
}
2423

2524
#define AMDGPU_ATOMIC_SUB(FUNC_NAME, TYPE, TYPE_MANGLED, BUILTIN) \
@@ -39,4 +38,6 @@ AMDGPU_ATOMIC_SUB(_Z21__spirv_AtomicFSubEXT, float, f, __atomic_fetch_sub)
3938
#undef AMDGPU_ATOMIC_IMPL
4039
#undef AMDGPU_ATOMIC_SUB
4140
#undef AMDGPU_ATOMIC_SUB_IMPL
41+
#undef AMDGPU_ARCH_GEQ
42+
#undef AMDGPU_ARCH_BETWEEN
4243
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_xchg.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,4 +22,6 @@ AMDGPU_ATOMIC(_Z22__spirv_AtomicExchange, float, f, __hip_atomic_exchange)
2222

2323
#undef AMDGPU_ATOMIC
2424
#undef AMDGPU_ATOMIC_IMPL
25+
#undef AMDGPU_ARCH_GEQ
26+
#undef AMDGPU_ARCH_BETWEEN
2527
#undef GET_ATOMIC_SCOPE_AND_ORDER

‎libclc/amdgcn-amdhsa/libspirv/atomic/atomic_xor.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,4 +17,6 @@ AMDGPU_ATOMIC(_Z17__spirv_AtomicXor, unsigned long, m, __hip_atomic_fetch_xor)
1717

1818
#undef AMDGPU_ATOMIC
1919
#undef AMDGPU_ATOMIC_IMPL
20+
#undef AMDGPU_ARCH_GEQ
21+
#undef AMDGPU_ARCH_BETWEEN
2022
#undef GET_ATOMIC_SCOPE_AND_ORDER

0 commit comments

Comments
 (0)