Skip to content

Commit af29982

Browse files
elizabethandrewselizabethandrews
authored
[SYCL] Save user specified names in lambda object (#5772)
This patch retains user specified names for lambda captures in lambda object. As a result, the name of openCL kernel arguments generated for SYCL kernel specified as a lamdba, now includes the user names in kernel argument name (matches current behavior for SYCL kernel specified as a functor object). Signed-off-by: Elizabeth Andrews <[email protected]>
1 parent b2ee289 commit af29982

30 files changed

+188
-163
lines changed

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7730,6 +7730,8 @@ let CategoryName = "Lambda Issue" in {
77307730
"%select{| explicitly}1 captured here">;
77317731
def err_implicit_this_capture : Error<
77327732
"implicit capture of 'this' is not allowed for kernel functions">;
7733+
def err_lambda_member_access : Error<
7734+
"invalid attempt to access member of lambda">;
77337735

77347736
// C++14 lambda init-captures.
77357737
def warn_cxx11_compat_init_capture : Warning<

‎clang/lib/Sema/SemaAccess.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1903,7 +1903,13 @@ void Sema::CheckLookupAccess(const LookupResult &R) {
19031903
AccessTarget Entity(Context, AccessedEntity::Member,
19041904
R.getNamingClass(), I.getPair(),
19051905
R.getBaseObjectType());
1906-
Entity.setDiag(diag::err_access);
1906+
// This is to avoid leaking implementation details of lambda object.
1907+
// We do not want to generate 'private member access' diagnostic for
1908+
// lambda object.
1909+
if ((R.getNamingClass())->isLambda())
1910+
Diag(R.getNameLoc(), diag::err_lambda_member_access);
1911+
else
1912+
Entity.setDiag(diag::err_access);
19071913
CheckAccess(*this, R.getNameLoc(), Entity);
19081914
}
19091915
}

‎clang/lib/Sema/SemaLambda.cpp

Lines changed: 13 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1702,12 +1702,23 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,
17021702
const sema::Capture &Capture) {
17031703
SourceLocation Loc = Capture.getLocation();
17041704
QualType FieldType = Capture.getCaptureType();
1705+
IdentifierInfo *Id = nullptr;
17051706

17061707
TypeSourceInfo *TSI = nullptr;
17071708
if (Capture.isVariableCapture()) {
17081709
auto *Var = Capture.getVariable();
17091710
if (Var->isInitCapture())
17101711
TSI = Capture.getVariable()->getTypeSourceInfo();
1712+
1713+
// TODO: Upstream this behavior to LLVM project to save
1714+
// user speciifed names for all lambdas.
1715+
// For SYCL compilations, save user specified names for
1716+
// lambda capture.
1717+
if (getLangOpts().SYCLIsDevice || getLangOpts().SYCLIsHost) {
1718+
StringRef CaptureName = Var->getName();
1719+
if (!CaptureName.empty())
1720+
Id = &Context.Idents.get(CaptureName.str());
1721+
}
17111722
}
17121723

17131724
// FIXME: Should we really be doing this? A null TypeSourceInfo seems more
@@ -1717,8 +1728,8 @@ FieldDecl *Sema::BuildCaptureField(RecordDecl *RD,
17171728

17181729
// Build the non-static data member.
17191730
FieldDecl *Field =
1720-
FieldDecl::Create(Context, RD, /*StartLoc=*/Loc, /*IdLoc=*/Loc,
1721-
/*Id=*/nullptr, FieldType, TSI, /*BW=*/nullptr,
1731+
FieldDecl::Create(Context, RD, /*StartLoc=*/Loc, /*IdLoc=*/Loc, Id,
1732+
FieldType, TSI, /*BW=*/nullptr,
17221733
/*Mutable=*/false, ICIS_NoInit);
17231734
// If the variable being captured has an invalid type, mark the class as
17241735
// invalid as well.

‎clang/test/CXX/expr/expr.prim/expr.prim.lambda/p11-1y.cpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,12 @@
1-
// RUN: %clang_cc1 -std=c++1y %s -verify
1+
// TODO: The SYCL changes in this test can be removed if/when changes (i.e.
2+
// PR to save user-specified names in lambda class) is upstreamed to LLVM
3+
// project.
24

3-
const char *has_no_member = [x("hello")] {}.x; // expected-error {{no member named 'x'}}
5+
// RUN: %clang_cc1 -std=c++1y %s -verify=notsycl,expected
6+
// RUN: %clang_cc1 -fsycl-is-device -std=c++1y %s -verify=sycl,expected
7+
8+
const char *has_no_member = [x("hello")] {}.x; // notsycl-error {{no member named 'x'}}
9+
// sycl-error@-1 {{invalid attempt to access member of lambda}}
410

511
double f;
612
auto with_float = [f(1.0f)] {

‎clang/test/CodeGenSYCL/accessor-readonly.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -14,8 +14,8 @@ void f0(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
1414

1515
// CHECK: spir_kernel{{.*}}f1_kernel
1616
// CHECK-NOT: readonly
17-
// CHECK-SAME: %_arg_{{.*}}%_arg_1{{.*}}%_arg_2{{.*}}%_arg_3
18-
// CHECK-SAME: readonly align 4 %_arg_4
17+
// CHECK-SAME: %_arg_write_acc{{.*}}%_arg_write_acc1{{.*}}%_arg_write_acc2{{.*}}%_arg_write_acc3
18+
// CHECK-SAME: readonly align 4 %_arg_read_acc
1919
void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
2020
myQueue.submit([&](cl::sycl::handler &cgh) {
2121
auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
@@ -25,9 +25,9 @@ void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
2525
}
2626

2727
// CHECK: spir_kernel{{.*}}f2_kernel
28-
// CHECK-SAME: readonly align 4 %_arg_
28+
// CHECK-SAME: readonly align 4 %_arg_read_acc
2929
// CHECK-NOT: readonly
30-
// CHECK-SAME: %_arg_8
30+
// CHECK-SAME: %_arg_write_acc
3131
void f2(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
3232
myQueue.submit([&](cl::sycl::handler &cgh) {
3333
auto read_acc = in_buf.get_access<cl::sycl::access::mode::read>(cgh);

‎clang/test/CodeGenSYCL/accessor_no_alias_property.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ int main() {
2020
accessorB;
2121

2222
// Check that noalias parameter attribute is emitted when no_alias accessor property is used
23-
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_, {{.*}})
23+
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_accessorA, {{.*}})
2424
cl::sycl::kernel_single_task<class kernel_function1>(
2525
[=]() {
2626
accessorA.use();

‎clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,8 +21,8 @@ int main() {
2121

2222
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_function
2323
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]],
24-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]],
25-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]],
24+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+1]],
25+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+2]],
2626
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
2727
// Check alloca for pointer argument
2828
// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*

‎clang/test/CodeGenSYCL/device-variables.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,10 +32,10 @@ int main() {
3232
// CHECK: store i32 1, i32 addrspace(4)* %b
3333
foo(local_value);
3434
// Local variables and constexprs captured by lambda
35-
// CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0
35+
// CHECK: [[GEP:%[a-z_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0
3636
// CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* noundef align 4 dereferenceable(4) [[GEP]])
3737
int some_device_local_var = some_local_var;
38-
// CHECK: [[GEP1:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 1
38+
// CHECK: [[GEP1:%[a-z_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 1
3939
// CHECK: [[LOAD1:%[0-9]+]] = load i32, i32 addrspace(4)* [[GEP1]]
4040
// CHECK: store i32 [[LOAD1]], i32 addrspace(4)* %some_device_local_var
4141
});

