diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 9187a972bd6fb..0aaf4ebd0e89e 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -124,6 +124,11 @@ template struct TargetToAS { access::address_space::global_space; }; +template <> struct TargetToAS { + constexpr static access::address_space AS = + access::address_space::global_device_space; +}; + template <> struct TargetToAS { constexpr static access::address_space AS = access::address_space::local_space; diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp index 6c0be13b2c523..3509f5541f99f 100644 --- a/sycl/include/CL/sycl/atomic.hpp +++ b/sycl/include/CL/sycl/atomic.hpp @@ -173,6 +173,7 @@ namespace sycl { template class atomic { + friend class atomic; static_assert(detail::IsValidAtomicType::value, "Invalid SYCL atomic type. Valid types are: int, " "unsigned int, long, unsigned long, long long, unsigned " @@ -197,6 +198,23 @@ class atomic { "T and pointerT must be same size"); } + // Create atomic in global_space with one from global_device_space + template ::type> + atomic(const atomic &RHS) { + Ptr = RHS.Ptr; + } + + template ::type> + atomic(atomic &&RHS) { + Ptr = RHS.Ptr; + } + void store(T Operand, memory_order Order = memory_order::relaxed) { __spirv_AtomicStore( Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand); diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 43642b75a3497..f9801782c54fa 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -506,7 +506,7 @@ class __SYCL_EXPORT handler { access::placeholder IsPH> detail::enable_if_t readFromFirstAccElement(accessor Src) const { - atomic AtomicSrc = Src; + atomic AtomicSrc = Src; return AtomicSrc.load(); } @@ -529,7 +529,7 @@ class __SYCL_EXPORT handler { access::placeholder IsPH> detail::enable_if_t writeToFirstAccElement(accessor Dst, T V) const { - atomic AtomicDst = Dst; + atomic AtomicDst = Dst; AtomicDst.store(V); } diff --git a/sycl/include/CL/sycl/intel/atomic_ref.hpp b/sycl/include/CL/sycl/intel/atomic_ref.hpp index 89202cb7ca08c..f6e8d4ff68616 100644 --- a/sycl/include/CL/sycl/intel/atomic_ref.hpp +++ b/sycl/include/CL/sycl/intel/atomic_ref.hpp @@ -43,7 +43,8 @@ using IsValidAtomicType = template using IsValidAtomicAddressSpace = bool_constant; + AS == access::address_space::local_space || + AS == access::address_space::global_device_space>; // DefaultOrder parameter is limited to read-modify-write orders template @@ -138,7 +139,7 @@ class atomic_ref_base { "intel::atomic_ref does not yet support pointer types"); static_assert(detail::IsValidAtomicAddressSpace::value, "Invalid atomic address_space. Valid address spaces are: " - "global_space, local_space"); + "global_space, local_space, global_device_space"); static_assert( detail::IsValidDefaultOrder::value, "Invalid default memory_order for atomics. Valid defaults are: " diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index 1a59113d9fc18..489e34a277ff2 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -275,6 +275,22 @@ template class multi_ptr { return multi_ptr(m_Pointer - r); } + // Explicit conversion to global_space + // Only available if Space == address_space::global_device_space || + // Space == address_space::global_host_space + template ::type> + explicit + operator multi_ptr() const { + using global_pointer_t = typename detail::PtrValueType< + ElementType, access::address_space::global_space>::type *; + return multi_ptr( + reinterpret_cast(m_Pointer)); + } + // Only if Space == global_space template void testMultPtr() { global_ptr ptr_8 = global_ptr(ptr_7); host_ptr ptr_9((void *)RawPtr); global_ptr ptr_10 = global_ptr(ptr_9); + // TODO: need propagation of a7b763b26 patch to acl tool before testing + // these conversions - otherwise the test would fail on accelerator + // device during reversed translation from SPIR-V to LLVM IR + // device_ptr ptr_11(accessorData_1); + // global_ptr ptr_12 = global_ptr(ptr_11); innerFunc(wiID.get(0), ptr_1, ptr_2, local_ptr); }); diff --git a/sycl/test/regression/implicit_atomic_conversion.cpp b/sycl/test/regression/implicit_atomic_conversion.cpp new file mode 100644 index 0000000000000..625fe0b50d049 --- /dev/null +++ b/sycl/test/regression/implicit_atomic_conversion.cpp @@ -0,0 +1,29 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +#include + +using namespace cl::sycl; + +void test_conversion(queue q) { + int init = 0; + { + buffer in_buf(&init, 1); + + q.submit([&](handler &cgh) { + auto in = in_buf.template get_access(cgh); + cgh.single_task([=]() { + cl::sycl::atomic atm = in[0]; + atm.store(42); + }); + }); + } + assert(init == 42 && "verification failed"); +} + +int main() { + queue q; + test_conversion(q); +}