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"

‎clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp

Lines changed: 16 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -93,21 +93,21 @@ int main() {
9393
// Check kernel_A parameters
9494
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
9595
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
96-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
97-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
98-
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]],
99-
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+_4]],
100-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]],
101-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]],
102-
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]])
96+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
97+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
98+
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
99+
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
100+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
101+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
102+
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
103103
// CHECK-SAME: !kernel_arg_runtime_aligned !5
104104

105105
// Check kernel_readOnlyAcc parameters
106106
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_readOnlyAcc
107107
// CHECK-SAME: i32 addrspace(1)* noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
108-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
109-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
110-
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]]
108+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
109+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
110+
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
111111
// CHECK-SAME: !kernel_arg_runtime_aligned !14
112112

113113
// Check kernel_B parameters
@@ -127,17 +127,17 @@ int main() {
127127

128128
// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
129129
// CHECK-SAME: float addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
130-
// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
131-
// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
132-
// CHECK-SAME: %"struct.cl::sycl::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]]
130+
// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
131+
// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
132+
// CHECK-SAME: %"struct.cl::sycl::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
133133
// CHECK-SAME: !kernel_arg_runtime_aligned !14
134134

135135
// Check kernel_acc_raw_ptr parameters
136136
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr
137137
// CHECK-SAME: i32 addrspace(1)* noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
138-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
139-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
140-
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]]
138+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
139+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
140+
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
141141
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
142142
// CHECK-SAME: !kernel_arg_runtime_aligned !26
143143

‎clang/test/CodeGenSYCL/kernel-handler.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ void test(int val) {
2323
}
2424

2525
// ALL: define dso_local{{ spir_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
26-
// ALL-SAME: (i32 noundef %_arg_, i8 addrspace(1)* noundef align 1 %_arg__specialization_constants_buffer)
26+
// ALL-SAME: (i32 noundef %_arg_a, i8 addrspace(1)* noundef align 1 %_arg__specialization_constants_buffer)
2727
// ALL: %kh = alloca %"class.cl::sycl::kernel_handler", align 1
2828

2929
// NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)** %_arg__specialization_constants_buffer.addr, align 8

‎clang/test/CodeGenSYCL/kernel-param-acc-array.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -26,13 +26,13 @@ int main() {
2626
// Check kernel_A parameters
2727
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
2828
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
29-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
30-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
31-
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]],
32-
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+_4]],
33-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]],
34-
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]],
35-
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]])
29+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
30+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
31+
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
32+
// CHECK-SAME: i32 addrspace(1)* noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
33+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
34+
// CHECK-SAME: %"struct.cl::sycl::range"* noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
35+
// CHECK-SAME: %"struct.cl::sycl::id"* noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
3636

3737
// CHECK alloca for pointer arguments
3838
// CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8

‎clang/test/CodeGenSYCL/pointers-in-structs.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -45,4 +45,4 @@ int main() {
4545
// CHECK-SAME: %[[WRAPPER_F]]* noundef byval(%[[WRAPPER_F]]) align 8 %_arg_F,
4646
// CHECK-SAME: %[[WRAPPER_F4_1]]* noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4
4747
// CHECK-SAME: %[[WRAPPER_F4_2]]* noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41
48-
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_)
48+
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Ptr)

‎clang/test/CodeGenSYCL/sampler.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
88
// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8*
99
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
10-
// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANONCAST]], i32 0, i32 0
10+
// CHECK-NEXT: [[GEP:%[a-zA-z0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANONCAST]], i32 0, i32 0
1111
// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
1212
// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
1313
//
@@ -25,13 +25,13 @@
2525
// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4
2626

2727
// Initialize 'a'
28-
// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0
28+
// CHECK: [[GEP_LAMBDA:%[a-zA-z0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0
2929
// CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.sampler_wrapper, %struct.sampler_wrapper addrspace(4)* [[GEP_LAMBDA]], i32 0, i32 1
3030
// CHECK: [[LOAD_A:%[0-9]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4
3131
// CHECK: store i32 [[LOAD_A]], i32 addrspace(4)* [[GEP_A]], align 8
3232

3333
// Initialize wrapped sampler 'smpl'
34-
// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0
34+
// CHECK: [[GEP_LAMBDA_0:%[a-zA-z0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0
3535
// CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.sampler_wrapper, %struct.sampler_wrapper addrspace(4)* [[GEP_LAMBDA_0]], i32 0, i32 0
3636
// CHECK: [[LOAD_SMPL:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8
3737
// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SMPL]])

‎clang/test/CodeGenSYCL/spir-enum.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ void test(enum_type val)
2020

2121
int main() {
2222

23-
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 noundef %_arg_)
23+
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 noundef %_arg_val)
2424

2525
// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)*
2626
// CHECK: call spir_func void @_ZZ4test9enum_typeENKUlvE_clEv(%class.anon addrspace(4)* {{[^,]*}} %{{.+}})

0 commit comments

Comments
 (0)