[SYCL][FPGA] Change address space for USM pointers as kernel args (#2095)

Query for address space of USM pointer before adding the appropriate
pointer argument of OpenCL kernel. If this address space is
global_device or global_host - leave it as is.

Signed-off-by: Dmitry Sidorov <[email protected]>

Author: MrSidims

clang/lib/Sema/SemaSYCL.cpp

@@ -1244,7 +1244,11 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     // space, because OpenCL requires it.
     QualType PointeeTy = FieldTy->getPointeeType();
     Qualifiers Quals = PointeeTy.getQualifiers();
-    Quals.setAddressSpace(LangAS::opencl_global);
+    auto AS = Quals.getAddressSpace();
+    // Leave global_device and global_host address spaces as is to help FPGA
+    // device in memory allocations
+    if (AS != LangAS::opencl_global_device && AS != LangAS::opencl_global_host)
+      Quals.setAddressSpace(LangAS::opencl_global);
     PointeeTy = SemaRef.getASTContext().getQualifiedType(
         PointeeTy.getUnqualifiedType(), Quals);
     QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy);

clang/test/CodeGenSYCL/kernel-device-space-arg.cpp

@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -disable-llvm-passes -o - | FileCheck %s
+
+// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function(i32 addrspace(5)* {{.*}} i32 addrspace(6)* {{.*}}
+
+#include "sycl.hpp"
+
+int main() {
+  __attribute__((opencl_global_device)) int *GLOBDEV = nullptr;
+  __attribute__((opencl_global_host)) int *GLOBHOST = nullptr;
+  cl::sycl::kernel_single_task<class kernel_function>(
+      [=]() {
+        __attribute__((opencl_global_device)) int *DevPtr = GLOBDEV;
+        __attribute__((opencl_global_host)) int *HostPtr = GLOBHOST;
+      });
+  return 0;
+}