[SYCL][FPGA] Change address space for USM pointers as kernel args

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 -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;
+}