From 27e7b6ab5f07e66fc019d2abff2e53a7a88ca2a8 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 10 Jun 2020 13:17:11 +0300 Subject: [PATCH 1/7] [SYCL] Add device_ptr and host_ptr Justification: Currently a device backend can't trace from where a pointer allocated by USM comes: it can be either allocated on host or on device (it's just a pointer in OpenCL global address space). On FPGAs at least we can generate more efficient hardware code if the user tells us where the pointer can point. With this change users can create multi_ptr with specialized address space global_host or global_device that will proved to the compiler additional information to process load-store optimizations. Accessor pointers shall be also moved to global_device address spaces - otherwise backend would assume, that a pointer in global address space can access both host and device memory. Previously there were added global_device in global_host address spaces for OpenCL/SYCL in clang. With this patch device_space and host_space were added in the SYCL headers and are mapped into the new address spaces and aliases to multi_ptr instantiated with the space: device_ptr and host_ptr. Added explicit conversion operator that allows to convert device_ptr/host_ptr to global_ptr. Conversion in the opposite direction is disallowed. Also accessor to global_buffer pointer was moved to global_device address spaces. Signed-off-by: Dmitry Sidorov --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 9 ++-- .../test/CodeGenSYCL/basic-kernel-wrapper.cpp | 12 ++--- clang/test/CodeGenSYCL/kernel-metadata.cpp | 2 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 6 ++- clang/test/SemaSYCL/accessors-targets.cpp | 2 +- clang/test/SemaSYCL/basic-kernel-wrapper.cpp | 6 +-- clang/test/SemaSYCL/fake-accessors.cpp | 6 +-- clang/test/SemaSYCL/wrapped-accessor.cpp | 6 +-- sycl/include/CL/sycl/access/access.hpp | 33 +++++++++++- sycl/include/CL/sycl/atomic.hpp | 14 +++-- .../CL/sycl/detail/generic_type_lists.hpp | 29 +++++------ sycl/include/CL/sycl/handler.hpp | 4 +- sycl/include/CL/sycl/multi_ptr.hpp | 51 +++++++++++++------ sycl/include/CL/sycl/pointers.hpp | 6 +++ .../check_device_code/kernel_arguments_as.cpp | 2 +- sycl/test/check_device_code/usm_pointers.cpp | 41 +++++++++++++++ sycl/test/multi_ptr/multi_ptr.cpp | 17 +++++++ 17 files changed, 185 insertions(+), 61 deletions(-) create mode 100644 sycl/test/check_device_code/usm_pointers.cpp diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 3184c58edcbf..f1ad520ee248 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -57,7 +57,9 @@ enum class address_space : int { private_space = 0, global_space, constant_space, - local_space + local_space, + device_space, + host_space }; } // namespace access @@ -139,8 +141,9 @@ class accessor { _ImplT impl; private: - void __init(__attribute__((opencl_global)) dataT *Ptr, range AccessRange, - range MemRange, id Offset) {} + void __init(__attribute__((opencl_global_device)) dataT *Ptr, + range AccessRange, range MemRange, + id Offset) {} }; template diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index 31795fc73b77..ceab1dee6860 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -20,12 +20,12 @@ int main() { } // CHECK: define spir_kernel void @{{.*}}kernel_function -// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]], +// CHECK-SAME: i32 addrspace(11)* [[MEM_ARG:%[a-zA-Z0-9_]+]], // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]], // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]], // CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]]) // Check alloca for pointer argument -// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)* +// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(11)* // Check lambda object alloca // CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon" // Check allocas for ranges @@ -34,7 +34,7 @@ int main() { // CHECK: [[OID:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::id" // // Check store of kernel pointer argument to alloca -// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8 +// CHECK: store i32 addrspace(11)* [[MEM_ARG]], i32 addrspace(11)** [[MEM_ARG]].addr, align 8 // Check for default constructor of accessor // CHECK: call spir_func {{.*}}accessor @@ -43,12 +43,12 @@ int main() { // CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0 // Check load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr +// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(11)*, i32 addrspace(11)** [[MEM_ARG]].addr // Check accessor __init method call -// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) +// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) // CHECK: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) // Check lambda "()" operator call // CHECK-OLD: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]]) diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp index dd502fa1dd84..d7f1f7cfe59f 100644 --- a/clang/test/CodeGenSYCL/kernel-metadata.cpp +++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_addr_space ![[MDAS:[0-9]+]] !kernel_arg_access_qual ![[MDAC:[0-9]+]] !kernel_arg_type ![[MDAT:[0-9]+]] !kernel_arg_base_type ![[MDAT:[0-9]+]] !kernel_arg_type_qual ![[MDATQ:[0-9]+]] -// CHECK: ![[MDAS]] = !{i32 1, i32 0, i32 0, i32 0} +// CHECK: ![[MDAS]] = !{i32 11, i32 0, i32 0, i32 0} // CHECK: ![[MDAC]] = !{!"none", !"none", !"none", !"none"} // CHECK: ![[MDAT]] = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"} // CHECK: ![[MDATQ]] = !{!"", !"", !"", !""} diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 5bd37447ce81..670c57c65c89 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -33,7 +33,9 @@ enum class address_space : int { private_space = 0, global_space, constant_space, - local_space + local_space, + device_space, + host_space }; } // namespace access @@ -57,7 +59,7 @@ struct DeviceValueType; template struct DeviceValueType { - using type = __attribute__((opencl_global)) dataT; + using type = __attribute__((opencl_global_device)) dataT; }; template diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp index dbaab2664e95..b03f4dbaa11b 100644 --- a/clang/test/SemaSYCL/accessors-targets.cpp +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -37,5 +37,5 @@ int main() { }); } // CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' -// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_global{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index 1f500eff0a88..07420f7f296f 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -23,11 +23,11 @@ int main() { // Check declaration of the kernel -// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel -// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *' +// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global_device int *' // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>' @@ -47,7 +47,7 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var // CHECK-NEXT: ImplicitCastExpr {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global_device int *' // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp index 24d36a6ba54b..1c911ae2f123 100644 --- a/clang/test/SemaSYCL/fake-accessors.cpp +++ b/clang/test/SemaSYCL/fake-accessors.cpp @@ -51,6 +51,6 @@ int main() { }); return 0; } -// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: fake_accessors{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 83bb3ff2448f..18fe7e70d959 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -23,11 +23,11 @@ int main() { } // Check declaration of the kernel -// CHECK: wrapped_access{{.*}} 'void (AccWrapper>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: wrapped_access{{.*}} 'void (AccWrapper>, __global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper>':'AccWrapper>' -// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *' +// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global_device int *' // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>' @@ -49,7 +49,7 @@ int main() { // Parameters of the _init method // CHECK-NEXT: ImplicitCastExpr {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global_device int *' // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 10101d02435f..7e8348d274fe 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -45,7 +45,9 @@ enum class address_space : int { private_space = 0, global_space, constant_space, - local_space + local_space, + device_space = 11, + host_space = 12 }; } // namespace access @@ -103,11 +105,15 @@ constexpr bool modeWritesNewData(access::mode m) { #ifdef __SYCL_DEVICE_ONLY__ #define __OPENCL_GLOBAL_AS__ __attribute__((opencl_global)) +#define __OPENCL_GLOBAL_DEVICE_AS__ __attribute__((opencl_global_device)) +#define __OPENCL_GLOBAL_HOST_AS__ __attribute__((opencl_global_host)) #define __OPENCL_LOCAL_AS__ __attribute__((opencl_local)) #define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant)) #define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private)) #else #define __OPENCL_GLOBAL_AS__ +#define __OPENCL_GLOBAL_DEVICE_AS__ +#define __OPENCL_GLOBAL_HOST_AS__ #define __OPENCL_LOCAL_AS__ #define __OPENCL_CONSTANT_AS__ #define __OPENCL_PRIVATE_AS__ @@ -118,6 +124,11 @@ template struct TargetToAS { access::address_space::global_space; }; +template <> struct TargetToAS { + constexpr static access::address_space AS = + access::address_space::device_space; +}; + template <> struct TargetToAS { constexpr static access::address_space AS = access::address_space::local_space; @@ -141,6 +152,16 @@ struct PtrValueType { using type = __OPENCL_GLOBAL_AS__ ElementType; }; +template +struct PtrValueType { + using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType; +}; + +template +struct PtrValueType { + using type = __OPENCL_GLOBAL_HOST_AS__ ElementType; +}; + template struct PtrValueType { // Current implementation of address spaces handling leads to possibility @@ -171,6 +192,14 @@ struct remove_AS<__OPENCL_GLOBAL_AS__ T> { typedef T type; }; +template struct remove_AS<__OPENCL_GLOBAL_DEVICE_AS__ T> { + typedef T type; +}; + +template struct remove_AS<__OPENCL_GLOBAL_HOST_AS__ T> { + typedef T type; +}; + template struct remove_AS<__OPENCL_PRIVATE_AS__ T> { typedef T type; @@ -188,6 +217,8 @@ struct remove_AS<__OPENCL_CONSTANT_AS__ T> { #endif #undef __OPENCL_GLOBAL_AS__ +#undef __OPENCL_GLOBAL_DEVICE_AS__ +#undef __OPENCL_GLOBAL_HOST_AS__ #undef __OPENCL_LOCAL_AS__ #undef __OPENCL_CONSTANT_AS__ #undef __OPENCL_PRIVATE_AS__ diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp index da9daa465efd..5244ebb637dd 100644 --- a/sycl/include/CL/sycl/atomic.hpp +++ b/sycl/include/CL/sycl/atomic.hpp @@ -47,7 +47,8 @@ template struct IsValidAtomicType { template struct IsValidAtomicAddressSpace { static constexpr bool value = (AS == access::address_space::global_space || - AS == access::address_space::local_space); + AS == access::address_space::local_space || + AS == access::address_space::device_space); }; // Type trait to translate a cl::sycl::access::address_space to @@ -56,6 +57,9 @@ template struct GetSpirvMemoryScope {}; template <> struct GetSpirvMemoryScope { static constexpr auto scope = __spv::Scope::Device; }; +template <> struct GetSpirvMemoryScope { + static constexpr auto scope = __spv::Scope::Device; +}; template <> struct GetSpirvMemoryScope { static constexpr auto scope = __spv::Scope::Workgroup; }; @@ -168,12 +172,12 @@ template class atomic { static_assert(detail::IsValidAtomicType::value, - "Invalid SYCL atomic type. Valid types are: int, " - "unsigned int, long, unsigned long, long long, unsigned " + "Invalid SYCL atomic type. Valid types are: int, " + "unsigned int, long, unsigned long, long long, unsigned " "long long, float"); static_assert(detail::IsValidAtomicAddressSpace::value, - "Invalid SYCL atomic address_space. Valid address spaces are: " - "global_space, local_space"); + "Invalid SYCL atomic address_space. Valid address spaces are: " + "global_space, local_space, device_space"); static constexpr auto SpirvScope = detail::GetSpirvMemoryScope::scope; diff --git a/sycl/include/CL/sycl/detail/generic_type_lists.hpp b/sycl/include/CL/sycl/detail/generic_type_lists.hpp index 191b52765c52..64e28b469566 100644 --- a/sycl/include/CL/sycl/detail/generic_type_lists.hpp +++ b/sycl/include/CL/sycl/detail/generic_type_lists.hpp @@ -361,21 +361,20 @@ using nan_list = type_list; - -using nonconst_address_space_list = - address_space_list; - -using nonlocal_address_space_list = - address_space_list; +using all_address_space_list = address_space_list< + access::address_space::local_space, access::address_space::global_space, + access::address_space::private_space, access::address_space::constant_space, + access::address_space::device_space, access::address_space::host_space>; + +using nonconst_address_space_list = address_space_list< + access::address_space::local_space, access::address_space::global_space, + access::address_space::private_space, access::address_space::device_space, + access::address_space::host_space>; + +using nonlocal_address_space_list = address_space_list< + access::address_space::global_space, access::address_space::private_space, + access::address_space::constant_space, access::address_space::device_space, + access::address_space::host_space>; } // namespace gvl } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 5ff72f711f0d..c88ca284ab50 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -505,7 +505,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(); } @@ -528,7 +528,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/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index 4495c654ecb3..48577933dce1 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -108,17 +108,18 @@ template class multi_ptr { return reinterpret_cast(m_Pointer)[index]; } - // Only if Space == global_space + // Only if Space == global_space || device_space template ::type> + (Space == access::address_space::global_space || + Space == access::address_space::device_space)>::type> multi_ptr(accessor Accessor) { - m_Pointer = (pointer_t)(Accessor.get_pointer().m_Pointer); + m_Pointer = (pointer_t)(Accessor.get_pointer().get()); } // Only if Space == local_space @@ -152,14 +153,16 @@ template class multi_ptr { // 2. from multi_ptr to multi_ptr - // Only if Space == global_space and element type is const - template < - int dimensions, access::mode Mode, access::placeholder isPlaceholder, - access::address_space _Space = Space, typename ET = ElementType, - typename = typename std::enable_if< - _Space == Space && Space == access::address_space::global_space && - std::is_const::value && - std::is_same::value>::type> + // Only if Space == global_space || device_space and element type is const + template ::value && + std::is_same::value>::type> multi_ptr(accessor::type, dimensions, Mode, access::target::global_buffer, isPlaceholder> Accessor) @@ -271,6 +274,22 @@ template class multi_ptr { return multi_ptr(m_Pointer - r); } + // Explicit conversion to global_space + // Only available if Space == address_space::device_space || + // Space == address_space::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( + (global_pointer_t)m_Pointer); + } + // Only if Space == global_space template class multi_ptr { return *this; } - // Only if Space == global_space + // Only if Space == global_space || device_space template ::type> + (Space == access::address_space::global_space || + Space == access::address_space::device_space)>::type> multi_ptr( accessor @@ -466,12 +486,13 @@ class multi_ptr { return *this; } - // Only if Space == global_space + // Only if Space == global_space || device_space template ::type> + (Space == access::address_space::global_space || + Space == access::address_space::device_space)>::type> multi_ptr( accessor diff --git a/sycl/include/CL/sycl/pointers.hpp b/sycl/include/CL/sycl/pointers.hpp index 9f91ba70ee6b..3a4fb3beda66 100644 --- a/sycl/include/CL/sycl/pointers.hpp +++ b/sycl/include/CL/sycl/pointers.hpp @@ -19,6 +19,12 @@ template class multi_ptr; template using global_ptr = multi_ptr; +template +using device_ptr = multi_ptr; + +template +using host_ptr = multi_ptr; + template using local_ptr = multi_ptr; diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp index 0c4c4a1dd2b4..42d7f6cf4649 100644 --- a/sycl/test/check_device_code/kernel_arguments_as.cpp +++ b/sycl/test/check_device_code/kernel_arguments_as.cpp @@ -4,7 +4,7 @@ // Check the address space of the pointer in accessor class. // // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" } -// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(1)* } +// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(11)* } // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" } // CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* } // diff --git a/sycl/test/check_device_code/usm_pointers.cpp b/sycl/test/check_device_code/usm_pointers.cpp new file mode 100644 index 000000000000..0bf07764d242 --- /dev/null +++ b/sycl/test/check_device_code/usm_pointers.cpp @@ -0,0 +1,41 @@ +// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning +// RUN: FileCheck %s --input-file %t.ll +// +// Check the address space of the pointer in multi_ptr class +// +// CHECK: %[[DEVPTR_T:.*]] = type { i8 addrspace(11)* } +// CHECK: %[[HOSTPTR_T:.*]] = type { i8 addrspace(12)* } +// +// CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}} +// CHECK: %m_Pointer = getelementptr inbounds %[[DEVPTR_T]] +// CHECK-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(11)*, i8 addrspace(11)* addrspace(4)* %m_Pointer +// CHECK-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(11)* %[[DEVLOAD]] to i8 addrspace(4)* +// ret i8 addrspace(4)* %[[DEVCAST]] +// +// CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}} +// CHECK: %m_Pointer = getelementptr inbounds %[[HOSTPTR_T]] +// CHECK-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(12)*, i8 addrspace(12)* addrspace(4)* %m_Pointer +// CHECK-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(12)* %[[HOSTLOAD]] to i8 addrspace(4)* +// ret i8 addrspace(4)* %[[HOSTCAST]] + +#include + +using namespace cl::sycl; + +int main() { + cl::sycl::queue queue; + { + queue.submit([&](cl::sycl::handler &cgh) { + cgh.single_task([=]() { + void *Ptr = nullptr; + device_ptr DevPtr(Ptr); + host_ptr HostPtr(Ptr); + global_ptr GlobPtr = global_ptr(DevPtr); + GlobPtr = global_ptr(HostPtr); + }); + }); + queue.wait(); + } + + return 0; +} diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp index c2e44f461e1b..04d978ca8e69 100644 --- a/sycl/test/multi_ptr/multi_ptr.cpp +++ b/sycl/test/multi_ptr/multi_ptr.cpp @@ -82,6 +82,7 @@ template void testMultPtr() { auto local_ptr = make_ptr( localAccessor.get_pointer()); + // General conversions in multi_ptr class T *RawPtr = nullptr; global_ptr ptr_4(RawPtr); ptr_4 = RawPtr; @@ -92,6 +93,12 @@ template void testMultPtr() { ptr_6 = (void *)RawPtr; + // Explicit conversions for device_ptr/host_ptr to global_ptr + device_ptr ptr_7((void *)RawPtr); + global_ptr ptr_8 = global_ptr(ptr_7); + host_ptr ptr_9((void *)RawPtr); + global_ptr ptr_10 = global_ptr(ptr_9); + innerFunc(wiID.get(0), ptr_1, ptr_2, local_ptr); }); }); @@ -109,12 +116,14 @@ void testMultPtrArrowOperator() { point data_1[1] = {1}; point data_2[1] = {2}; point data_3[1] = {3}; + point data_4[1] = {4}; { range<1> numOfItems{1}; buffer, 1> bufferData_1(data_1, numOfItems); buffer, 1> bufferData_2(data_2, numOfItems); buffer, 1> bufferData_3(data_3, numOfItems); + buffer, 1> bufferData_4(data_4, numOfItems); queue myQueue; myQueue.submit([&](handler &cgh) { accessor, 1, access::mode::read, access::target::global_buffer, @@ -126,6 +135,9 @@ void testMultPtrArrowOperator() { accessor, 1, access::mode::read_write, access::target::local, access::placeholder::false_t> accessorData_3(1, cgh); + accessor, 1, access::mode::read, access::target::global_buffer, + access::placeholder::false_t> + accessorData_4(bufferData_4, cgh); cgh.single_task>([=]() { auto ptr_1 = make_ptr, access::address_space::global_space>( @@ -134,10 +146,13 @@ void testMultPtrArrowOperator() { accessorData_2.get_pointer()); auto ptr_3 = make_ptr, access::address_space::local_space>( accessorData_3.get_pointer()); + auto ptr_4 = make_ptr, access::address_space::device_space>( + accessorData_4.get_pointer()); auto x1 = ptr_1->x; auto x2 = ptr_2->x; auto x3 = ptr_3->x; + auto x4 = ptr_4 -> x; static_assert(std::is_same::value, "Expected decltype(ptr_1->x) == T"); @@ -145,6 +160,8 @@ void testMultPtrArrowOperator() { "Expected decltype(ptr_2->x) == T"); static_assert(std::is_same::value, "Expected decltype(ptr_3->x) == T"); + static_assert(std::is_same::value, + "Expected decltype(ptr_4->x) == T"); }); }); } From 23a8adafaff1a2a28d896f76d18fd1d6f68c5450 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Fri, 19 Jun 2020 00:54:09 +0300 Subject: [PATCH 2/7] Remove assigned values Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/access/access.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 7e8348d274fe..dab04441cba0 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -46,8 +46,8 @@ enum class address_space : int { global_space, constant_space, local_space, - device_space = 11, - host_space = 12 + device_space, + host_space }; } // namespace access From 5aa55f390d6624176a99fde35d42ee3b72e4b49c Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 22 Jun 2020 18:35:19 +0300 Subject: [PATCH 3/7] Rename device/host_space to global_device/host_space Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/access/access.hpp | 10 +++---- sycl/include/CL/sycl/atomic.hpp | 12 +++++---- .../CL/sycl/detail/generic_type_lists.hpp | 27 +++++++++++-------- sycl/include/CL/sycl/handler.hpp | 4 +-- sycl/include/CL/sycl/multi_ptr.hpp | 25 ++++++++--------- sycl/include/CL/sycl/pointers.hpp | 6 +++-- sycl/test/multi_ptr/multi_ptr.cpp | 2 +- 7 files changed, 48 insertions(+), 38 deletions(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index dab04441cba0..0aaf4ebd0e89 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -46,8 +46,8 @@ enum class address_space : int { global_space, constant_space, local_space, - device_space, - host_space + global_device_space, + global_host_space }; } // namespace access @@ -126,7 +126,7 @@ template struct TargetToAS { template <> struct TargetToAS { constexpr static access::address_space AS = - access::address_space::device_space; + access::address_space::global_device_space; }; template <> struct TargetToAS { @@ -153,12 +153,12 @@ struct PtrValueType { }; template -struct PtrValueType { +struct PtrValueType { using type = __OPENCL_GLOBAL_DEVICE_AS__ ElementType; }; template -struct PtrValueType { +struct PtrValueType { using type = __OPENCL_GLOBAL_HOST_AS__ ElementType; }; diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp index 5244ebb637dd..6c0be13b2c52 100644 --- a/sycl/include/CL/sycl/atomic.hpp +++ b/sycl/include/CL/sycl/atomic.hpp @@ -46,9 +46,10 @@ template struct IsValidAtomicType { }; template struct IsValidAtomicAddressSpace { - static constexpr bool value = (AS == access::address_space::global_space || - AS == access::address_space::local_space || - AS == access::address_space::device_space); + static constexpr bool value = + (AS == access::address_space::global_space || + AS == access::address_space::local_space || + AS == access::address_space::global_device_space); }; // Type trait to translate a cl::sycl::access::address_space to @@ -57,7 +58,8 @@ template struct GetSpirvMemoryScope {}; template <> struct GetSpirvMemoryScope { static constexpr auto scope = __spv::Scope::Device; }; -template <> struct GetSpirvMemoryScope { +template <> +struct GetSpirvMemoryScope { static constexpr auto scope = __spv::Scope::Device; }; template <> struct GetSpirvMemoryScope { @@ -177,7 +179,7 @@ class atomic { "long long, float"); static_assert(detail::IsValidAtomicAddressSpace::value, "Invalid SYCL atomic address_space. Valid address spaces are: " - "global_space, local_space, device_space"); + "global_space, local_space, global_device_space"); static constexpr auto SpirvScope = detail::GetSpirvMemoryScope::scope; diff --git a/sycl/include/CL/sycl/detail/generic_type_lists.hpp b/sycl/include/CL/sycl/detail/generic_type_lists.hpp index 64e28b469566..9965ea66eee9 100644 --- a/sycl/include/CL/sycl/detail/generic_type_lists.hpp +++ b/sycl/include/CL/sycl/detail/generic_type_lists.hpp @@ -364,17 +364,22 @@ namespace gvl { using all_address_space_list = address_space_list< access::address_space::local_space, access::address_space::global_space, access::address_space::private_space, access::address_space::constant_space, - access::address_space::device_space, access::address_space::host_space>; - -using nonconst_address_space_list = address_space_list< - access::address_space::local_space, access::address_space::global_space, - access::address_space::private_space, access::address_space::device_space, - access::address_space::host_space>; - -using nonlocal_address_space_list = address_space_list< - access::address_space::global_space, access::address_space::private_space, - access::address_space::constant_space, access::address_space::device_space, - access::address_space::host_space>; + access::address_space::global_device_space, + access::address_space::global_host_space>; + +using nonconst_address_space_list = + address_space_list; + +using nonlocal_address_space_list = + address_space_list; } // namespace gvl } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index c88ca284ab50..b06cbc412a0b 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -505,7 +505,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(); } @@ -528,7 +528,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/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index 48577933dce1..764efac6b0e4 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -108,14 +108,14 @@ template class multi_ptr { return reinterpret_cast(m_Pointer)[index]; } - // Only if Space == global_space || device_space + // Only if Space == global_space || global_device_space template ::type> + Space == access::address_space::global_device_space)>::type> multi_ptr(accessor Accessor) { @@ -153,14 +153,15 @@ template class multi_ptr { // 2. from multi_ptr to multi_ptr - // Only if Space == global_space || device_space and element type is const + // Only if Space == global_space || global_device_space and element type is + // const template ::value && std::is_same::value>::type> multi_ptr(accessor::type, dimensions, Mode, @@ -275,13 +276,13 @@ template class multi_ptr { } // Explicit conversion to global_space - // Only available if Space == address_space::device_space || - // Space == address_space::host_space + // Only available if Space == address_space::global_device_space || + // Space == address_space::global_host_space template ::type> + (Space == access::address_space::global_device_space || + Space == access::address_space::global_host_space)>::type> explicit operator multi_ptr() const { using global_pointer_t = typename detail::PtrValueType< @@ -364,13 +365,13 @@ template class multi_ptr { return *this; } - // Only if Space == global_space || device_space + // Only if Space == global_space || global_device_space template ::type> + Space == access::address_space::global_device_space)>::type> multi_ptr( accessor @@ -486,13 +487,13 @@ class multi_ptr { return *this; } - // Only if Space == global_space || device_space + // Only if Space == global_space || global_device_space template ::type> + Space == access::address_space::global_device_space)>::type> multi_ptr( accessor diff --git a/sycl/include/CL/sycl/pointers.hpp b/sycl/include/CL/sycl/pointers.hpp index 3a4fb3beda66..efec74e0fd3a 100644 --- a/sycl/include/CL/sycl/pointers.hpp +++ b/sycl/include/CL/sycl/pointers.hpp @@ -20,10 +20,12 @@ template using global_ptr = multi_ptr; template -using device_ptr = multi_ptr; +using device_ptr = + multi_ptr; template -using host_ptr = multi_ptr; +using host_ptr = + multi_ptr; template using local_ptr = multi_ptr; diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp index 04d978ca8e69..bd394dbd559d 100644 --- a/sycl/test/multi_ptr/multi_ptr.cpp +++ b/sycl/test/multi_ptr/multi_ptr.cpp @@ -146,7 +146,7 @@ void testMultPtrArrowOperator() { accessorData_2.get_pointer()); auto ptr_3 = make_ptr, access::address_space::local_space>( accessorData_3.get_pointer()); - auto ptr_4 = make_ptr, access::address_space::device_space>( + auto ptr_4 = make_ptr, access::address_space::global_device_space>( accessorData_4.get_pointer()); auto x1 = ptr_1->x; From 2c832ee9466b3ab990c670f3619ff2542198c4f6 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 24 Jun 2020 20:50:51 +0300 Subject: [PATCH 4/7] Remove clang testing Signed-off-by: Dmitry Sidorov --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 9 +++------ clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp | 12 ++++++------ clang/test/CodeGenSYCL/kernel-metadata.cpp | 2 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 6 ++---- clang/test/SemaSYCL/accessors-targets.cpp | 2 +- clang/test/SemaSYCL/basic-kernel-wrapper.cpp | 6 +++--- clang/test/SemaSYCL/fake-accessors.cpp | 6 +++--- clang/test/SemaSYCL/wrapped-accessor.cpp | 6 +++--- 8 files changed, 22 insertions(+), 27 deletions(-) diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index f1ad520ee248..3184c58edcbf 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -57,9 +57,7 @@ enum class address_space : int { private_space = 0, global_space, constant_space, - local_space, - device_space, - host_space + local_space }; } // namespace access @@ -141,9 +139,8 @@ class accessor { _ImplT impl; private: - void __init(__attribute__((opencl_global_device)) dataT *Ptr, - range AccessRange, range MemRange, - id Offset) {} + void __init(__attribute__((opencl_global)) dataT *Ptr, range AccessRange, + range MemRange, id Offset) {} }; template diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index ceab1dee6860..31795fc73b77 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -20,12 +20,12 @@ int main() { } // CHECK: define spir_kernel void @{{.*}}kernel_function -// CHECK-SAME: i32 addrspace(11)* [[MEM_ARG:%[a-zA-Z0-9_]+]], +// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]], // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]], // CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]], // CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]]) // Check alloca for pointer argument -// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(11)* +// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)* // Check lambda object alloca // CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon" // Check allocas for ranges @@ -34,7 +34,7 @@ int main() { // CHECK: [[OID:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::id" // // Check store of kernel pointer argument to alloca -// CHECK: store i32 addrspace(11)* [[MEM_ARG]], i32 addrspace(11)** [[MEM_ARG]].addr, align 8 +// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8 // Check for default constructor of accessor // CHECK: call spir_func {{.*}}accessor @@ -43,12 +43,12 @@ int main() { // CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0 // Check load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(11)*, i32 addrspace(11)** [[MEM_ARG]].addr +// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr // Check accessor __init method call -// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) +// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) // CHECK: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)* -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(11)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) // Check lambda "()" operator call // CHECK-OLD: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]]) diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp index d7f1f7cfe59f..dd502fa1dd84 100644 --- a/clang/test/CodeGenSYCL/kernel-metadata.cpp +++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_addr_space ![[MDAS:[0-9]+]] !kernel_arg_access_qual ![[MDAC:[0-9]+]] !kernel_arg_type ![[MDAT:[0-9]+]] !kernel_arg_base_type ![[MDAT:[0-9]+]] !kernel_arg_type_qual ![[MDATQ:[0-9]+]] -// CHECK: ![[MDAS]] = !{i32 11, i32 0, i32 0, i32 0} +// CHECK: ![[MDAS]] = !{i32 1, i32 0, i32 0, i32 0} // CHECK: ![[MDAC]] = !{!"none", !"none", !"none", !"none"} // CHECK: ![[MDAT]] = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"} // CHECK: ![[MDATQ]] = !{!"", !"", !"", !""} diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 670c57c65c89..5bd37447ce81 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -33,9 +33,7 @@ enum class address_space : int { private_space = 0, global_space, constant_space, - local_space, - device_space, - host_space + local_space }; } // namespace access @@ -59,7 +57,7 @@ struct DeviceValueType; template struct DeviceValueType { - using type = __attribute__((opencl_global_device)) dataT; + using type = __attribute__((opencl_global)) dataT; }; template diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp index b03f4dbaa11b..dbaab2664e95 100644 --- a/clang/test/SemaSYCL/accessors-targets.cpp +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -37,5 +37,5 @@ int main() { }); } // CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' -// CHECK: {{.*}}use_global{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index 07420f7f296f..1f500eff0a88 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -23,11 +23,11 @@ int main() { // Check declaration of the kernel -// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel -// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global_device int *' +// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *' // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>' @@ -47,7 +47,7 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var // CHECK-NEXT: ImplicitCastExpr {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global_device int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *' // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp index 1c911ae2f123..24d36a6ba54b 100644 --- a/clang/test/SemaSYCL/fake-accessors.cpp +++ b/clang/test/SemaSYCL/fake-accessors.cpp @@ -51,6 +51,6 @@ int main() { }); return 0; } -// CHECK: fake_accessors{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_typedef{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_alias{{.*}} 'void (__global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 18fe7e70d959..83bb3ff2448f 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -23,11 +23,11 @@ int main() { } // Check declaration of the kernel -// CHECK: wrapped_access{{.*}} 'void (AccWrapper>, __global_device int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: wrapped_access{{.*}} 'void (AccWrapper>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper>':'AccWrapper>' -// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global_device int *' +// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *' // CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' // CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>' @@ -49,7 +49,7 @@ int main() { // Parameters of the _init method // CHECK-NEXT: ImplicitCastExpr {{.*}} -// CHECK-NEXT: DeclRefExpr {{.*}} '__global_device int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global_device int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global int *' // CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue From 2ef86a0c0c86fed2e887724798fb39defca9dffd Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Thu, 25 Jun 2020 17:42:30 +0300 Subject: [PATCH 5/7] Update ASi values Signed-off-by: Dmitry Sidorov --- sycl/test/check_device_code/kernel_arguments_as.cpp | 2 +- sycl/test/check_device_code/usm_pointers.cpp | 12 ++++++------ 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp index 42d7f6cf4649..7faef5ec7244 100644 --- a/sycl/test/check_device_code/kernel_arguments_as.cpp +++ b/sycl/test/check_device_code/kernel_arguments_as.cpp @@ -4,7 +4,7 @@ // Check the address space of the pointer in accessor class. // // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" } -// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(11)* } +// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", i32 addrspace(5)* } // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" } // CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* } // diff --git a/sycl/test/check_device_code/usm_pointers.cpp b/sycl/test/check_device_code/usm_pointers.cpp index 0bf07764d242..aa0a0ed58045 100644 --- a/sycl/test/check_device_code/usm_pointers.cpp +++ b/sycl/test/check_device_code/usm_pointers.cpp @@ -3,19 +3,19 @@ // // Check the address space of the pointer in multi_ptr class // -// CHECK: %[[DEVPTR_T:.*]] = type { i8 addrspace(11)* } -// CHECK: %[[HOSTPTR_T:.*]] = type { i8 addrspace(12)* } +// CHECK: %[[DEVPTR_T:.*]] = type { i8 addrspace(5)* } +// CHECK: %[[HOSTPTR_T:.*]] = type { i8 addrspace(6)* } // // CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}} // CHECK: %m_Pointer = getelementptr inbounds %[[DEVPTR_T]] -// CHECK-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(11)*, i8 addrspace(11)* addrspace(4)* %m_Pointer -// CHECK-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(11)* %[[DEVLOAD]] to i8 addrspace(4)* +// CHECK-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(5)*, i8 addrspace(5)* addrspace(4)* %m_Pointer +// CHECK-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(5)* %[[DEVLOAD]] to i8 addrspace(4)* // ret i8 addrspace(4)* %[[DEVCAST]] // // CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}} // CHECK: %m_Pointer = getelementptr inbounds %[[HOSTPTR_T]] -// CHECK-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(12)*, i8 addrspace(12)* addrspace(4)* %m_Pointer -// CHECK-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(12)* %[[HOSTLOAD]] to i8 addrspace(4)* +// CHECK-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(6)*, i8 addrspace(6)* addrspace(4)* %m_Pointer +// CHECK-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(6)* %[[HOSTLOAD]] to i8 addrspace(4)* // ret i8 addrspace(4)* %[[HOSTCAST]] #include From fc5aa11379a3e97a6d26318a09d29f1d96bc84b9 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 1 Jul 2020 16:48:03 +0300 Subject: [PATCH 6/7] Move accessor pointers back to global space Otherwise it breaks atomics. Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/access/access.hpp | 5 ----- sycl/include/CL/sycl/handler.hpp | 4 ++-- sycl/include/CL/sycl/multi_ptr.hpp | 16 ---------------- .../check_device_code/kernel_arguments_as.cpp | 2 +- 4 files changed, 3 insertions(+), 24 deletions(-) diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 0aaf4ebd0e89..9187a972bd6f 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -124,11 +124,6 @@ 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/handler.hpp b/sycl/include/CL/sycl/handler.hpp index b06cbc412a0b..5ff72f711f0d 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -505,7 +505,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(); } @@ -528,7 +528,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/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index 764efac6b0e4..1a59113d9fc1 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -275,22 +275,6 @@ 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( - (global_pointer_t)m_Pointer); - } - // Only if Space == global_space template Date: Wed, 1 Jul 2020 16:52:31 +0300 Subject: [PATCH 7/7] Ignore clang-format concern Signed-off-by: Dmitry Sidorov --- sycl/test/multi_ptr/multi_ptr.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp index bd394dbd559d..9ebb33046a45 100644 --- a/sycl/test/multi_ptr/multi_ptr.cpp +++ b/sycl/test/multi_ptr/multi_ptr.cpp @@ -152,7 +152,7 @@ void testMultPtrArrowOperator() { auto x1 = ptr_1->x; auto x2 = ptr_2->x; auto x3 = ptr_3->x; - auto x4 = ptr_4 -> x; + auto x4 = ptr_4->x; static_assert(std::is_same::value, "Expected decltype(ptr_1->x) == T");