diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp index d5b1c2e564938..e09320ec6c29a 100644 --- a/clang/lib/Sema/SemaLambda.cpp +++ b/clang/lib/Sema/SemaLambda.cpp @@ -1714,7 +1714,9 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD, TypeSourceInfo *TSI = nullptr; if (Capture.isVariableCapture()) { - const auto *Var = dyn_cast_or_null(Capture.getVariable()); + ValueDecl *Val = Capture.getVariable(); + const auto *Var = dyn_cast_or_null(Val); + if (Var && Var->isInitCapture()) TSI = Var->getTypeSourceInfo(); @@ -1723,7 +1725,7 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD, // For SYCL compilations, save user specified names for // lambda capture. if (getLangOpts().SYCLIsDevice || getLangOpts().SYCLIsHost) { - StringRef CaptureName = Var->getName(); + StringRef CaptureName = Val ? Val->getName() : ""; if (!CaptureName.empty()) Id = &Context.Idents.get(CaptureName.str()); } diff --git a/clang/test/CodeGenSYCL/kernel_binding_decls.cpp b/clang/test/CodeGenSYCL/kernel_binding_decls.cpp new file mode 100644 index 0000000000000..f6434057f98ec --- /dev/null +++ b/clang/test/CodeGenSYCL/kernel_binding_decls.cpp @@ -0,0 +1,61 @@ +// 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 + +#include "sycl.hpp" + +// 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_x, float {{.*}} %_arg_f2) +// CHECK: entry: + +// Check alloca of the captured types +// CHECK: %_arg_x.addr = alloca i32, align 4 +// CHECK: %_arg_f2.addr = alloca float, align 4 +// CHECK: %__SYCLKernel = alloca %class.anon, align 4 + +// Copy the parameters into the alloca-ed addresses +// CHECK: store i32 %_arg_x, ptr addrspace(4) %_arg_x.addr +// CHECK: store float %_arg_f2, ptr addrspace(4) %_arg_f2.addr + +// Store the int and the float into the struct created +// CHECK: %x = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 0 +// CHECK: %0 = load i32, ptr addrspace(4) %_arg_x.addr +// CHECK: store i32 %0, ptr addrspace(4) %x +// CHECK: %f2 = getelementptr inbounds %class.anon, ptr addrspace(4) %__SYCLKernel{{.*}}, i32 0, i32 1 +// CHECK: %1 = load float, ptr addrspace(4) %_arg_f2.addr +// CHECK: store float %1, ptr addrspace(4) %f2 + +// 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: %x = getelementptr inbounds %class.anon, ptr addrspace(4) %this1, i32 0, i32 0 +// CHECK: store i32 10, ptr addrspace(4) %x + +// Check the store of 2.3f into the float value +// CHECK: %f2 = getelementptr inbounds %class.anon, ptr addrspace(4) %this1, i32 0, i32 1 +// CHECK: store float 0x4002666660000000, ptr addrspace(4) %f2 +// CHECK: ret void diff --git a/clang/test/CodeGenSYCL/no-opaque-ptr-kernel_binding_decls.cpp b/clang/test/CodeGenSYCL/no-opaque-ptr-kernel_binding_decls.cpp new file mode 100644 index 0000000000000..d82f7caf54657 --- /dev/null +++ b/clang/test/CodeGenSYCL/no-opaque-ptr-kernel_binding_decls.cpp @@ -0,0 +1,61 @@ +// 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 + +#include "sycl.hpp" + +// 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_x, float {{.*}} %_arg_f2) +// CHECK: entry: + +// Check alloca of the captured types +// CHECK: %_arg_x.addr = alloca i32, align 4 +// CHECK: %_arg_f2.addr = alloca float, align 4 +// CHECK: %__SYCLKernel = alloca %class.anon, align 4 + +// Copy the parameters into the alloca-ed addresses +// CHECK: store i32 %_arg_x, i32 addrspace(4)* %_arg_x.addr +// CHECK: store float %_arg_f2, float addrspace(4)* %_arg_f2.addr + +// Store the int and the float into the struct created +// CHECK: %x = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %__SYCLKernel{{.*}}, i32 0, i32 0 +// CHECK: %1 = load i32, i32 addrspace(4)* %_arg_x.addr +// CHECK: store i32 %1, i32 addrspace(4)* %x +// CHECK: %f2 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %__SYCLKernel{{.*}}, i32 0, i32 1 +// CHECK: %2 = load float, float addrspace(4)* %_arg_f2.addr +// CHECK: store float %2, float addrspace(4)* %f2 + +// 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: %x = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %this1, i32 0, i32 0 +// CHECK: store i32 10, i32 addrspace(4)* %x + +// Check the store of 2.3f into the float value +// CHECK: %f2 = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %this1, i32 0, i32 1 +// CHECK: store float 0x4002666660000000, float addrspace(4)* %f2 +// CHECK: ret void diff --git a/clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp b/clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp new file mode 100644 index 0000000000000..e9dc6cd5527fd --- /dev/null +++ b/clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp @@ -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(Lambda); +} + +// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_bind_x 'int'