From beae6471dbaef20bb3bbbb86de5f0bfa34d695e8 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Wed, 15 Dec 2021 11:35:33 +0300 Subject: [PATCH 1/5] [SYCL] Add basic support for the generic_space address space The basic support is added to let the sycl::atomic_ref class template be instantiated with the address_space::generic_space address space. --- sycl/include/CL/__spirv/spirv_ops.hpp | 3 +- sycl/include/CL/sycl/access/access.hpp | 5 ++ sycl/include/CL/sycl/atomic_ref.hpp | 5 -- sycl/include/CL/sycl/multi_ptr.hpp | 79 ++++++++++++------- sycl/include/CL/sycl/pointers.hpp | 4 + .../basic_tests/atomic-ref-instantiation.cpp | 3 +- 6 files changed, 64 insertions(+), 35 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 55b0764aefe1f..252e51e1ccf37 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -239,7 +239,8 @@ extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType, #define __SPIRV_ATOMICS(macro, Arg) \ macro(__attribute__((opencl_global)), Arg) \ - macro(__attribute__((opencl_local)), Arg) + macro(__attribute__((opencl_local)), Arg) \ + macro(, Arg) __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, float) __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, double) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 37c8547616b04..5d5ec93b4ae4f 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -162,6 +162,11 @@ struct DecoratedType { using type = __OPENCL_PRIVATE_AS__ ElementType; }; +template +struct DecoratedType { + using type = ElementType; +}; + template struct DecoratedType { using type = __OPENCL_GLOBAL_AS__ ElementType; diff --git a/sycl/include/CL/sycl/atomic_ref.hpp b/sycl/include/CL/sycl/atomic_ref.hpp index 725dc8e7f27e1..c9fe30175488c 100644 --- a/sycl/include/CL/sycl/atomic_ref.hpp +++ b/sycl/include/CL/sycl/atomic_ref.hpp @@ -114,11 +114,6 @@ template <> struct bit_equal { template class atomic_ref_base { - static_assert( - AddressSpace != access::address_space::generic_space, - "access::address_space::generic_space is a valid address space but the " - "address space is not supported yet."); - static_assert( detail::IsValidAtomicRefType::value, "Invalid atomic type. Valid types are int, unsigned int, long, " diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index 07c506897b930..eefc0d0c23512 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -51,6 +51,13 @@ template class multi_ptr { multi_ptr(const multi_ptr &rhs) = default; multi_ptr(multi_ptr &&) = default; #ifdef __SYCL_DEVICE_ONLY__ + // The generic address space have no corresponding 'opencl_...' attribute and + // this constructor is considered as a duplicate for the + // multi_ptr(ElementType *pointer) one, so the check is required. + template < + access::address_space _Space = Space, + typename = typename detail::enable_if_t< + _Space == Space && Space != access::address_space::generic_space>> multi_ptr(pointer_t pointer) : m_Pointer(pointer) {} #endif @@ -71,6 +78,13 @@ template class multi_ptr { multi_ptr &operator=(multi_ptr &&) = default; #ifdef __SYCL_DEVICE_ONLY__ + // The generic address space have no corresponding 'opencl_...' attribute and + // this operator is considered as a duplicate for the + // multi_ptr &operator=(ElementType *pointer) one, so the check is required. + template < + access::address_space _Space = Space, + typename = typename detail::enable_if_t< + _Space == Space && Space != access::address_space::generic_space>> multi_ptr &operator=(pointer_t pointer) { m_Pointer = pointer; return *this; @@ -109,13 +123,14 @@ template class multi_ptr { return reinterpret_cast(m_Pointer)[index]; } - // Only if Space == global_space || global_device_space + // Only if Space == global_space || global_device_space || generic_space template > multi_ptr(accessor @@ -123,12 +138,13 @@ template class multi_ptr { m_Pointer = (pointer_t)(Accessor.get_pointer().get()); } - // Only if Space == local_space - template > + // Only if Space == local_space || generic_space + template < + int dimensions, access::mode Mode, access::placeholder isPlaceholder, + typename PropertyListT, access::address_space _Space = Space, + typename = typename detail::enable_if_t< + _Space == Space && (Space == access::address_space::generic_space || + Space == access::address_space::local_space)>> multi_ptr(accessor Accessor) @@ -154,15 +170,16 @@ template class multi_ptr { // 2. from multi_ptr to multi_ptr - // Only if Space == global_space || global_device_space and element type is - // const + // Only if Space == global_space || global_device_space || generic_space and + // element type is const template < int dimensions, access::mode Mode, access::placeholder isPlaceholder, typename PropertyListT, access::address_space _Space = Space, typename ET = ElementType, typename = typename detail::enable_if_t< _Space == Space && - (Space == access::address_space::global_space || + (Space == access::address_space::generic_space || + Space == access::address_space::global_space || Space == access::address_space::global_device_space) && std::is_const::value && std::is_same::value>> multi_ptr(accessor, dimensions, Mode, @@ -170,13 +187,15 @@ template class multi_ptr { Accessor) : multi_ptr(Accessor.get_pointer()) {} - // Only if Space == local_space and element type is const + // Only if Space == local_space || generic_space and element type is const template < int dimensions, access::mode Mode, access::placeholder isPlaceholder, typename PropertyListT, access::address_space _Space = Space, typename ET = ElementType, typename = typename detail::enable_if_t< - _Space == Space && Space == access::address_space::local_space && + _Space == Space && + (Space == access::address_space::generic_space || + Space == access::address_space::local_space) && std::is_const::value && std::is_same::value>> multi_ptr(accessor, dimensions, Mode, access::target::local, isPlaceholder, PropertyListT> @@ -373,23 +392,26 @@ template class multi_ptr { return *this; } - // Only if Space == global_space || global_device_space + // Only if Space == global_space || global_device_space || generic_space template > multi_ptr(accessor Accessor) : multi_ptr(Accessor.get_pointer()) {} - // Only if Space == local_space - template > + // Only if Space == local_space || generic_space + template < + typename ElementType, int dimensions, access::mode Mode, + typename PropertyListT, access::address_space _Space = Space, + typename = typename detail::enable_if_t< + _Space == Space && (Space == access::address_space::generic_space || + Space == access::address_space::local_space)>> multi_ptr(accessor Accessor) @@ -493,23 +515,26 @@ class multi_ptr { return *this; } - // Only if Space == global_space || global_device_space + // Only if Space == global_space || global_device_space || generic_space template > multi_ptr(accessor Accessor) : multi_ptr(Accessor.get_pointer()) {} - // Only if Space == local_space - template > + // Only if Space == local_space || generic_space + template < + typename ElementType, int dimensions, access::mode Mode, + typename PropertyListT, access::address_space _Space = Space, + typename = typename detail::enable_if_t< + _Space == Space && (Space == access::address_space::generic_space || + Space == access::address_space::local_space)>> multi_ptr(accessor Accessor) diff --git a/sycl/include/CL/sycl/pointers.hpp b/sycl/include/CL/sycl/pointers.hpp index efec74e0fd3a6..38b14ada09143 100644 --- a/sycl/include/CL/sycl/pointers.hpp +++ b/sycl/include/CL/sycl/pointers.hpp @@ -16,6 +16,10 @@ namespace sycl { template class multi_ptr; // Template specialization aliases for different pointer address spaces +template +using generic_ptr = + multi_ptr; + template using global_ptr = multi_ptr; diff --git a/sycl/test/basic_tests/atomic-ref-instantiation.cpp b/sycl/test/basic_tests/atomic-ref-instantiation.cpp index 920c89e804ae5..f00d377e61d70 100644 --- a/sycl/test/basic_tests/atomic-ref-instantiation.cpp +++ b/sycl/test/basic_tests/atomic-ref-instantiation.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -o %t.out -Xclang -verify-ignore-unexpected=note - +// expected-no-diagnostics #include struct A {}; @@ -20,6 +20,5 @@ int main() { A* p = &a; auto ref_p = sycl::atomic_ref(p); - // expected-error@CL/sycl/atomic_ref.hpp:* {{"access::address_space::generic_space is a valid address space but the address space is not supported yet."}} return 0; } From 9c4dc1e058197f095b2062a544723a155d63ce2b Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Wed, 15 Dec 2021 13:36:41 +0300 Subject: [PATCH 2/5] [SYCL][NFC] Fix formatting issues --- sycl/include/CL/__spirv/spirv_ops.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 252e51e1ccf37..8f205e8a68578 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -239,8 +239,7 @@ extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType, #define __SPIRV_ATOMICS(macro, Arg) \ macro(__attribute__((opencl_global)), Arg) \ - macro(__attribute__((opencl_local)), Arg) \ - macro(, Arg) + macro(__attribute__((opencl_local)), Arg) macro(, Arg) __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, float) __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, double) From e2df757da3dbdc742f73132e883e72cb8996e0d4 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Thu, 16 Dec 2021 17:02:19 +0300 Subject: [PATCH 3/5] [SYCL][CUDA] Add __spirv_Atomic operations for generic address space --- .../include/spirv/atomic/atomic_load.h | 13 +++-- .../include/spirv/atomic/atomic_store.h | 13 +++-- .../generic/libspirv/atomic/atomic_cmpxchg.cl | 28 ++++++++++ libclc/generic/libspirv/atomic/atomic_load.cl | 48 +++++++++------- libclc/generic/libspirv/atomic/atomic_max.cl | 26 +++++---- libclc/generic/libspirv/atomic/atomic_min.cl | 35 +++++++----- .../generic/libspirv/atomic/atomic_store.cl | 55 ++++++++++++------- libclc/generic/libspirv/atomic/atomic_xchg.cl | 45 +++++++++------ 8 files changed, 169 insertions(+), 94 deletions(-) diff --git a/libclc/generic/include/spirv/atomic/atomic_load.h b/libclc/generic/include/spirv/atomic/atomic_load.h index 30d7d4baf3af1..71dff989681b5 100644 --- a/libclc/generic/include/spirv/atomic/atomic_load.h +++ b/libclc/generic/include/spirv/atomic/atomic_load.h @@ -7,14 +7,15 @@ //===----------------------------------------------------------------------===// // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define DECL(TYPE, TYPE_MANGLED, AS, AS_MANGLED) \ - _CLC_DECL TYPE \ - _Z18__spirv_AtomicLoadPU3##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ +#define DECL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED) \ + _CLC_DECL TYPE \ + _Z18__spirv_AtomicLoadP##AS_PREFIX##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ volatile AS const TYPE *, enum Scope, enum MemorySemanticsMask); -#define DECL_AS(TYPE, TYPE_MANGLED) \ -DECL(TYPE, TYPE_MANGLED, global, AS1) \ -DECL(TYPE, TYPE_MANGLED, local, AS3) +#define DECL_AS(TYPE, TYPE_MANGLED) \ + DECL(TYPE, TYPE_MANGLED, U3, global, AS1) \ + DECL(TYPE, TYPE_MANGLED, U3, local, AS3) \ + DECL(TYPE, TYPE_MANGLED, , , ) DECL_AS(int, i) DECL_AS(unsigned int, j) diff --git a/libclc/generic/include/spirv/atomic/atomic_store.h b/libclc/generic/include/spirv/atomic/atomic_store.h index 59dde9749b054..061ca2ce39c65 100644 --- a/libclc/generic/include/spirv/atomic/atomic_store.h +++ b/libclc/generic/include/spirv/atomic/atomic_store.h @@ -7,14 +7,15 @@ //===----------------------------------------------------------------------===// // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define DECL(TYPE, TYPE_MANGLED, AS, AS_MANGLED) \ - _CLC_DECL void \ - _Z19__spirv_AtomicStorePU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ +#define DECL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED) \ + _CLC_DECL void \ + _Z19__spirv_AtomicStoreP##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ volatile AS TYPE *, enum Scope, enum MemorySemanticsMask, TYPE); -#define DECL_AS(TYPE, TYPE_MANGLED) \ -DECL(TYPE, TYPE_MANGLED, global, AS1) \ -DECL(TYPE, TYPE_MANGLED, local, AS3) +#define DECL_AS(TYPE, TYPE_MANGLED) \ + DECL(TYPE, TYPE_MANGLED, U3, global, AS1) \ + DECL(TYPE, TYPE_MANGLED, U3, local, AS3) \ + DECL(TYPE, TYPE_MANGLED, , , ) DECL_AS(int, i) DECL_AS(unsigned int, j) diff --git a/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl b/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl index c090e33fed32f..5c5f4428394dc 100644 --- a/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl +++ b/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl @@ -24,6 +24,13 @@ _Z29__spirv_AtomicCompareExchangePU3AS1iN5__spv5Scope4FlagENS1_19MemorySemantics return __sync_val_compare_and_swap(p, cmp, val); } +_CLC_DEF int +_Z29__spirv_AtomicCompareExchangePiN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_ii( + volatile int *p, enum Scope scope, enum MemorySemanticsMask eq, + enum MemorySemanticsMask neq, int val, int cmp) { + return __sync_val_compare_and_swap(p, cmp, val); +} + _CLC_DEF uint _Z29__spirv_AtomicCompareExchangePU3AS3jN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_jj( volatile local uint *p, enum Scope scope, enum MemorySemanticsMask eq, @@ -38,6 +45,13 @@ _Z29__spirv_AtomicCompareExchangePU3AS1jN5__spv5Scope4FlagENS1_19MemorySemantics return __sync_val_compare_and_swap(p, cmp, val); } +_CLC_DEF uint +_Z29__spirv_AtomicCompareExchangePjN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_jj( + volatile uint *p, enum Scope scope, enum MemorySemanticsMask eq, + enum MemorySemanticsMask neq, uint val, uint cmp) { + return __sync_val_compare_and_swap(p, cmp, val); +} + #ifdef cl_khr_int64_base_atomics _CLC_DEF long _Z29__spirv_AtomicCompareExchangePU3AS3lN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_ll( @@ -53,6 +67,13 @@ _Z29__spirv_AtomicCompareExchangePU3AS1lN5__spv5Scope4FlagENS1_19MemorySemantics return __sync_val_compare_and_swap_8(p, cmp, val); } +_CLC_DEF long +_Z29__spirv_AtomicCompareExchangePlN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_ll( + volatile long *p, enum Scope scope, enum MemorySemanticsMask eq, + enum MemorySemanticsMask neq, long val, long cmp) { + return __sync_val_compare_and_swap_8(p, cmp, val); +} + _CLC_DEF ulong _Z29__spirv_AtomicCompareExchangePU3AS3mN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_mm( volatile local ulong *p, enum Scope scope, enum MemorySemanticsMask eq, @@ -66,4 +87,11 @@ _Z29__spirv_AtomicCompareExchangePU3AS1mN5__spv5Scope4FlagENS1_19MemorySemantics enum MemorySemanticsMask neq, ulong val, ulong cmp) { return __sync_val_compare_and_swap_8(p, cmp, val); } + +_CLC_DEF ulong +_Z29__spirv_AtomicCompareExchangePmN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_mm( + volatile ulong *p, enum Scope scope, enum MemorySemanticsMask eq, + enum MemorySemanticsMask neq, ulong val, ulong cmp) { + return __sync_val_compare_and_swap_8(p, cmp, val); +} #endif diff --git a/libclc/generic/libspirv/atomic/atomic_load.cl b/libclc/generic/libspirv/atomic/atomic_load.cl index 6d8b1447261b5..3f7033bd94ad1 100644 --- a/libclc/generic/libspirv/atomic/atomic_load.cl +++ b/libclc/generic/libspirv/atomic/atomic_load.cl @@ -10,29 +10,35 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define FDECL(TYPE, PREFIX, AS, BYTE_SIZE, MEM_ORDER) \ -TYPE __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *); - -#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, PREFIX, BYTE_SIZE) \ - FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \ - FDECL(TYPE, PREFIX, AS, BYTE_SIZE, acquire) \ - FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \ - _CLC_DEF TYPE \ - _Z18__spirv_AtomicLoadPU3##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ - volatile AS const TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics) { \ - if (semantics & Acquire) { \ - return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_acquire(p); \ - } \ - if (semantics & SequentiallyConsistent) { \ - return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_seq_cst(p); \ - } \ - return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_unordered(p); \ +#define FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, MEM_ORDER) \ +TYPE __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *); + +#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_, AS_MANGLED, PREFIX, \ + BYTE_SIZE) \ + FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, unordered) \ + FDECL(TYPE, PREFIX, , , BYTE_SIZE, unordered) \ + FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, acquire) \ + FDECL(TYPE, PREFIX, , , BYTE_SIZE, acquire) \ + FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, seq_cst) \ + FDECL(TYPE, PREFIX, , , BYTE_SIZE, seq_cst) \ + _CLC_DEF TYPE \ + _Z18__spirv_AtomicLoadP##AS_PREFIX##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ + volatile AS const TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics) { \ + if (semantics & Acquire) { \ + return __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_acquire( \ + p); \ + } \ + if (semantics & SequentiallyConsistent) { \ + return __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_seq_cst(p); \ + } \ + return __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_unordered(p); \ } -#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ -IMPL(TYPE, TYPE_MANGLED, global, AS1, PREFIX, BYTE_SIZE) \ -IMPL(TYPE, TYPE_MANGLED, local, AS3, PREFIX, BYTE_SIZE) +#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ + IMPL(TYPE, TYPE_MANGLED, U3, global, global_, AS1, PREFIX, BYTE_SIZE) \ + IMPL(TYPE, TYPE_MANGLED, U3, local, local_, AS3, PREFIX, BYTE_SIZE) \ + IMPL(TYPE, TYPE_MANGLED, , , , , PREFIX, BYTE_SIZE) IMPL_AS(int, i, , 4) IMPL_AS(unsigned int, j, u, 4) diff --git a/libclc/generic/libspirv/atomic/atomic_max.cl b/libclc/generic/libspirv/atomic/atomic_max.cl index a2c37c258760c..39e8880b70a6b 100644 --- a/libclc/generic/libspirv/atomic/atomic_max.cl +++ b/libclc/generic/libspirv/atomic/atomic_max.cl @@ -10,28 +10,34 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, NAME, PREFIX, SUFFIX) \ +#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED, NAME, PREFIX, SUFFIX) \ _CLC_DEF TYPE \ - _Z18##NAME##PU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + _Z18##NAME##P##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ volatile AS TYPE *p, enum Scope scope, \ enum MemorySemanticsMask semantics, TYPE val) { \ return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \ } -IMPL(int, i, global, AS1, __spirv_AtomicSMax, , max) -IMPL(unsigned int, j, global, AS1, __spirv_AtomicUMax, , umax) -IMPL(int, i, local, AS3, __spirv_AtomicSMax, , max) -IMPL(unsigned int, j, local, AS3, __spirv_AtomicUMax, , umax) +IMPL(int, i, U3, global, AS1, __spirv_AtomicSMax, , max) +IMPL(unsigned int, j, U3, global, AS1, __spirv_AtomicUMax, , umax) +IMPL(int, i, U3, local, AS3, __spirv_AtomicSMax, , max) +IMPL(unsigned int, j, U3, local, AS3, __spirv_AtomicUMax, , umax) +IMPL(int, i, , , , __spirv_AtomicSMax, , max) +IMPL(unsigned int, j, , , , __spirv_AtomicUMax, , umax) #ifdef cl_khr_int64_extended_atomics unsigned long __clc__sync_fetch_and_max_local_8(volatile local long *, long); unsigned long __clc__sync_fetch_and_max_global_8(volatile global long *, long); +unsigned long __clc__sync_fetch_and_max_8(volatile long *, long); unsigned long __clc__sync_fetch_and_umax_local_8(volatile local unsigned long *, unsigned long); unsigned long __clc__sync_fetch_and_umax_global_8(volatile global unsigned long *, unsigned long); +unsigned long __clc__sync_fetch_and_umax_8(volatile unsigned long *, unsigned long); -IMPL(long, l, global, AS1, __spirv_AtomicSMax, __clc, max_global_8) -IMPL(unsigned long, m, global, AS1, __spirv_AtomicUMax, __clc, umax_global_8) -IMPL(long, l, local, AS3, __spirv_AtomicSMax, __clc, max_local_8) -IMPL(unsigned long, m, local, AS3, __spirv_AtomicUMax, __clc, umax_local_8) +IMPL(long, l, U3, global, AS1, __spirv_AtomicSMax, __clc, max_global_8) +IMPL(unsigned long, m, U3, global, AS1, __spirv_AtomicUMax, __clc, umax_global_8) +IMPL(long, l, U3, local, AS3, __spirv_AtomicSMax, __clc, max_local_8) +IMPL(unsigned long, m, U3, local, AS3, __spirv_AtomicUMax, __clc, umax_local_8) +IMPL(long, l, , , , __spirv_AtomicSMax, __clc, max_8) +IMPL(unsigned long, m, , , , __spirv_AtomicUMax, __clc, umax_8) #endif #undef IMPL diff --git a/libclc/generic/libspirv/atomic/atomic_min.cl b/libclc/generic/libspirv/atomic/atomic_min.cl index 3e3c4dfdf727c..ac4c77d372df4 100644 --- a/libclc/generic/libspirv/atomic/atomic_min.cl +++ b/libclc/generic/libspirv/atomic/atomic_min.cl @@ -10,28 +10,35 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, NAME, PREFIX, SUFFIX) \ - _CLC_DEF TYPE \ - _Z18##NAME##PU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ - volatile AS TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics, TYPE val) { \ - return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \ +#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED, NAME, PREFIX, \ + SUFFIX) \ + _CLC_DEF TYPE \ + _Z18##NAME##P##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + volatile AS TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics, TYPE val) { \ + return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \ } -IMPL(int, i, global, AS1, __spirv_AtomicSMin, , min) -IMPL(unsigned int, j, global, AS1, __spirv_AtomicUMin, , umin) -IMPL(int, i, local, AS3, __spirv_AtomicSMin, , min) -IMPL(unsigned int, j, local, AS3, __spirv_AtomicUMin, , umin) +IMPL(int, i, U3, global, AS1, __spirv_AtomicSMin, , min) +IMPL(unsigned int, j, U3, global, AS1, __spirv_AtomicUMin, , umin) +IMPL(int, i, U3, local, AS3, __spirv_AtomicSMin, , min) +IMPL(unsigned int, j, U3, local, AS3, __spirv_AtomicUMin, , umin) +IMPL(int, i, , , , __spirv_AtomicSMin, , min) +IMPL(unsigned int, j, , , , __spirv_AtomicUMin, , umin) #ifdef cl_khr_int64_extended_atomics unsigned long __clc__sync_fetch_and_min_local_8(volatile local long *, long); unsigned long __clc__sync_fetch_and_min_global_8(volatile global long *, long); +unsigned long __clc__sync_fetch_and_min_8(volatile long *, long); unsigned long __clc__sync_fetch_and_umin_local_8(volatile local unsigned long *, unsigned long); unsigned long __clc__sync_fetch_and_umin_global_8(volatile global unsigned long *, unsigned long); +unsigned long __clc__sync_fetch_and_umin_8(volatile unsigned long *, unsigned long); -IMPL(long, l, global, AS1, __spirv_AtomicSMin, __clc, min_global_8) -IMPL(unsigned long, m, global, AS1, __spirv_AtomicUMin, __clc, umin_global_8) -IMPL(long, l, local, AS3, __spirv_AtomicSMin, __clc, min_local_8) -IMPL(unsigned long, m, local, AS3, __spirv_AtomicUMin, __clc, umin_local_8) +IMPL(long, l, U3, global, AS1, __spirv_AtomicSMin, __clc, min_global_8) +IMPL(unsigned long, m, U3, global, AS1, __spirv_AtomicUMin, __clc, umin_global_8) +IMPL(long, l, U3, local, AS3, __spirv_AtomicSMin, __clc, min_local_8) +IMPL(unsigned long, m, U3, local, AS3, __spirv_AtomicUMin, __clc, umin_local_8) +IMPL(long, l, , , , __spirv_AtomicSMin, __clc, min_8) +IMPL(unsigned long, m, , , , __spirv_AtomicUMin, __clc, umin_8) #endif #undef IMPL diff --git a/libclc/generic/libspirv/atomic/atomic_store.cl b/libclc/generic/libspirv/atomic/atomic_store.cl index d24672957ddc4..685504bf28224 100644 --- a/libclc/generic/libspirv/atomic/atomic_store.cl +++ b/libclc/generic/libspirv/atomic/atomic_store.cl @@ -26,29 +26,42 @@ _Z19__spirv_AtomicStorePU3AS3fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE (volatile local uint *)p, scope, semantics, as_uint(val)); } -#define FDECL(TYPE, PREFIX, AS, BYTE_SIZE, MEM_ORDER) \ -TYPE __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *, TYPE); - -#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, PREFIX, BYTE_SIZE) \ - FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \ - FDECL(TYPE, PREFIX, AS, BYTE_SIZE, release) \ - FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \ - _CLC_DEF void \ - _Z19__spirv_AtomicStorePU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ - volatile AS TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics, TYPE val) { \ - if (semantics == Release) { \ - __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_release(p, val); \ - } else if (semantics == SequentiallyConsistent) { \ - __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_seq_cst(p, val); \ - } else { \ - __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_unordered(p, val); \ - } \ +_CLC_DEF void +_Z19__spirv_AtomicStorePfN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf( + volatile float *p, enum Scope scope, + enum MemorySemanticsMask semantics, float val) { + _Z19__spirv_AtomicStorePjN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEj( + (volatile uint *)p, scope, semantics, as_uint(val)); +} + +#define FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, MEM_ORDER) \ +TYPE __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *, TYPE); + +#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_, AS_MANGLED, PREFIX, \ + BYTE_SIZE) \ + FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, unordered) \ + FDECL(TYPE, PREFIX, , , BYTE_SIZE, unordered) \ + FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, release) \ + FDECL(TYPE, PREFIX, , , BYTE_SIZE, release) \ + FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, seq_cst) \ + FDECL(TYPE, PREFIX, , , BYTE_SIZE, seq_cst) \ + _CLC_DEF void \ + _Z19__spirv_AtomicStoreP##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + volatile AS TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics, TYPE val) { \ + if (semantics == Release) { \ + __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_release(p, val); \ + } else if (semantics == SequentiallyConsistent) { \ + __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_seq_cst(p, val); \ + } else { \ + __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_unordered(p, val); \ + } \ } -#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ -IMPL(TYPE, TYPE_MANGLED, global, AS1, PREFIX, BYTE_SIZE) \ -IMPL(TYPE, TYPE_MANGLED, local, AS3, PREFIX, BYTE_SIZE) +#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ + IMPL(TYPE, TYPE_MANGLED, U3, global, global_, AS1, PREFIX, BYTE_SIZE) \ + IMPL(TYPE, TYPE_MANGLED, U3, local, local_, AS3, PREFIX, BYTE_SIZE) \ + IMPL(TYPE, TYPE_MANGLED, , , , , PREFIX, BYTE_SIZE) IMPL_AS(int, i, , 4) IMPL_AS(unsigned int, j, u, 4) diff --git a/libclc/generic/libspirv/atomic/atomic_xchg.cl b/libclc/generic/libspirv/atomic/atomic_xchg.cl index 6f22977caa530..f2b9c6b24d92f 100644 --- a/libclc/generic/libspirv/atomic/atomic_xchg.cl +++ b/libclc/generic/libspirv/atomic/atomic_xchg.cl @@ -10,6 +10,21 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. +#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED, FN_NAME) \ + _CLC_DEF TYPE \ + _Z22__spirv_AtomicExchangeP##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + volatile AS TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics, TYPE val) { \ + return FN_NAME(p, val); \ + } + +IMPL(int, i, U3, global, AS1, __sync_swap_4) +IMPL(unsigned int, j, U3, global, AS1, __sync_swap_4) +IMPL(int, i, U3, local, AS3, __sync_swap_4) +IMPL(unsigned int, j, U3, local, AS3, __sync_swap_4) +IMPL(int, i, , , , __sync_swap_4) +IMPL(unsigned int, j, , , , __sync_swap_4) + _CLC_DEF float _Z22__spirv_AtomicExchangePU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf( volatile global float *p, enum Scope scope, @@ -28,23 +43,21 @@ _Z22__spirv_AtomicExchangePU3AS3fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4Fl (volatile local uint *)p, scope, semantics, as_uint(val))); } -#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, FN_NAME) \ - _CLC_DEF TYPE \ - _Z22__spirv_AtomicExchangePU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ - volatile AS TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics, TYPE val) { \ - return FN_NAME(p, val); \ - } - -IMPL(int, i, global, AS1, __sync_swap_4) -IMPL(unsigned int, j, global, AS1, __sync_swap_4) -IMPL(int, i, local, AS3, __sync_swap_4) -IMPL(unsigned int, j, local, AS3, __sync_swap_4) +_CLC_DEF float +_Z22__spirv_AtomicExchangePfN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf( + volatile float *p, enum Scope scope, + enum MemorySemanticsMask semantics, float val) { + return as_float( + _Z22__spirv_AtomicExchangePjN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEj( + (volatile uint *)p, scope, semantics, as_uint(val))); +} #ifdef cl_khr_int64_base_atomics -IMPL(long, l, global, AS1, __sync_swap_8) -IMPL(unsigned long, m, global, AS1, __sync_swap_8) -IMPL(long, l, local, AS3, __sync_swap_8) -IMPL(unsigned long, m, local, AS3, __sync_swap_8) +IMPL(long, l, U3, global, AS1, __sync_swap_8) +IMPL(unsigned long, m, U3, global, AS1, __sync_swap_8) +IMPL(long, l, U3, local, AS3, __sync_swap_8) +IMPL(unsigned long, m, U3, local, AS3, __sync_swap_8) +IMPL(long, l, , , , __sync_swap_8) +IMPL(unsigned long, m, , , , __sync_swap_8) #endif #undef IMPL From 3fe6c188ea7723b86088d33fdc9d5f082411c6ff Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Mon, 20 Dec 2021 10:35:03 +0300 Subject: [PATCH 4/5] Revert "[SYCL][CUDA] Add __spirv_Atomic operations for generic address space" This reverts commit e2df757da3dbdc742f73132e883e72cb8996e0d4. --- .../include/spirv/atomic/atomic_load.h | 13 ++--- .../include/spirv/atomic/atomic_store.h | 13 ++--- .../generic/libspirv/atomic/atomic_cmpxchg.cl | 28 ---------- libclc/generic/libspirv/atomic/atomic_load.cl | 48 +++++++--------- libclc/generic/libspirv/atomic/atomic_max.cl | 26 ++++----- libclc/generic/libspirv/atomic/atomic_min.cl | 35 +++++------- .../generic/libspirv/atomic/atomic_store.cl | 55 +++++++------------ libclc/generic/libspirv/atomic/atomic_xchg.cl | 45 ++++++--------- 8 files changed, 94 insertions(+), 169 deletions(-) diff --git a/libclc/generic/include/spirv/atomic/atomic_load.h b/libclc/generic/include/spirv/atomic/atomic_load.h index 71dff989681b5..30d7d4baf3af1 100644 --- a/libclc/generic/include/spirv/atomic/atomic_load.h +++ b/libclc/generic/include/spirv/atomic/atomic_load.h @@ -7,15 +7,14 @@ //===----------------------------------------------------------------------===// // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define DECL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED) \ - _CLC_DECL TYPE \ - _Z18__spirv_AtomicLoadP##AS_PREFIX##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ +#define DECL(TYPE, TYPE_MANGLED, AS, AS_MANGLED) \ + _CLC_DECL TYPE \ + _Z18__spirv_AtomicLoadPU3##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ volatile AS const TYPE *, enum Scope, enum MemorySemanticsMask); -#define DECL_AS(TYPE, TYPE_MANGLED) \ - DECL(TYPE, TYPE_MANGLED, U3, global, AS1) \ - DECL(TYPE, TYPE_MANGLED, U3, local, AS3) \ - DECL(TYPE, TYPE_MANGLED, , , ) +#define DECL_AS(TYPE, TYPE_MANGLED) \ +DECL(TYPE, TYPE_MANGLED, global, AS1) \ +DECL(TYPE, TYPE_MANGLED, local, AS3) DECL_AS(int, i) DECL_AS(unsigned int, j) diff --git a/libclc/generic/include/spirv/atomic/atomic_store.h b/libclc/generic/include/spirv/atomic/atomic_store.h index 061ca2ce39c65..59dde9749b054 100644 --- a/libclc/generic/include/spirv/atomic/atomic_store.h +++ b/libclc/generic/include/spirv/atomic/atomic_store.h @@ -7,15 +7,14 @@ //===----------------------------------------------------------------------===// // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define DECL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED) \ - _CLC_DECL void \ - _Z19__spirv_AtomicStoreP##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ +#define DECL(TYPE, TYPE_MANGLED, AS, AS_MANGLED) \ + _CLC_DECL void \ + _Z19__spirv_AtomicStorePU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ volatile AS TYPE *, enum Scope, enum MemorySemanticsMask, TYPE); -#define DECL_AS(TYPE, TYPE_MANGLED) \ - DECL(TYPE, TYPE_MANGLED, U3, global, AS1) \ - DECL(TYPE, TYPE_MANGLED, U3, local, AS3) \ - DECL(TYPE, TYPE_MANGLED, , , ) +#define DECL_AS(TYPE, TYPE_MANGLED) \ +DECL(TYPE, TYPE_MANGLED, global, AS1) \ +DECL(TYPE, TYPE_MANGLED, local, AS3) DECL_AS(int, i) DECL_AS(unsigned int, j) diff --git a/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl b/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl index 5c5f4428394dc..c090e33fed32f 100644 --- a/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl +++ b/libclc/generic/libspirv/atomic/atomic_cmpxchg.cl @@ -24,13 +24,6 @@ _Z29__spirv_AtomicCompareExchangePU3AS1iN5__spv5Scope4FlagENS1_19MemorySemantics return __sync_val_compare_and_swap(p, cmp, val); } -_CLC_DEF int -_Z29__spirv_AtomicCompareExchangePiN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_ii( - volatile int *p, enum Scope scope, enum MemorySemanticsMask eq, - enum MemorySemanticsMask neq, int val, int cmp) { - return __sync_val_compare_and_swap(p, cmp, val); -} - _CLC_DEF uint _Z29__spirv_AtomicCompareExchangePU3AS3jN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_jj( volatile local uint *p, enum Scope scope, enum MemorySemanticsMask eq, @@ -45,13 +38,6 @@ _Z29__spirv_AtomicCompareExchangePU3AS1jN5__spv5Scope4FlagENS1_19MemorySemantics return __sync_val_compare_and_swap(p, cmp, val); } -_CLC_DEF uint -_Z29__spirv_AtomicCompareExchangePjN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_jj( - volatile uint *p, enum Scope scope, enum MemorySemanticsMask eq, - enum MemorySemanticsMask neq, uint val, uint cmp) { - return __sync_val_compare_and_swap(p, cmp, val); -} - #ifdef cl_khr_int64_base_atomics _CLC_DEF long _Z29__spirv_AtomicCompareExchangePU3AS3lN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_ll( @@ -67,13 +53,6 @@ _Z29__spirv_AtomicCompareExchangePU3AS1lN5__spv5Scope4FlagENS1_19MemorySemantics return __sync_val_compare_and_swap_8(p, cmp, val); } -_CLC_DEF long -_Z29__spirv_AtomicCompareExchangePlN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_ll( - volatile long *p, enum Scope scope, enum MemorySemanticsMask eq, - enum MemorySemanticsMask neq, long val, long cmp) { - return __sync_val_compare_and_swap_8(p, cmp, val); -} - _CLC_DEF ulong _Z29__spirv_AtomicCompareExchangePU3AS3mN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_mm( volatile local ulong *p, enum Scope scope, enum MemorySemanticsMask eq, @@ -87,11 +66,4 @@ _Z29__spirv_AtomicCompareExchangePU3AS1mN5__spv5Scope4FlagENS1_19MemorySemantics enum MemorySemanticsMask neq, ulong val, ulong cmp) { return __sync_val_compare_and_swap_8(p, cmp, val); } - -_CLC_DEF ulong -_Z29__spirv_AtomicCompareExchangePmN5__spv5Scope4FlagENS0_19MemorySemanticsMask4FlagES4_mm( - volatile ulong *p, enum Scope scope, enum MemorySemanticsMask eq, - enum MemorySemanticsMask neq, ulong val, ulong cmp) { - return __sync_val_compare_and_swap_8(p, cmp, val); -} #endif diff --git a/libclc/generic/libspirv/atomic/atomic_load.cl b/libclc/generic/libspirv/atomic/atomic_load.cl index 3f7033bd94ad1..6d8b1447261b5 100644 --- a/libclc/generic/libspirv/atomic/atomic_load.cl +++ b/libclc/generic/libspirv/atomic/atomic_load.cl @@ -10,35 +10,29 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, MEM_ORDER) \ -TYPE __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *); - -#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_, AS_MANGLED, PREFIX, \ - BYTE_SIZE) \ - FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, unordered) \ - FDECL(TYPE, PREFIX, , , BYTE_SIZE, unordered) \ - FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, acquire) \ - FDECL(TYPE, PREFIX, , , BYTE_SIZE, acquire) \ - FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, seq_cst) \ - FDECL(TYPE, PREFIX, , , BYTE_SIZE, seq_cst) \ - _CLC_DEF TYPE \ - _Z18__spirv_AtomicLoadP##AS_PREFIX##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ - volatile AS const TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics) { \ - if (semantics & Acquire) { \ - return __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_acquire( \ - p); \ - } \ - if (semantics & SequentiallyConsistent) { \ - return __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_seq_cst(p); \ - } \ - return __clc__atomic_##PREFIX##load_##AS_##BYTE_SIZE##_unordered(p); \ +#define FDECL(TYPE, PREFIX, AS, BYTE_SIZE, MEM_ORDER) \ +TYPE __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *); + +#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, PREFIX, BYTE_SIZE) \ + FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \ + FDECL(TYPE, PREFIX, AS, BYTE_SIZE, acquire) \ + FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \ + _CLC_DEF TYPE \ + _Z18__spirv_AtomicLoadPU3##AS_MANGLED##K##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE( \ + volatile AS const TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics) { \ + if (semantics & Acquire) { \ + return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_acquire(p); \ + } \ + if (semantics & SequentiallyConsistent) { \ + return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_seq_cst(p); \ + } \ + return __clc__atomic_##PREFIX##load_##AS##_##BYTE_SIZE##_unordered(p); \ } -#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ - IMPL(TYPE, TYPE_MANGLED, U3, global, global_, AS1, PREFIX, BYTE_SIZE) \ - IMPL(TYPE, TYPE_MANGLED, U3, local, local_, AS3, PREFIX, BYTE_SIZE) \ - IMPL(TYPE, TYPE_MANGLED, , , , , PREFIX, BYTE_SIZE) +#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ +IMPL(TYPE, TYPE_MANGLED, global, AS1, PREFIX, BYTE_SIZE) \ +IMPL(TYPE, TYPE_MANGLED, local, AS3, PREFIX, BYTE_SIZE) IMPL_AS(int, i, , 4) IMPL_AS(unsigned int, j, u, 4) diff --git a/libclc/generic/libspirv/atomic/atomic_max.cl b/libclc/generic/libspirv/atomic/atomic_max.cl index 39e8880b70a6b..a2c37c258760c 100644 --- a/libclc/generic/libspirv/atomic/atomic_max.cl +++ b/libclc/generic/libspirv/atomic/atomic_max.cl @@ -10,34 +10,28 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED, NAME, PREFIX, SUFFIX) \ +#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, NAME, PREFIX, SUFFIX) \ _CLC_DEF TYPE \ - _Z18##NAME##P##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + _Z18##NAME##PU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ volatile AS TYPE *p, enum Scope scope, \ enum MemorySemanticsMask semantics, TYPE val) { \ return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \ } -IMPL(int, i, U3, global, AS1, __spirv_AtomicSMax, , max) -IMPL(unsigned int, j, U3, global, AS1, __spirv_AtomicUMax, , umax) -IMPL(int, i, U3, local, AS3, __spirv_AtomicSMax, , max) -IMPL(unsigned int, j, U3, local, AS3, __spirv_AtomicUMax, , umax) -IMPL(int, i, , , , __spirv_AtomicSMax, , max) -IMPL(unsigned int, j, , , , __spirv_AtomicUMax, , umax) +IMPL(int, i, global, AS1, __spirv_AtomicSMax, , max) +IMPL(unsigned int, j, global, AS1, __spirv_AtomicUMax, , umax) +IMPL(int, i, local, AS3, __spirv_AtomicSMax, , max) +IMPL(unsigned int, j, local, AS3, __spirv_AtomicUMax, , umax) #ifdef cl_khr_int64_extended_atomics unsigned long __clc__sync_fetch_and_max_local_8(volatile local long *, long); unsigned long __clc__sync_fetch_and_max_global_8(volatile global long *, long); -unsigned long __clc__sync_fetch_and_max_8(volatile long *, long); unsigned long __clc__sync_fetch_and_umax_local_8(volatile local unsigned long *, unsigned long); unsigned long __clc__sync_fetch_and_umax_global_8(volatile global unsigned long *, unsigned long); -unsigned long __clc__sync_fetch_and_umax_8(volatile unsigned long *, unsigned long); -IMPL(long, l, U3, global, AS1, __spirv_AtomicSMax, __clc, max_global_8) -IMPL(unsigned long, m, U3, global, AS1, __spirv_AtomicUMax, __clc, umax_global_8) -IMPL(long, l, U3, local, AS3, __spirv_AtomicSMax, __clc, max_local_8) -IMPL(unsigned long, m, U3, local, AS3, __spirv_AtomicUMax, __clc, umax_local_8) -IMPL(long, l, , , , __spirv_AtomicSMax, __clc, max_8) -IMPL(unsigned long, m, , , , __spirv_AtomicUMax, __clc, umax_8) +IMPL(long, l, global, AS1, __spirv_AtomicSMax, __clc, max_global_8) +IMPL(unsigned long, m, global, AS1, __spirv_AtomicUMax, __clc, umax_global_8) +IMPL(long, l, local, AS3, __spirv_AtomicSMax, __clc, max_local_8) +IMPL(unsigned long, m, local, AS3, __spirv_AtomicUMax, __clc, umax_local_8) #endif #undef IMPL diff --git a/libclc/generic/libspirv/atomic/atomic_min.cl b/libclc/generic/libspirv/atomic/atomic_min.cl index ac4c77d372df4..3e3c4dfdf727c 100644 --- a/libclc/generic/libspirv/atomic/atomic_min.cl +++ b/libclc/generic/libspirv/atomic/atomic_min.cl @@ -10,35 +10,28 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED, NAME, PREFIX, \ - SUFFIX) \ - _CLC_DEF TYPE \ - _Z18##NAME##P##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ - volatile AS TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics, TYPE val) { \ - return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \ +#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, NAME, PREFIX, SUFFIX) \ + _CLC_DEF TYPE \ + _Z18##NAME##PU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + volatile AS TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics, TYPE val) { \ + return PREFIX##__sync_fetch_and_##SUFFIX(p, val); \ } -IMPL(int, i, U3, global, AS1, __spirv_AtomicSMin, , min) -IMPL(unsigned int, j, U3, global, AS1, __spirv_AtomicUMin, , umin) -IMPL(int, i, U3, local, AS3, __spirv_AtomicSMin, , min) -IMPL(unsigned int, j, U3, local, AS3, __spirv_AtomicUMin, , umin) -IMPL(int, i, , , , __spirv_AtomicSMin, , min) -IMPL(unsigned int, j, , , , __spirv_AtomicUMin, , umin) +IMPL(int, i, global, AS1, __spirv_AtomicSMin, , min) +IMPL(unsigned int, j, global, AS1, __spirv_AtomicUMin, , umin) +IMPL(int, i, local, AS3, __spirv_AtomicSMin, , min) +IMPL(unsigned int, j, local, AS3, __spirv_AtomicUMin, , umin) #ifdef cl_khr_int64_extended_atomics unsigned long __clc__sync_fetch_and_min_local_8(volatile local long *, long); unsigned long __clc__sync_fetch_and_min_global_8(volatile global long *, long); -unsigned long __clc__sync_fetch_and_min_8(volatile long *, long); unsigned long __clc__sync_fetch_and_umin_local_8(volatile local unsigned long *, unsigned long); unsigned long __clc__sync_fetch_and_umin_global_8(volatile global unsigned long *, unsigned long); -unsigned long __clc__sync_fetch_and_umin_8(volatile unsigned long *, unsigned long); -IMPL(long, l, U3, global, AS1, __spirv_AtomicSMin, __clc, min_global_8) -IMPL(unsigned long, m, U3, global, AS1, __spirv_AtomicUMin, __clc, umin_global_8) -IMPL(long, l, U3, local, AS3, __spirv_AtomicSMin, __clc, min_local_8) -IMPL(unsigned long, m, U3, local, AS3, __spirv_AtomicUMin, __clc, umin_local_8) -IMPL(long, l, , , , __spirv_AtomicSMin, __clc, min_8) -IMPL(unsigned long, m, , , , __spirv_AtomicUMin, __clc, umin_8) +IMPL(long, l, global, AS1, __spirv_AtomicSMin, __clc, min_global_8) +IMPL(unsigned long, m, global, AS1, __spirv_AtomicUMin, __clc, umin_global_8) +IMPL(long, l, local, AS3, __spirv_AtomicSMin, __clc, min_local_8) +IMPL(unsigned long, m, local, AS3, __spirv_AtomicUMin, __clc, umin_local_8) #endif #undef IMPL diff --git a/libclc/generic/libspirv/atomic/atomic_store.cl b/libclc/generic/libspirv/atomic/atomic_store.cl index 685504bf28224..d24672957ddc4 100644 --- a/libclc/generic/libspirv/atomic/atomic_store.cl +++ b/libclc/generic/libspirv/atomic/atomic_store.cl @@ -26,42 +26,29 @@ _Z19__spirv_AtomicStorePU3AS3fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE (volatile local uint *)p, scope, semantics, as_uint(val)); } -_CLC_DEF void -_Z19__spirv_AtomicStorePfN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf( - volatile float *p, enum Scope scope, - enum MemorySemanticsMask semantics, float val) { - _Z19__spirv_AtomicStorePjN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEj( - (volatile uint *)p, scope, semantics, as_uint(val)); -} - -#define FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, MEM_ORDER) \ -TYPE __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *, TYPE); - -#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_, AS_MANGLED, PREFIX, \ - BYTE_SIZE) \ - FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, unordered) \ - FDECL(TYPE, PREFIX, , , BYTE_SIZE, unordered) \ - FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, release) \ - FDECL(TYPE, PREFIX, , , BYTE_SIZE, release) \ - FDECL(TYPE, PREFIX, AS, AS_, BYTE_SIZE, seq_cst) \ - FDECL(TYPE, PREFIX, , , BYTE_SIZE, seq_cst) \ - _CLC_DEF void \ - _Z19__spirv_AtomicStoreP##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ - volatile AS TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics, TYPE val) { \ - if (semantics == Release) { \ - __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_release(p, val); \ - } else if (semantics == SequentiallyConsistent) { \ - __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_seq_cst(p, val); \ - } else { \ - __clc__atomic_##PREFIX##store_##AS_##BYTE_SIZE##_unordered(p, val); \ - } \ +#define FDECL(TYPE, PREFIX, AS, BYTE_SIZE, MEM_ORDER) \ +TYPE __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_##MEM_ORDER(volatile AS const TYPE *, TYPE); + +#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, PREFIX, BYTE_SIZE) \ + FDECL(TYPE, PREFIX, AS, BYTE_SIZE, unordered) \ + FDECL(TYPE, PREFIX, AS, BYTE_SIZE, release) \ + FDECL(TYPE, PREFIX, AS, BYTE_SIZE, seq_cst) \ + _CLC_DEF void \ + _Z19__spirv_AtomicStorePU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + volatile AS TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics, TYPE val) { \ + if (semantics == Release) { \ + __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_release(p, val); \ + } else if (semantics == SequentiallyConsistent) { \ + __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_seq_cst(p, val); \ + } else { \ + __clc__atomic_##PREFIX##store_##AS##_##BYTE_SIZE##_unordered(p, val); \ + } \ } -#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ - IMPL(TYPE, TYPE_MANGLED, U3, global, global_, AS1, PREFIX, BYTE_SIZE) \ - IMPL(TYPE, TYPE_MANGLED, U3, local, local_, AS3, PREFIX, BYTE_SIZE) \ - IMPL(TYPE, TYPE_MANGLED, , , , , PREFIX, BYTE_SIZE) +#define IMPL_AS(TYPE, TYPE_MANGLED, PREFIX, BYTE_SIZE) \ +IMPL(TYPE, TYPE_MANGLED, global, AS1, PREFIX, BYTE_SIZE) \ +IMPL(TYPE, TYPE_MANGLED, local, AS3, PREFIX, BYTE_SIZE) IMPL_AS(int, i, , 4) IMPL_AS(unsigned int, j, u, 4) diff --git a/libclc/generic/libspirv/atomic/atomic_xchg.cl b/libclc/generic/libspirv/atomic/atomic_xchg.cl index f2b9c6b24d92f..6f22977caa530 100644 --- a/libclc/generic/libspirv/atomic/atomic_xchg.cl +++ b/libclc/generic/libspirv/atomic/atomic_xchg.cl @@ -10,21 +10,6 @@ // TODO: Stop manually mangling this name. Need C++ namespaces to get the exact mangling. -#define IMPL(TYPE, TYPE_MANGLED, AS_PREFIX, AS, AS_MANGLED, FN_NAME) \ - _CLC_DEF TYPE \ - _Z22__spirv_AtomicExchangeP##AS_PREFIX##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ - volatile AS TYPE *p, enum Scope scope, \ - enum MemorySemanticsMask semantics, TYPE val) { \ - return FN_NAME(p, val); \ - } - -IMPL(int, i, U3, global, AS1, __sync_swap_4) -IMPL(unsigned int, j, U3, global, AS1, __sync_swap_4) -IMPL(int, i, U3, local, AS3, __sync_swap_4) -IMPL(unsigned int, j, U3, local, AS3, __sync_swap_4) -IMPL(int, i, , , , __sync_swap_4) -IMPL(unsigned int, j, , , , __sync_swap_4) - _CLC_DEF float _Z22__spirv_AtomicExchangePU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf( volatile global float *p, enum Scope scope, @@ -43,21 +28,23 @@ _Z22__spirv_AtomicExchangePU3AS3fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4Fl (volatile local uint *)p, scope, semantics, as_uint(val))); } -_CLC_DEF float -_Z22__spirv_AtomicExchangePfN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf( - volatile float *p, enum Scope scope, - enum MemorySemanticsMask semantics, float val) { - return as_float( - _Z22__spirv_AtomicExchangePjN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEj( - (volatile uint *)p, scope, semantics, as_uint(val))); -} +#define IMPL(TYPE, TYPE_MANGLED, AS, AS_MANGLED, FN_NAME) \ + _CLC_DEF TYPE \ + _Z22__spirv_AtomicExchangePU3##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \ + volatile AS TYPE *p, enum Scope scope, \ + enum MemorySemanticsMask semantics, TYPE val) { \ + return FN_NAME(p, val); \ + } + +IMPL(int, i, global, AS1, __sync_swap_4) +IMPL(unsigned int, j, global, AS1, __sync_swap_4) +IMPL(int, i, local, AS3, __sync_swap_4) +IMPL(unsigned int, j, local, AS3, __sync_swap_4) #ifdef cl_khr_int64_base_atomics -IMPL(long, l, U3, global, AS1, __sync_swap_8) -IMPL(unsigned long, m, U3, global, AS1, __sync_swap_8) -IMPL(long, l, U3, local, AS3, __sync_swap_8) -IMPL(unsigned long, m, U3, local, AS3, __sync_swap_8) -IMPL(long, l, , , , __sync_swap_8) -IMPL(unsigned long, m, , , , __sync_swap_8) +IMPL(long, l, global, AS1, __sync_swap_8) +IMPL(unsigned long, m, global, AS1, __sync_swap_8) +IMPL(long, l, local, AS3, __sync_swap_8) +IMPL(unsigned long, m, local, AS3, __sync_swap_8) #endif #undef IMPL From 5c5c737155c6826c394fd686019f5fda8ace6619 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Tue, 21 Dec 2021 10:43:08 +0300 Subject: [PATCH 5/5] [SYCL] Trigger precommit tests