Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[SYCL] Null VarDecl dereference when a binding declaration is captured #6924

Merged
merged 7 commits into from
Oct 7, 2022
Merged

[SYCL] Null VarDecl dereference when a binding declaration is captured #6924

Show file tree
Hide file tree
Changes from 5 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
6f290fd
[SYCL] Null VarDecl dereference when a binding declaration is captured
premanandrao Sep 30, 2022
d819062
Address code-review comments
premanandrao Oct 4, 2022
e3313ea
Address code-review comments (missed one case)
premanandrao Oct 4, 2022
ebc1521
Add no-opaque-ptr variant
premanandrao Oct 4, 2022
10525b7
Get the name of the binding declaration
premanandrao Oct 5, 2022
dc44388
Use name from ValueDecl
premanandrao Oct 5, 2022
2f5f80e
Fix CodeGenSYCL tests to use includes properly
premanandrao Oct 6, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 8 additions & 2 deletions clang/lib/Sema/SemaLambda.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1714,7 +1714,9 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,

TypeSourceInfo *TSI = nullptr;
if (Capture.isVariableCapture()) {
const auto *Var = dyn_cast_or_null<VarDecl>(Capture.getVariable());
ValueDecl *Val = Capture.getVariable();
const auto *Var = dyn_cast_or_null<VarDecl>(Val);

if (Var && Var->isInitCapture())
TSI = Var->getTypeSourceInfo();

Expand All @@ -1723,7 +1725,11 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,
// For SYCL compilations, save user specified names for
// lambda capture.
if (getLangOpts().SYCLIsDevice || getLangOpts().SYCLIsHost) {
StringRef CaptureName = Var->getName();
BindingDecl *BD = nullptr;
if (!Var)
BD = dyn_cast_or_null<BindingDecl>(Val);
Fznamznon marked this conversation as resolved.
Show resolved Hide resolved

StringRef CaptureName = Var ? Var->getName() : BD ? BD->getName() : "";
if (!CaptureName.empty())
Id = &Context.Idents.get(CaptureName.str());
}
Expand Down
61 changes: 61 additions & 0 deletions clang/test/CodeGenSYCL/kernel_binding_decls.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,61 @@
// RUN: %clang_cc1 -internal-isystem Inputs -fsycl-is-device -std=c++20 -triple spir64-unknown-unknown -disable-llvm-passes -opaque-pointers -emit-llvm %s -o - | FileCheck %s
premanandrao marked this conversation as resolved.
Show resolved Hide resolved

#include "Inputs/sycl.hpp"
premanandrao marked this conversation as resolved.
Show resolved Hide resolved

// This test checks that we correctly capture binding declarations.

void foo() {
sycl::handler h;
int a[2] = {1, 2};
auto [x, y] = a;
struct S {
float b[3] = { 0, 3.0f, 4.0 };
} s;
auto [f1, f2, f3] = s.b;
auto Lambda = [=]() { x = 10; f2 = 2.3f; };
h.single_task(Lambda);
}

// CHECK: %class.anon = type { i32, float }

// Check the sycl kernel arguments - one int and one float parameter
// CHECK: define {{.*}} spir_kernel void @{{.*}}foov{{.*}}(i32 {{.*}} %_arg_, float {{.*}} %_arg_1)
// CHECK: entry:

// Check alloca of the captured types
// CHECK: %_arg_.addr = alloca i32, align 4
// CHECK: %_arg_.addr2 = alloca float, align 4
// CHECK: %__SYCLKernel = alloca %class.anon, align 4

// Copy the parameters into the alloca-ed addresses
// CHECK: store i32 %_arg_, ptr addrspace(4) %_arg_.addr
// CHECK: store float %_arg_1, ptr addrspace(4) %_arg_.addr2

// Store the int and the float into the struct created
// CHECK: %0 = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 0
// CHECK: %1 = load i32, ptr addrspace(4) %_arg_.addr
// CHECK: store i32 %1, ptr addrspace(4) %0
// CHECK: %2 = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 1
// CHECK: %3 = load float, ptr addrspace(4) %_arg_.addr2
// CHECK: store float %3, ptr addrspace(4) %2

// Call the lambda
// CHECK: call spir_func void @{{.*}}foo{{.*}}(ptr addrspace(4) {{.*}} %__SYCLKernel{{.*}})
// CHECK: ret void

// Check the lambda call
// CHECK: define {{.*}} spir_func void @{{.*}}foo{{.*}}(ptr addrspace(4) {{.*}} %this)
// CHECK: entry:
// CHECK: %this.addr = alloca ptr addrspace(4)
// CHECK: %this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
// CHECK: store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast
// CHECK: %this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast

// Check the store of 10 into the int value
// CHECK: %0 = getelementptr inbounds %class.anon, ptr addrspace(4) %this1, i32 0, i32 0
// CHECK: store i32 10, ptr addrspace(4) %0

// Check the store of 2.3f into the float value
// CHECK: %1 = getelementptr inbounds %class.anon, ptr addrspace(4) %this1, i32 0, i32 1
// CHECK: store float 0x4002666660000000, ptr addrspace(4) %1
// CHECK: ret void
61 changes: 61 additions & 0 deletions clang/test/CodeGenSYCL/no-opaque-ptr-kernel_binding_decls.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,61 @@
// RUN: %clang_cc1 -internal-isystem Inputs -fsycl-is-device -std=c++20 -triple spir64-unknown-unknown -disable-llvm-passes -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s
premanandrao marked this conversation as resolved.
Show resolved Hide resolved

#include "Inputs/sycl.hpp"
premanandrao marked this conversation as resolved.
Show resolved Hide resolved

// This test checks that we correctly capture binding declarations.

void foo() {
sycl::handler h;
int a[2] = {1, 2};
auto [x, y] = a;
struct S {
float b[3] = { 0, 3.0f, 4.0 };
} s;
auto [f1, f2, f3] = s.b;
auto Lambda = [=]() { x = 10; f2 = 2.3f; };
h.single_task(Lambda);
}

// CHECK: %class.anon = type { i32, float }

// Check the sycl kernel arguments - one int and one float parameter
// CHECK: define {{.*}} spir_kernel void @{{.*}}foov{{.*}}(i32 {{.*}} %_arg_, float {{.*}} %_arg_1)
// CHECK: entry:

// Check alloca of the captured types
// CHECK: %_arg_.addr = alloca i32, align 4
// CHECK: %_arg_.addr2 = alloca float, align 4
// CHECK: %__SYCLKernel = alloca %class.anon, align 4

// Copy the parameters into the alloca-ed addresses
// CHECK: store i32 %_arg_, i32 addrspace(4)* %_arg_.addr
// CHECK: store float %_arg_1, float addrspace(4)* %_arg_.addr2

// Store the int and the float into the struct created
// CHECK: %1 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %__SYCLKernel{{.*}}, i32 0, i32 0
// CHECK: %2 = load i32, i32 addrspace(4)* %_arg_.addr
// CHECK: store i32 %2, i32 addrspace(4)* %1
// CHECK: %3 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %__SYCLKernel{{.*}}, i32 0, i32 1
// CHECK: %4 = load float, float addrspace(4)* %_arg_.addr2
// CHECK: store float %4, float addrspace(4)* %3

// Call the lambda
// CHECK: call spir_func void @{{.*}}foo{{.*}}(%class.anon addrspace(4)* {{.*}} %__SYCLKernel{{.*}})
// CHECK: ret void

// Check the lambda call
// CHECK: define {{.*}} spir_func void @{{.*}}foo{{.*}}(%class.anon addrspace(4)* {{.*}} %this)
// CHECK: entry:
// CHECK: %this.addr = alloca %class.anon addrspace(4)*
// CHECK: %this.addr.ascast = addrspacecast %class.anon addrspace(4)** %this.addr to %class.anon addrspace(4)* addrspace(4)*
// CHECK: store %class.anon addrspace(4)* %this, %class.anon addrspace(4)* addrspace(4)* %this.addr.ascast
// CHECK: %this1 = load %class.anon addrspace(4)*, %class.anon addrspace(4)* addrspace(4)* %this.addr.ascast

// Check the store of 10 into the int value
// CHECK: %0 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %this1, i32 0, i32 0
// CHECK: store i32 10, i32 addrspace(4)* %0

// Check the store of 2.3f into the float value
// CHECK: %1 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %this1, i32 0, i32 1
// CHECK: store float 0x4002666660000000, float addrspace(4)* %1
// CHECK: ret void
22 changes: 22 additions & 0 deletions clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,22 @@
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -std=c++20 -fsyntax-only %s -verify=device -ast-dump | FileCheck %s
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-host -std=c++20 -fsyntax-only %s -verify=host

// This test checks that when a binding declaration is captured that
// we don't dereference the null VarDecl. Also checks that the kernel
// parameter has the name of the binding declaration associated with it.

#include "sycl.hpp"

// host-no-diagnostics
// device-no-diagnostics

void foo() {
int a[2] = {1, 2};
auto [bind_x, bind_y] = a;
auto Lambda = [=]() { bind_x = 10; };
sycl::handler h;
h.single_task<class C>(Lambda);
}

// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int)'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_bind_x 'int'