Skip to content

Commit 883a2f0

Browse files
baderbader
authored
[SYCL] Adjust initialization of array kernel parameters. (#19778)
If a kernel parameter is an array of pointers, current initialization method causes the address space information for kernel parameters to be lost in LLVM IR. Using implicit cast with address space conversion instead of "taking address of a pointer -> bitcasting to SYCL kernel field type -> dereferencing casted pointer" enables LLVM IR to propagate the address space of array of pointers. Some pseudo-code expressing the idea: ```c++ __global T* arr[N] // parameter type ``` Old initialization method: ```c++ {*(reinterpret_cast<T*>)(&arr[0]), *(reinterpret_cast<T*>)(&arr[1]), ...} ``` Type conversion: __global T* __generic* -> __generic T* __generic* New initialization method: ```c++ {arr[0], arr[1], ...} ``` Type conversion: __global T* -> __generic T* Since pointers in LLVM are typeless, the old initalization method doesn't covert address spaces.
1 parent ffc2512 commit 883a2f0

File tree

4 files changed

+97
-52
lines changed

4 files changed

+97
-52
lines changed

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 7 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -4097,10 +4097,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
40974097
}
40984098

40994099
void addArrayElementInit(FieldDecl *FD, QualType T) {
4100-
Expr *RCE = createReinterpretCastExpr(
4101-
createGetAddressOf(ArrayParamBases.pop_back_val()),
4102-
SemaSYCLRef.getASTContext().getPointerType(T));
4103-
Expr *Initializer = createDerefOp(RCE);
4100+
Expr *Initializer = ArrayParamBases.pop_back_val();
4101+
if (!T->isPointerType()) {
4102+
Expr *RCE = createReinterpretCastExpr(
4103+
createGetAddressOf(Initializer),
4104+
SemaSYCLRef.getASTContext().getPointerType(T));
4105+
Initializer = createDerefOp(RCE);
4106+
}
41044107
addFieldInit(FD, T, Initializer);
41054108
}
41064109

‎clang/test/CodeGenSYCL/array-kernel-param.cpp‎

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
2+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
3+
4+
#include "sycl.hpp"
5+
6+
sycl::queue myQueue;
7+
8+
using namespace sycl;
9+
10+
// CHECK: %struct.__wrapper_class = type { [2 x i32] }
11+
// CHECK: %class.anon = type { [2 x i32] }
12+
// CHECK: %struct.__wrapper_class.0 = type { [2 x ptr addrspace(1)] }
13+
// CHECK: %class.anon.1 = type { [2 x ptr addrspace(4)] }
14+
15+
int main() {
16+
int Array[2];
17+
myQueue.submit([&](sycl::handler &h) {
18+
h.single_task<class IntArray>(
19+
[=] {
20+
int local = Array[1];
21+
});
22+
});
23+
24+
// CHECK-LABEL: @{{.*}}IntArray(ptr {{.*}}byval(%struct.__wrapper_class)
25+
// CHECK: %__SYCLKernel = alloca %class.anon, align 4
26+
// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
27+
// CHECK: %_arg_Array.ascast = addrspacecast ptr %_arg_Array to ptr addrspace(4)
28+
// CHECK: %Array = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
29+
// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class, ptr addrspace(4) %_arg_Array.ascast, i32 0, i32 0
30+
// CHECK: %arrayinit.begin = getelementptr inbounds [2 x i32], ptr addrspace(4) %Array, i64 0, i64 0
31+
// CHECK: br label %arrayinit.body
32+
// CHECK: arrayinit.body: ; preds = %arrayinit.body, %entry
33+
// CHECK: %arrayinit.index = phi i64 [ 0, %entry ], [ %arrayinit.next, %arrayinit.body ]
34+
// CHECK: %1 = getelementptr inbounds i32, ptr addrspace(4) %arrayinit.begin, i64 %arrayinit.index
35+
// CHECK: %arrayidx = getelementptr inbounds nuw [2 x i32], ptr addrspace(4) %0, i64 0, i64 %arrayinit.index
36+
// CHECK: %2 = load i32, ptr addrspace(4) %arrayidx, align 4
37+
// CHECK: store i32 %2, ptr addrspace(4) %1, align 4
38+
// CHECK: %arrayinit.next = add nuw i64 %arrayinit.index, 1
39+
// CHECK: %arrayinit.done = icmp eq i64 %arrayinit.next, 2
40+
// CHECK: br i1 %arrayinit.done, label %arrayinit.end, label %arrayinit.body
41+
// CHECK: arrayinit.end: ; preds = %arrayinit.body
42+
43+
int *ArrayOfPointers[2];
44+
myQueue.submit([&](sycl::handler &h) {
45+
h.single_task<class PtrArray>(
46+
[=] {
47+
int local = *ArrayOfPointers[1];
48+
});
49+
});
50+
// CHECK-LABEL: @{{.*}}PtrArray(ptr {{.*}}byval(%struct.__wrapper_class.0)
51+
// CHECK: %__SYCLKernel = alloca %class.anon.1, align 8
52+
// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
53+
// CHECK: %_arg_ArrayOfPointers.ascast = addrspacecast ptr %_arg_ArrayOfPointers to ptr addrspace(4)
54+
// CHECK: %ArrayOfPointers = getelementptr inbounds nuw %class.anon.1, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
55+
// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
56+
// CHECK: %arrayidx = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %0, i64 0, i64 0
57+
// CHECK: %1 = load ptr addrspace(1), ptr addrspace(4) %arrayidx, align 8
58+
// CHECK: %2 = addrspacecast ptr addrspace(1) %1 to ptr addrspace(4)
59+
// CHECK: store ptr addrspace(4) %2, ptr addrspace(4) %ArrayOfPointers, align 8
60+
// CHECK: %arrayinit.element = getelementptr inbounds ptr addrspace(4), ptr addrspace(4) %ArrayOfPointers, i64 1
61+
// CHECK: %3 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
62+
// CHECK: %arrayidx1 = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %3, i64 0, i64 1
63+
// CHECK: %4 = load ptr addrspace(1), ptr addrspace(4) %arrayidx1, align 8
64+
// CHECK: %5 = addrspacecast ptr addrspace(1) %4 to ptr addrspace(4)
65+
// CHECK: store ptr addrspace(4) %5, ptr addrspace(4) %arrayinit.element, align 8
66+
}

