Skip to content

Commit 9f2fd98

Browse files
vladimirlazvladimirlaz
authored
[SYCL] Rename helper functions to avoid collision (#3665)
* [SYCL] Rename helper function to avoid collision Helper functions have the same names but different signatures to the ones to be mapped to SPIR-V operations. That makes some backends generate external calls for them which causes undefined reference on JIT compilation. The problem is seen only when compiling with -O0 or -O1 optimization options because when -O3 or -O2 are used helper functions are inlined and definitions are removed from final code eliminating the issue.
1 parent d15b841 commit 9f2fd98

File tree

3 files changed

+58
-40
lines changed

3 files changed

+58
-40
lines changed

‎sycl/include/CL/__spirv/spirv_ops.hpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -199,18 +199,18 @@ __spirv_GenericCastToPtrExplicit_ToLocal(const void *Ptr,
199199

200200
template <typename dataT>
201201
extern __attribute__((opencl_global)) dataT *
202-
__spirv_GenericCastToPtrExplicit_ToGlobal(
203-
const void *Ptr, __spv::StorageClass::Flag S) noexcept {
204-
return (__attribute__((opencl_global))
205-
dataT *)__spirv_GenericCastToPtrExplicit_ToGlobal(Ptr, S);
202+
__SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
203+
return (__attribute__((opencl_global)) dataT *)
204+
__spirv_GenericCastToPtrExplicit_ToGlobal(
205+
Ptr, __spv::StorageClass::CrossWorkgroup);
206206
}
207207

208208
template <typename dataT>
209209
extern __attribute__((opencl_local)) dataT *
210-
__spirv_GenericCastToPtrExplicit_ToLocal(const void *Ptr,
211-
__spv::StorageClass::Flag S) noexcept {
212-
return (__attribute__((opencl_local))
213-
dataT *)__spirv_GenericCastToPtrExplicit_ToLocal(Ptr, S);
210+
__SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
211+
return (__attribute__((opencl_local)) dataT *)
212+
__spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
213+
__spv::StorageClass::Workgroup);
214214
}
215215

216216
template <typename dataT>

‎sycl/include/CL/sycl/ONEAPI/sub_group.hpp

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -244,13 +244,11 @@ struct sub_group {
244244
#ifdef __NVPTX__
245245
return src[get_local_id()[0]];
246246
#else // __NVPTX__
247-
auto l = __spirv_GenericCastToPtrExplicit_ToLocal<T>(
248-
src, __spv::StorageClass::Workgroup);
247+
auto l = __SYCL_GenericCastToPtrExplicit_ToLocal<T>(src);
249248
if (l)
250249
return load(l);
251250

252-
auto g = __spirv_GenericCastToPtrExplicit_ToGlobal<T>(
253-
src, __spv::StorageClass::CrossWorkgroup);
251+
auto g = __SYCL_GenericCastToPtrExplicit_ToGlobal<T>(src);
254252
if (g)
255253
return load(g);
256254

@@ -400,15 +398,13 @@ struct sub_group {
400398
#ifdef __NVPTX__
401399
dst[get_local_id()[0]] = x;
402400
#else // __NVPTX__
403-
auto l = __spirv_GenericCastToPtrExplicit_ToLocal<T>(
404-
dst, __spv::StorageClass::Workgroup);
401+
auto l = __SYCL_GenericCastToPtrExplicit_ToLocal<T>(dst);
405402
if (l) {
406403
store(l, x);
407404
return;
408405
}
409406

410-
auto g = __spirv_GenericCastToPtrExplicit_ToGlobal<T>(
411-
dst, __spv::StorageClass::CrossWorkgroup);
407+
auto g = __SYCL_GenericCastToPtrExplicit_ToGlobal<T>(dst);
412408
if (g) {
413409
store(g, x);
414410
return;

‎sycl/test/extensions/sub_group_as.cpp

Lines changed: 46 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,7 @@
1-
// RUN: %clangxx -fsycl -fsycl-device-only -O3 -S -emit-llvm -x c++ %s -o - | FileCheck %s
2-
1+
// RUN: %clangxx -fsycl -fsycl-device-only -O3 -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefix CHECK-O3
2+
// RUN: %clangxx -fsycl -fsycl-device-only -O0 -S -emit-llvm -x c++ %s -o - | FileCheck %s --check-prefix CHECK-O0
3+
// Test compilation with -O3 when all methods are inlined in kernel function
4+
// and -O0 when helper methods are preserved.
35
#include <CL/sycl.hpp>
46
#include <cassert>
57
#include <cstdint>
@@ -43,46 +45,66 @@ int main(int argc, char *argv[]) {
4345
local[i] = i;
4446
}
4547
}
46-
// CHECK: call spir_func void {{.*}}spirv_ControlBarrierjjj
48+
// CHECK-O3: call spir_func void {{.*}}spirv_ControlBarrierjjj
4749
it.barrier();
4850

4951
int i = (it.get_global_id(0) / sg.get_max_local_range()[0]) *
5052
sg.get_max_local_range()[0];
5153

52-
// load for global address space
53-
// CHECK: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
54-
// CHECK: {{.*}}SubgroupLocalInvocationId
55-
// CHECK: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
56-
// CHECK: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
57-
// CHECK: call spir_func void {{.*}}assert
54+
// load() for global address space
55+
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
56+
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
57+
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
58+
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
59+
// CHECK-O3: call spir_func void {{.*}}assert
60+
5861
auto x = sg.load(&global[i]);
5962

6063
// load() for local address space
61-
// CHECK: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
62-
// CHECK: {{.*}}SubgroupLocalInvocationId
63-
// CHECK: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
64-
// CHECK: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
65-
// CHECK: call spir_func void {{.*}}assert
64+
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
65+
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
66+
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
67+
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
68+
// CHECK-O3: call spir_func void {{.*}}assert
6669
auto y = sg.load(&local[i]);
6770

6871
// load() for private address space
69-
// CHECK: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
70-
// CHECK: {{.*}}SubgroupLocalInvocationId
71-
// CHECK: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
72-
// CHECK: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
73-
// CHECK: call spir_func void {{.*}}assert
72+
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
73+
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
74+
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
75+
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
76+
// CHECK-O3: call spir_func void {{.*}}assert
7477
auto z = sg.load(v + i);
7578

7679
// store() for global address space
77-
// CHECK: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
78-
// CHECK: {{.*}}SubgroupLocalInvocationId
79-
// CHECK: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
80-
// CHECK: call spir_func void {{.*}}spirv_SubgroupBlockWriteINTEL{{.*}}(i32 addrspace(1)*
81-
// CHECK: call spir_func void {{.*}}assert
80+
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
81+
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
82+
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
83+
// CHECK-O3: call spir_func void {{.*}}spirv_SubgroupBlockWriteINTEL{{.*}}(i32 addrspace(1)*
84+
// CHECK-O3: call spir_func void {{.*}}assert
8285
sg.store(&global[i], x + y + z);
8386
});
8487
});
8588
}
89+
// load() accepting raw pointers method
90+
// CHECK-O0: define{{.*}}spir_func i32 {{.*}}cl4sycl6ONEAPI9sub_group4load{{.*}}addrspace(4)* %src)
91+
// CHECK-O0: call spir_func i32 addrspace(3)* {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
92+
// CHECK-O0: call spir_func i32 {{.*}}sycl6ONEAPI9sub_group4load{{.*}}i32 addrspace(3)* %
93+
// CHECK-O0: call spir_func i32 addrspace(1)* {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
94+
// CHECK-O0: call spir_func i32 {{.*}}sycl6ONEAPI9sub_group4load{{.*}}i32 addrspace(1)* %
95+
// CHECK-O0: call spir_func void {{.*}}assert
96+
97+
// store() accepting raw pointers method
98+
// CHECK-O0: define{{.*}}spir_func void {{.*}}cl4sycl6ONEAPI9sub_group5store{{.*}}i32 addrspace(4)* %dst
99+
// CHECK-O0: call spir_func i32 addrspace(3)* {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
100+
// CHECK-O0: call spir_func void {{.*}}cl4sycl6ONEAPI9sub_group5store{{.*}}, i32 addrspace(3)* %
101+
// CHECK-O0: call spir_func i32 addrspace(1)* {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
102+
// CHECK-O0: call spir_func void {{.*}}cl4sycl6ONEAPI9sub_group5store{{.*}}, i32 addrspace(1)* %
103+
// CHECK-O0: call spir_func void {{.*}}assert
104+
// CHECK-O0: define {{.*}}spir_func i32 addrspace(3)* {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)* %
105+
// CHECK-O0: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
106+
// CHECK-O0: define {{.*}}spir_func i32 addrspace(1)* {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)* %
107+
// CHECK-O0: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
86108

87109
return 0;
88110
}

0 commit comments

Comments
 (0)