[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.

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<VarDecl>(Capture.getVariable());
+    ValueDecl *Val = Capture.getVariable();
+    const auto *Var = dyn_cast_or_null<VarDecl>(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());
     }

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

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

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<class C>(Lambda);
+}
+
+// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int)'
+// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_bind_x 'int'