Skip to content

Commit 0e455c9

Browse files
premanandraopremanandrao
authored
[SYCL] Null VarDecl dereference when a binding declaration is captured (#6924)
After the support for capturing structured bindings in lambdas, variables captured in BuildCaptureField need not be VarDecls. A previous patch needs to now account for a possible null pointer before dereferencing the pointer to get its name string.
1 parent 08b2022 commit 0e455c9

File tree

4 files changed

+148
-2
lines changed

4 files changed

+148
-2
lines changed

‎clang/lib/Sema/SemaLambda.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1714,7 +1714,9 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,
17141714

17151715
TypeSourceInfo *TSI = nullptr;
17161716
if (Capture.isVariableCapture()) {
1717-
const auto *Var = dyn_cast_or_null<VarDecl>(Capture.getVariable());
1717+
ValueDecl *Val = Capture.getVariable();
1718+
const auto *Var = dyn_cast_or_null<VarDecl>(Val);
1719+
17181720
if (Var && Var->isInitCapture())
17191721
TSI = Var->getTypeSourceInfo();
17201722

@@ -1723,7 +1725,7 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,
17231725
// For SYCL compilations, save user specified names for
17241726
// lambda capture.
17251727
if (getLangOpts().SYCLIsDevice || getLangOpts().SYCLIsHost) {
1726-
StringRef CaptureName = Var->getName();
1728+
StringRef CaptureName = Val ? Val->getName() : "";
17271729
if (!CaptureName.empty())
17281730
Id = &Context.Idents.get(CaptureName.str());
17291731
}

‎clang/test/CodeGenSYCL/kernel_binding_decls.cpp

Lines changed: 61 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,61 @@
1+
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -std=c++20 -triple spir64-unknown-unknown -disable-llvm-passes -opaque-pointers -emit-llvm %s -o - | FileCheck %s
2+
3+
#include "sycl.hpp"
4+
5+
// This test checks that we correctly capture binding declarations.
6+
7+
void foo() {
8+
sycl::handler h;
9+
int a[2] = {1, 2};
10+
auto [x, y] = a;
11+
struct S {
12+
float b[3] = { 0, 3.0f, 4.0 };
13+
} s;
14+
auto [f1, f2, f3] = s.b;
15+
auto Lambda = [=]() { x = 10; f2 = 2.3f; };
16+
h.single_task(Lambda);
17+
}
18+
19+
// CHECK: %class.anon = type { i32, float }
20+
21+
// Check the sycl kernel arguments - one int and one float parameter
22+
// CHECK: define {{.*}} spir_kernel void @{{.*}}foov{{.*}}(i32 {{.*}} %_arg_x, float {{.*}} %_arg_f2)
23+
// CHECK: entry:
24+
25+
// Check alloca of the captured types
26+
// CHECK: %_arg_x.addr = alloca i32, align 4
27+
// CHECK: %_arg_f2.addr = alloca float, align 4
28+
// CHECK: %__SYCLKernel = alloca %class.anon, align 4
29+
30+
// Copy the parameters into the alloca-ed addresses
31+
// CHECK: store i32 %_arg_x, ptr addrspace(4) %_arg_x.addr
32+
// CHECK: store float %_arg_f2, ptr addrspace(4) %_arg_f2.addr
33+
34+
// Store the int and the float into the struct created
35+
// CHECK: %x = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 0
36+
// CHECK: %0 = load i32, ptr addrspace(4) %_arg_x.addr
37+
// CHECK: store i32 %0, ptr addrspace(4) %x
38+
// CHECK: %f2 = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 1
39+
// CHECK: %1 = load float, ptr addrspace(4) %_arg_f2.addr
40+
// CHECK: store float %1, ptr addrspace(4) %f2
41+
42+
// Call the lambda
43+
// CHECK: call spir_func void @{{.*}}foo{{.*}}(ptr addrspace(4) {{.*}} %__SYCLKernel{{.*}})
44+
// CHECK: ret void
45+
46+
// Check the lambda call
47+
// CHECK: define {{.*}} spir_func void @{{.*}}foo{{.*}}(ptr addrspace(4) {{.*}} %this)
48+
// CHECK: entry:
49+
// CHECK: %this.addr = alloca ptr addrspace(4)
50+
// CHECK: %this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
51+
// CHECK: store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast
52+
// CHECK: %this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast
53+
54+
// Check the store of 10 into the int value
55+
// CHECK: %x = getelementptr inbounds %class.anon, ptr addrspace(4) %this1, i32 0, i32 0
56+
// CHECK: store i32 10, ptr addrspace(4) %x
57+
58+
// Check the store of 2.3f into the float value
59+
// CHECK: %f2 = getelementptr inbounds %class.anon, ptr addrspace(4) %this1, i32 0, i32 1
60+
// CHECK: store float 0x4002666660000000, ptr addrspace(4) %f2
61+
// CHECK: ret void

‎clang/test/CodeGenSYCL/no-opaque-ptr-kernel_binding_decls.cpp

Lines changed: 61 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,61 @@
1+
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -std=c++20 -triple spir64-unknown-unknown -disable-llvm-passes -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s
2+
3+
#include "sycl.hpp"
4+
5+
// This test checks that we correctly capture binding declarations.
6+
7+
void foo() {
8+
sycl::handler h;
9+
int a[2] = {1, 2};
10+
auto [x, y] = a;
11+
struct S {
12+
float b[3] = { 0, 3.0f, 4.0 };
13+
} s;
14+
auto [f1, f2, f3] = s.b;
15+
auto Lambda = [=]() { x = 10; f2 = 2.3f; };
16+
h.single_task(Lambda);
17+
}
18+
19+
// CHECK: %class.anon = type { i32, float }
20+
21+
// Check the sycl kernel arguments - one int and one float parameter
22+
// CHECK: define {{.*}} spir_kernel void @{{.*}}foov{{.*}}(i32 {{.*}} %_arg_x, float {{.*}} %_arg_f2)
23+
// CHECK: entry:
24+
25+
// Check alloca of the captured types
26+
// CHECK: %_arg_x.addr = alloca i32, align 4
27+
// CHECK: %_arg_f2.addr = alloca float, align 4
28+
// CHECK: %__SYCLKernel = alloca %class.anon, align 4
29+
30+
// Copy the parameters into the alloca-ed addresses
31+
// CHECK: store i32 %_arg_x, i32 addrspace(4)* %_arg_x.addr
32+
// CHECK: store float %_arg_f2, float addrspace(4)* %_arg_f2.addr
33+
34+
// Store the int and the float into the struct created
35+
// CHECK: %x = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %__SYCLKernel{{.*}}, i32 0, i32 0
36+
// CHECK: %1 = load i32, i32 addrspace(4)* %_arg_x.addr
37+
// CHECK: store i32 %1, i32 addrspace(4)* %x
38+
// CHECK: %f2 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %__SYCLKernel{{.*}}, i32 0, i32 1
39+
// CHECK: %2 = load float, float addrspace(4)* %_arg_f2.addr
40+
// CHECK: store float %2, float addrspace(4)* %f2
41+
42+
// Call the lambda
43+
// CHECK: call spir_func void @{{.*}}foo{{.*}}(%class.anon addrspace(4)* {{.*}} %__SYCLKernel{{.*}})
44+
// CHECK: ret void
45+
46+
// Check the lambda call
47+
// CHECK: define {{.*}} spir_func void @{{.*}}foo{{.*}}(%class.anon addrspace(4)* {{.*}} %this)
48+
// CHECK: entry:
49+
// CHECK: %this.addr = alloca %class.anon addrspace(4)*
50+
// CHECK: %this.addr.ascast = addrspacecast %class.anon addrspace(4)** %this.addr to %class.anon addrspace(4)* addrspace(4)*
51+
// CHECK: store %class.anon addrspace(4)* %this, %class.anon addrspace(4)* addrspace(4)* %this.addr.ascast
52+
// CHECK: %this1 = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* %this.addr.ascast
53+
54+
// Check the store of 10 into the int value
55+
// CHECK: %x = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %this1, i32 0, i32 0
56+
// CHECK: store i32 10, i32 addrspace(4)* %x
57+
58+
// Check the store of 2.3f into the float value
59+
// CHECK: %f2 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %this1, i32 0, i32 1
60+
// CHECK: store float 0x4002666660000000, float addrspace(4)* %f2
61+
// CHECK: ret void

‎clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -std=c++20 -fsyntax-only %s -verify=device -ast-dump | FileCheck %s
2+
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-host -std=c++20 -fsyntax-only %s -verify=host
3+
4+
// This test checks that when a binding declaration is captured that
5+
// we don't dereference the null VarDecl. Also checks that the kernel
6+
// parameter has the name of the binding declaration associated with it.
7+
8+
#include "sycl.hpp"
9+
10+
// host-no-diagnostics
11+
// device-no-diagnostics
12+
13+
void foo() {
14+
int a[2] = {1, 2};
15+
auto [bind_x, bind_y] = a;
16+
auto Lambda = [=]() { bind_x = 10; };
17+
sycl::handler h;
18+
h.single_task<class C>(Lambda);
19+
}
20+
21+
// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int)'
22+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_bind_x 'int'

0 commit comments

Comments
 (0)