‎clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ void test(int val) {
3030
// --- Attributes
3131
// CHECK: {{.*}} !kernel_arg_accessor_ptr ![[ACC_PTR_ATTR:[0-9]+]] !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
3232
// --- init_esimd call is expected instead of __init:
33-
// CHECK: call spir_func void @{{.*}}__init_esimd{{.*}}(%"{{.*}}sycl::accessor" addrspace(4)* {{[^,]*}} %{{[0-9]+}}, i32 addrspace(1)* noundef %{{[0-9]+}})
33+
// CHECK: call spir_func void @{{.*}}__init_esimd{{.*}}(%"{{.*}}sycl::accessor" addrspace(4)* {{[^,]*}} %{{[a-zA-Z0-9_]+}}, i32 addrspace(1)* noundef %{{[0-9]+}})
3434
// CHECK-LABEL: }
3535
// CHECK: ![[ACC_PTR_ATTR]] = !{i1 true, i1 false, i1 true}
3636
}

‎clang/test/CodeGenSYCL/image_accessor.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -8,27 +8,27 @@
88
//
99
// CHECK-1DRO: %opencl.image1d_ro_t = type opaque
1010
// CHECK-1DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
11-
// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
11+
// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
1212
//
1313
// CHECK-2DRO: %opencl.image2d_ro_t = type opaque
1414
// CHECK-2DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
15-
// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
15+
// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
1616
//
1717
// CHECK-3DRO: %opencl.image3d_ro_t = type opaque
1818
// CHECK-3DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
19-
// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
19+
// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
2020
//
2121
// CHECK-1DWO: %opencl.image1d_wo_t = type opaque
2222
// CHECK-1DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
23-
// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
23+
// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
2424
//
2525
// CHECK-2DWO: %opencl.image2d_wo_t = type opaque
2626
// CHECK-2DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
27-
// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
27+
// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
2828
//
2929
// CHECK-3DWO: %opencl.image3d_wo_t = type opaque
3030
// CHECK-3DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
31-
// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
31+
// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[a-zA-Z]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
3232
//
3333
// TODO: Add tests for the image_array opencl datatype support.
3434
#include "Inputs/sycl.hpp"

0 commit comments

Comments
 (0)