[SYCL] Fix handler set_arg(s) method.

Fixed incorrect storing and extractiong of global/constant accessor
in set_arg(s) logic.

Signed-off-by: Alexey Voronov <[email protected]>

Author: alexeyvoronov-intel

sycl/include/CL/sycl/handler2.hpp

@@ -217,6 +217,18 @@ class handler {
       void *Ptr = LambdaPtr + KernelArgs[I].offset;
       const detail::kernel_param_kind_t &Kind = KernelArgs[I].kind;
       const int &Size = KernelArgs[I].info;
+      if (Kind == detail::kernel_param_kind_t::kind_accessor) {
+        // For args kind of accessor Size is information about accessor.
+        // The first 11 bits of Size encodes the accessor target.
+        const access::target AccTarget =
+            static_cast<access::target>(Size & 0x7ff);
+        if (AccTarget == access::target::global_buffer ||
+            AccTarget == access::target::constant_buffer) {
+          detail::AccessorBaseHost *AccBase =
+              static_cast<detail::AccessorBaseHost *>(Ptr);
+          Ptr = detail::getSyclObjImpl(*AccBase).get();
+        }
+      }
       processArg(Ptr, Kind, Size, I, IndexShift, IsKernelCreatedFromSource);
     }
   }
@@ -264,9 +276,7 @@ class handler {
       switch (AccTarget) {
       case access::target::global_buffer:
       case access::target::constant_buffer: {
-        detail::AccessorBaseHost *AccBase =
-            static_cast<detail::AccessorBaseHost *>(Ptr);
-        detail::Requirement *AccImpl = detail::getSyclObjImpl(*AccBase).get();
+        detail::Requirement *AccImpl = static_cast<detail::Requirement *>(Ptr);
         MArgs.emplace_back(Kind, AccImpl, Size, Index + IndexShift);
         if (!IsKernelCreatedFromSource) {
           const size_t SizeAccField = sizeof(size_t) * AccImpl->MDims;
@@ -312,7 +322,6 @@ class handler {
       case access::target::image_array: {
         throw cl::sycl::invalid_parameter_error(
             "Unsupported accessor target case.");
-        assert(0);
         break;
       }
       }
@@ -445,7 +454,7 @@ class handler {
     // Store copy of the accessor.
     MAccStorage.push_back(std::move(AccImpl));
     // Add accessor to the list of arguments.
-    MArgs.emplace_back(detail::kernel_param_kind_t::kind_accessor, AccBase,
+    MArgs.emplace_back(detail::kernel_param_kind_t::kind_accessor, Req,
                        static_cast<int>(AccessTarget), ArgIndex);
   }
 

sycl/test/basic_tests/set_arg_interop.cpp

@@ -0,0 +1,81 @@
+// RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl -O3
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+#include <cassert>
+
+using namespace cl::sycl;
+
+int main() {
+  queue Queue;
+  if (!Queue.is_host()) {
+    context Context = Queue.get_context();
+
+    cl_context ClContext = Context.get();
+
+    const size_t CountSources = 2;
+    const char *Sources[CountSources] = {
+        "kernel void foo1(global float* Array, global int* Value) { *Array = "
+        "42; *Value = 1; }\n",
+        "kernel void foo2(global float* Array) { int id = get_global_id(0); "
+        "Array[id] = id; }\n",
+    };
+
+    cl_int Err;
+    cl_program ClProgram = clCreateProgramWithSource(ClContext, CountSources,
+                                                     Sources, nullptr, &Err);
+    assert(Err == CL_SUCCESS);
+
+    Err = clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr);
+    assert(Err == CL_SUCCESS);
+
+    cl_kernel FirstCLKernel = clCreateKernel(ClProgram, "foo1", &Err);
+    assert(Err == CL_SUCCESS);
+
+    cl_kernel SecondCLKernel = clCreateKernel(ClProgram, "foo2", &Err);
+    assert(Err == CL_SUCCESS);
+
+    const size_t Count = 100;
+    float Array[Count];
+
+    kernel FirstKernel(FirstCLKernel, Context);
+    kernel SecondKernel(SecondCLKernel, Context);
+    int Value;
+    {
+      buffer<float, 1> FirstBuffer(Array, range<1>(1));
+      buffer<int, 1> SecondBuffer(&Value, range<1>(1));
+      Queue.submit([&](handler &CGH) {
+        CGH.set_arg(0, FirstBuffer.get_access<access::mode::write>(CGH));
+        CGH.set_arg(1, SecondBuffer.get_access<access::mode::write>(CGH));
+        CGH.single_task(FirstKernel);
+      });
+    }
+    Queue.wait_and_throw();
+
+    assert(Array[0] == 42);
+    assert(Value == 1);
+
+    {
+      buffer<float, 1> FirstBuffer(Array, range<1>(Count));
+      Queue.submit([&](handler &CGH) {
+        auto Acc = FirstBuffer.get_access<access::mode::read_write>(CGH);
+        CGH.set_arg(0, FirstBuffer.get_access<access::mode::read_write>(CGH));
+        CGH.parallel_for(range<1>{Count}, SecondKernel);
+      });
+    }
+    Queue.wait_and_throw();
+
+    for (size_t I = 0; I < Count; ++I) {
+      assert(Array[I] == I);
+    }
+
+    clReleaseContext(ClContext);
+    clReleaseKernel(FirstCLKernel);
+    clReleaseKernel(SecondCLKernel);
+    clReleaseProgram(ClProgram);
+  }
+  return 0;
+}