‎clang/test/SemaSYCL/array-kernel-param.cpp‎

Lines changed: 20 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -159,20 +159,16 @@ int main() {
159159
// CHECK-NEXT: InitListExpr
160160
// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
161161
// Initializer for ArrayOfPointers[0]
162-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
163-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
164-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
165-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
162+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
163+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
166164
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
167165
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
168166
// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
169167
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers'
170168
// CHECK-NEXT: IntegerLiteral {{.*}} 0
171169
// Initializer for ArrayOfPointers[1]
172-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
173-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
174-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
175-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
170+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
171+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
176172
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
177173
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
178174
// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
@@ -317,10 +313,8 @@ int main() {
317313
// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][3]'
318314
// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]'
319315
// Initializer for ArrayOfPointers_2D[0][0]
320-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
321-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
322-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
323-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
316+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
317+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
324318
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
325319
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
326320
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
@@ -331,10 +325,8 @@ int main() {
331325
// CHECK-NEXT: IntegerLiteral {{.*}} 0
332326

333327
// Initializer for ArrayOfPointers_2D[0][1]
334-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
335-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
336-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
337-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
328+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
329+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
338330
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
339331
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
340332
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
@@ -345,10 +337,8 @@ int main() {
345337
// CHECK-NEXT: IntegerLiteral {{.*}} 1
346338

347339
// Initializer for ArrayOfPointers_2D[0][2]
348-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
349-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
350-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
351-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
340+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
341+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
352342
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
353343
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
354344
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
@@ -361,10 +351,8 @@ int main() {
361351
// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]'
362352

363353
// Initializer for ArrayOfPointers_2D[1][0]
364-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
365-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
366-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
367-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
354+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
355+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
368356
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
369357
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
370358
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
@@ -375,10 +363,8 @@ int main() {
375363
// CHECK-NEXT: IntegerLiteral {{.*}} 0
376364

377365
// Initializer for ArrayOfPointers_2D[1][1]
378-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
379-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
380-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
381-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
366+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
367+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
382368
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
383369
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
384370
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
@@ -389,10 +375,8 @@ int main() {
389375
// CHECK-NEXT: IntegerLiteral {{.*}} 1
390376

391377
// Initializer for ArrayOfPointers_2D[1][2]
392-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
393-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
394-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
395-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
378+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
379+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
396380
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
397381
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
398382
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *[3]' lvalue
@@ -405,21 +389,17 @@ int main() {
405389
// Initializer for ArrayOfPointers
406390
// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
407391
// Initializer for ArrayOfPointers[0]
408-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
409-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
410-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
411-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
392+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
393+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
412394
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
413395
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
414396
// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
415397
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPointers'
416398
// CHECK-NEXT: IntegerLiteral {{.*}} 0
417399

418400
// Initializer for ArrayOfPointers[1]
419-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
420-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
421-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
422-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
401+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
402+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
423403
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
424404
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
425405
// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .

‎clang/test/SemaSYCL/built-in-type-kernel-arg.cpp‎

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -126,19 +126,15 @@ int main() {
126126
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
127127
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_data_addr' '__global int *'
128128
// CHECK: InitListExpr {{.*}} 'int *[2]'
129-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
130-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
131-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
132-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
129+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
130+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
133131
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
134132
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
135133
// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .
136134
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array'
137135
// CHECK-NEXT: IntegerLiteral {{.*}} 0
138-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <LValueToRValue>
139-
// CHECK-NEXT: UnaryOperator {{.*}} 'int *' lvalue prefix '*' cannot overflow
140-
// CHECK-NEXT: CXXReinterpretCastExpr {{.*}} 'int **' reinterpret_cast<int **> <BitCast>
141-
// CHECK-NEXT: UnaryOperator {{.*}} '__global int **' prefix '&' cannot overflow
136+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
137+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
142138
// CHECK-NEXT: ArraySubscriptExpr {{.*}} '__global int *' lvalue
143139
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int **' <ArrayToPointerDecay>
144140
// CHECK-NEXT: MemberExpr {{.*}} '__global int *[2]' lvalue .

0 commit comments

Comments
 (0)