Skip to content

Commit 1e3136e

Browse files
authored
[SYCL] Fix builtins address space type (#4275)
This patch fixes the `core/GroupAsyncCopy.cl` and `ocl/prefetch.cl` lit tests in `libclc`. These started failing with the introduction of the SYCL address spaces in #3634. The issue is that the builtins are declared with the SYCL address spaces in the tablegen so when used in an OpenCL setting it ends up with address space mismatches between things like `opencl_global` and `sycl_global`. These two were the only tests affected because they're also the only tested builtins that do not have generic address space variants in tablegen, the other tests would just fall back on the generic address space variants. With this patch the builtins should automatically get the appropriate version of the address space depending on the context.
1 parent 46c9f3e commit 1e3136e

File tree

5 files changed

+99
-1
lines changed

5 files changed

+99
-1
lines changed

clang/include/clang/Basic/AddressSpaces.h

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -87,6 +87,40 @@ inline bool isPtrSizeAddressSpace(LangAS AS) {
8787
AS == LangAS::ptr64);
8888
}
8989

90+
inline LangAS asSYCLLangAS(LangAS AS) {
91+
switch (AS) {
92+
case LangAS::opencl_global:
93+
return LangAS::sycl_global;
94+
case LangAS::opencl_global_device:
95+
return LangAS::sycl_global_device;
96+
case LangAS::opencl_global_host:
97+
return LangAS::sycl_global_host;
98+
case LangAS::opencl_local:
99+
return LangAS::sycl_local;
100+
case LangAS::opencl_private:
101+
return LangAS::sycl_private;
102+
default:
103+
return AS;
104+
}
105+
}
106+
107+
inline LangAS asOpenCLLangAS(LangAS AS) {
108+
switch (AS) {
109+
case LangAS::sycl_global:
110+
return LangAS::opencl_global;
111+
case LangAS::sycl_global_device:
112+
return LangAS::opencl_global_device;
113+
case LangAS::sycl_global_host:
114+
return LangAS::opencl_global_host;
115+
case LangAS::sycl_local:
116+
return LangAS::opencl_local;
117+
case LangAS::sycl_private:
118+
return LangAS::opencl_private;
119+
default:
120+
return AS;
121+
}
122+
}
123+
90124
} // namespace clang
91125

92126
#endif // LLVM_CLANG_BASIC_ADDRESSSPACES_H

clang/lib/AST/MicrosoftMangle.cpp

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2184,6 +2184,21 @@ void MicrosoftCXXNameMangler::mangleAddressSpaceType(QualType T,
21842184
case LangAS::cuda_device:
21852185
Extra.mangleSourceName("_ASCUdevice");
21862186
break;
2187+
case LangAS::sycl_global:
2188+
Extra.mangleSourceName("_ASSYglobal");
2189+
break;
2190+
case LangAS::sycl_global_device:
2191+
Extra.mangleSourceName("_ASSYdevice");
2192+
break;
2193+
case LangAS::sycl_global_host:
2194+
Extra.mangleSourceName("_ASSYhost");
2195+
break;
2196+
case LangAS::sycl_local:
2197+
Extra.mangleSourceName("_ASSYlocal");
2198+
break;
2199+
case LangAS::sycl_private:
2200+
Extra.mangleSourceName("_ASSYprivate");
2201+
break;
21872202
case LangAS::cuda_constant:
21882203
Extra.mangleSourceName("_ASCUconstant");
21892204
break;
Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,18 @@
1+
// RUN: %clang_cc1 %s -x cl -fdeclare-spirv-builtins -fsyntax-only -emit-llvm -o - -O0 | FileCheck %s
2+
//
3+
// Check that SPIR-V builtins are declared with OpenCL address spaces rather
4+
// than SYCL address spaces when using them with OpenCL. OpenCL address spaces
5+
// are mangled with the CL prefix and SYCL address spaces are mangled with the
6+
// SY prefix.
7+
8+
// CHECK: __spirv_ocl_modf{{.*}}CLglobal
9+
void modf_global(float a, global float *ptr) { __spirv_ocl_modf(a, ptr); }
10+
11+
// CHECK: __spirv_ocl_modf{{.*}}CLlocal
12+
void modf_local(float a, local float *ptr) { __spirv_ocl_modf(a, ptr); }
13+
14+
// CHECK: __spirv_ocl_modf{{.*}}CLprivate
15+
void modf_private(float a) {
16+
float *ptr;
17+
__spirv_ocl_modf(a, ptr);
18+
}
Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,29 @@
1+
// RUN: %clang_cc1 %s -fsycl-is-device -fdeclare-spirv-builtins -fsyntax-only -emit-llvm -o - -O0 | FileCheck %s
2+
//
3+
// Check that SPIR-V builtins are declared with SYCL address spaces rather
4+
// than OpenCL address spaces when using them with SYCL. OpenCL address spaces
5+
// are mangled with the CL prefix and SYCL address spaces are mangled with the
6+
// SY prefix.
7+
//
8+
// The opencl_global, opencl_local, and opencl_private attributes get turned
9+
// into sycl_global, sycl_local and sycl_private address spaces by clang.
10+
11+
#include "Inputs/sycl.hpp"
12+
13+
// CHECK: __spirv_ocl_modf{{.*}}SYglobal
14+
void modf_global(float a) {
15+
__attribute__((opencl_global)) float *ptr = nullptr;
16+
sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
17+
}
18+
19+
// CHECK: __spirv_ocl_modf{{.*}}SYlocal
20+
void modf_local(float a) {
21+
__attribute__((opencl_local)) float *ptr = nullptr;
22+
sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
23+
}
24+
25+
// CHECK: __spirv_ocl_modf{{.*}}SYprivate
26+
void modf_private(float a) {
27+
__attribute__((opencl_private)) float *ptr = nullptr;
28+
sycl::kernel_single_task<class fake_kernel>([=]() { __spirv_ocl_modf(a, ptr); });
29+
}

clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -982,7 +982,9 @@ static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name);
982982
// [const|volatile] pointers, so this is ok to do it as a last step.
983983
if (Ty.IsPointer != 0) {
984984
for (unsigned Index = 0; Index < QT.size(); Index++) {
985-
QT[Index] = Context.getAddrSpaceQualType(QT[Index], Ty.AS);
985+
QT[Index] = Context.getAddrSpaceQualType(
986+
QT[Index], S.getLangOpts().SYCLIsDevice ? asSYCLLangAS(Ty.AS)
987+
: asOpenCLLangAS(Ty.AS));
986988
QT[Index] = Context.getPointerType(QT[Index]);
987989
}
988990
}

0 commit comments

Comments
 (0)