[SYCL] Fix builtins address space type

clang/include/clang/Basic/AddressSpaces.h

@@ -87,6 +87,40 @@ inline bool isPtrSizeAddressSpace(LangAS AS) {
           AS == LangAS::ptr64);
 }
 
+inline LangAS asSYCLLangAS(LangAS AS) {
+  switch (AS) {
+  case LangAS::opencl_global:
+    return LangAS::sycl_global;
+  case LangAS::opencl_global_device:
+    return LangAS::sycl_global_device;
+  case LangAS::opencl_global_host:
+    return LangAS::sycl_global_host;
+  case LangAS::opencl_local:
+    return LangAS::sycl_local;
+  case LangAS::opencl_private:
+    return LangAS::sycl_private;
+  default:
+    return AS;
+  }
+}
+
+inline LangAS asOpenCLLangAS(LangAS AS) {
+  switch (AS) {
+  case LangAS::sycl_global:
+    return LangAS::opencl_global;
+  case LangAS::sycl_global_device:
+    return LangAS::opencl_global_device;
+  case LangAS::sycl_global_host:
+    return LangAS::opencl_global_host;
+  case LangAS::sycl_local:
+    return LangAS::opencl_local;
+  case LangAS::sycl_private:
+    return LangAS::opencl_private;
+  default:
+    return AS;
+  }
+}
+
 } // namespace clang
 
 #endif // LLVM_CLANG_BASIC_ADDRESSSPACES_H

clang/lib/AST/MicrosoftMangle.cpp

@@ -2184,6 +2184,21 @@ void MicrosoftCXXNameMangler::mangleAddressSpaceType(QualType T,
     case LangAS::cuda_device:
       Extra.mangleSourceName("_ASCUdevice");
       break;
+    case LangAS::sycl_global:
+      Extra.mangleSourceName("_ASSYglobal");
+      break;
+    case LangAS::sycl_global_device:
+      Extra.mangleSourceName("_ASSYdevice");
+      break;
+    case LangAS::sycl_global_host:
+      Extra.mangleSourceName("_ASSYhost");
+      break;
+    case LangAS::sycl_local:
+      Extra.mangleSourceName("_ASSYlocal");
+      break;
+    case LangAS::sycl_private:
+      Extra.mangleSourceName("_ASSYprivate");
+      break;
     case LangAS::cuda_constant:
       Extra.mangleSourceName("_ASCUconstant");
       break;

clang/test/CodeGenOpenCL/spirv-builtins-addr-space.cl

@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 %s -x cl -fdeclare-spirv-builtins -fsyntax-only -emit-llvm -o - -O0 | FileCheck %s
+//
+// Check that SPIR-V builtins are declared with OpenCL address spaces rather
+// than SYCL address spaces when using them with OpenCL. OpenCL address spaces
+// are mangled with the CL prefix and SYCL address spaces are mangled with the
+// SY prefix.
+
+// CHECK: __spirv_ocl_modf{{.*}}CLglobal
+void modf_global(float a, global float *ptr) { __spirv_ocl_modf(a, ptr); }
+
+// CHECK: __spirv_ocl_modf{{.*}}CLlocal
+void modf_local(float a, local float *ptr) { __spirv_ocl_modf(a, ptr); }
+
+// CHECK: __spirv_ocl_modf{{.*}}CLprivate
+void modf_private(float a) {
+  float *ptr;
+  __spirv_ocl_modf(a, ptr);
+}

clang/test/CodeGenSYCL/spirv-builtins-addr-space.cpp

@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 %s -fsycl-is-device -fdeclare-spirv-builtins -fsyntax-only -emit-llvm -o - -O0 | FileCheck %s
+//
+// Check that SPIR-V builtins are declared with SYCL address spaces rather
+// than OpenCL address spaces when using them with SYCL. OpenCL address spaces
+// are mangled with the CL prefix and SYCL address spaces are mangled with the
+// SY prefix.
+//
+// The opencl_global, opencl_local, and opencl_private attributes get turned
+// into sycl_global, sycl_local and sycl_private address spaces by clang.
+
+#include "Inputs/sycl.hpp"
+
+// CHECK: __spirv_ocl_modf{{.*}}SYglobal
+void modf_global(float a) {
+  __attribute__((opencl_global)) float *ptr = nullptr;
+  sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
+}
+
+// CHECK: __spirv_ocl_modf{{.*}}SYlocal
+void modf_local(float a) {
+  __attribute__((opencl_local)) float *ptr = nullptr;
+  sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
+}
+
+// CHECK: __spirv_ocl_modf{{.*}}SYprivate
+void modf_private(float a) {
+  __attribute__((opencl_private)) float *ptr = nullptr;
+  sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
+}

clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp

@@ -958,7 +958,9 @@ static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name);
   // [const|volatile] pointers, so this is ok to do it as a last step.
   if (Ty.IsPointer != 0) {
     for (unsigned Index = 0; Index < QT.size(); Index++) {
-      QT[Index] = Context.getAddrSpaceQualType(QT[Index], Ty.AS);
+      QT[Index] = Context.getAddrSpaceQualType(
+          QT[Index], S.getLangOpts().SYCLIsDevice ? asSYCLLangAS(Ty.AS)
+                                                  : asOpenCLLangAS(Ty.AS));
       QT[Index] = Context.getPointerType(QT[Index]);
     }
   }