diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 55b0764aefe1f..8f205e8a68578 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -239,7 +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(__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; }