[SYCL] Specify address space for pointer to data in accessor class

Accessor class contains pointer to data which is currently in the
generic address space. This causes SPIR-V Environment specification
violation for cases when accessor is wrapped to some object because
structure that has a pointer in the generic address space is not
allowed as a kernel argument. Fix specifies concrete address space for
this pointer in accessor class.

Signed-off-by: Artur Gainullin <[email protected]>

Author: againull

sycl/include/CL/sycl/accessor.hpp

@@ -699,7 +699,7 @@ class accessor :
 
   detail::AccessorImplDevice<AdjustedDim> impl;
 
-  PtrType MData;
+  ConcreteASPtrType MData;
 
   void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
               range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
@@ -715,7 +715,7 @@ class accessor :
       MData += Offset[0];
   }
 
-  PtrType getQualifiedPtr() const { return MData; }
+  ConcreteASPtrType getQualifiedPtr() const { return MData; }
 
 public:
   // Default constructor for objects later initialized with __init member.
@@ -1030,9 +1030,9 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
       : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
 
 private:
-  PtrType getQualifiedPtr() const { return MData; }
+  ConcreteASPtrType getQualifiedPtr() const { return MData; }
 
-  PtrType MData;
+  ConcreteASPtrType MData;
 
 #else
 

sycl/test/check_device_code/kernel_arguments_as.cpp

@@ -0,0 +1,45 @@
+// RUN: %clangxx --sycl -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll
+// RUN: FileCheck %s --input-file %t.ll
+//
+// 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: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
+// CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
+//
+// Check that kernel arguments doesn't have generic address space.
+//
+// CHECK-NOT: define weak_odr dso_local spir_kernel void @"{{.*}}check_adress_space"({{.*}}addrspace(4){{.*}})
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+template <typename Acc> struct AccWrapper { Acc accessor; };
+
+int main() {
+
+  cl::sycl::queue queue;
+  int array[10] = {0};
+  {
+    cl::sycl::buffer<int, 1> buf((int *)array, cl::sycl::range<1>(10),
+                                 {cl::sycl::property::buffer::use_host_ptr()});
+    queue.submit([&](cl::sycl::handler &cgh) {
+      auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+      cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                         cl::sycl::access::target::local>
+          local_acc(cl::sycl::range<1>(10), cgh);
+      auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
+      auto local_acc_wrapped = AccWrapper<decltype(local_acc)>{local_acc};
+      cgh.parallel_for<class check_adress_space>(
+          cl::sycl::range<1>(buf.get_count()), [=](cl::sycl::item<1> it) {
+            auto idx = it.get_linear_id();
+            acc_wrapped.accessor[idx] = local_acc_wrapped.accessor[idx];
+          });
+    });
+    queue.wait();
+  }
+
+  return 0;
+}