Skip to content

Commit fa4ba4c

Browse files
Fznamznonbader
Fznamznon
authored and
bader
committed
[SYCL] Refactor OpenCL kernel generation (#227)
This change unifies handling of sampler and accessor types during OpenCL kernel generation. Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
1 parent e844b0c commit fa4ba4c

11 files changed

+262
-304
lines changed

‎clang/include/clang/Sema/Sema.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -11266,7 +11266,7 @@ class Sema {
1126611266
return *SyclIntHeader.get();
1126711267
}
1126811268

11269-
void ConstructSYCLKernel(FunctionDecl *KernelCallerFunc);
11269+
void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc);
1127011270
void MarkDevice(void);
1127111271
};
1127211272

‎clang/lib/Sema/SemaSYCL.cpp

+203-247
Large diffs are not rendered by default.

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -5523,7 +5523,7 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) {
55235523
// so we are checking for SYCL kernel attribute after instantination.
55245524
if (getLangOpts().SYCLIsDevice &&
55255525
CurFD->hasAttr<SYCLKernelAttr>()) {
5526-
ConstructSYCLKernel(CurFD);
5526+
ConstructOpenCLKernel(CurFD);
55275527
}
55285528
CurFD->setInstantiationIsPending(false);
55295529
}
@@ -5537,7 +5537,7 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) {
55375537
// so we are checking for SYCL kernel attribute after instantination.
55385538
if (getLangOpts().SYCLIsDevice &&
55395539
Function->hasAttr<SYCLKernelAttr>()) {
5540-
ConstructSYCLKernel(Function);
5540+
ConstructOpenCLKernel(Function);
55415541
}
55425542
Function->setInstantiationIsPending(false);
55435543
}

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

+1-15
Original file line numberDiff line numberDiff line change
@@ -26,9 +26,6 @@ int main() {
2626
// Check lambda object alloca
2727
// CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon"
2828
// Check allocas for ranges
29-
// CHECK: [[ACC_RANGE_COPY:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
30-
// CHECK: [[MEM_RANGE_COPY:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
31-
// CHECK: [[OFFSET_COPY:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
3229
//
3330
// Check store of kernel pointer argument to alloca
3431
// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8
@@ -39,19 +36,8 @@ int main() {
3936
// Check load from kernel pointer argument alloca
4037
// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr
4138

42-
// Check that ranges/offsets arguments are copied to allocas
43-
// CHECK: [[BITCAST1:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE_COPY]] to i8*
44-
// CHECK: [[BITCAST2:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE]] to i8*
45-
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[BITCAST1]], i8* align 4 [[BITCAST2]], i64 4, i1 false)
46-
// CHECK: [[BITCAST3:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE_COPY]] to i8*
47-
// CHECK: [[BITCAST4:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE]] to i8*
48-
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[BITCAST3]], i8* align 4 [[BITCAST4]], i64 4, i1 false)
49-
// CHECK: [[BITCAST5:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET_COPY]] to i8*
50-
// CHECK: [[BITCAST6:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET]] to i8*
51-
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[BITCAST5]], i8* align 4 [[BITCAST6]], i64 4, i1 false)
52-
5339
// Check accessor __init method call
54-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[ACC_RANGE_COPY]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[MEM_RANGE_COPY]], %"struct.{{.*}}.cl::sycl::id"* byval align 4 [[OFFSET_COPY]])
40+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[ACC_RANGE]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[MEM_RANGE]], %"struct.{{.*}}.cl::sycl::id"* byval align 4 [[OFFSET]])
5541

5642
// Check lambda "()" operator call
5743
// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]])

‎clang/test/CodeGenSYCL/sampler.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -3,11 +3,11 @@
33
// CHECK-NEXT: entry:
44
// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
55
// CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 8
6-
// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
6+
// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8
77
// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %"class.{{.*}}.anon"* [[ANON]] to i8*
88
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
99
// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0
10-
// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
10+
// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8
1111
// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler"* [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
1212
//
1313
#include "sycl.hpp"

‎clang/test/SemaSYCL/Inputs/sycl.hpp

+20-1
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,24 @@ struct _ImplT {
5656
id<dim> Offset;
5757
};
5858

59+
template <typename dataT, access::target accessTarget>
60+
struct DeviceValueType;
61+
62+
template <typename dataT>
63+
struct DeviceValueType<dataT, access::target::global_buffer> {
64+
using type = __global dataT;
65+
};
66+
67+
template <typename dataT>
68+
struct DeviceValueType<dataT, access::target::constant_buffer> {
69+
using type = __constant dataT;
70+
};
71+
72+
template <typename dataT>
73+
struct DeviceValueType<dataT, access::target::local> {
74+
using type = __local dataT;
75+
};
76+
5977
template <typename dataT, int dimensions, access::mode accessmode,
6078
access::target accessTarget = access::target::global_buffer,
6179
access::placeholder isPlaceholder = access::placeholder::false_t>
@@ -67,7 +85,8 @@ class accessor {
6785
_ImplT<dimensions> impl;
6886

6987
private:
70-
void __init(__global dataT *Ptr, range<dimensions> AccessRange,
88+
using PtrType = typename DeviceValueType<dataT, accessTarget>::type *;
89+
void __init(PtrType Ptr, range<dimensions> AccessRange,
7190
range<dimensions> MemRange, id<dimensions> Offset) {}
7291
};
7392

‎clang/test/SemaSYCL/accessors-targets.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,6 @@ int main() {
3636
constant_acc.use();
3737
});
3838
}
39-
// CHECK: {{.*}}use_local 'void (__local int *, range<1>, range<1>, id<1>)'
40-
// CHECK: {{.*}}use_global 'void (__global int *, range<1>, range<1>, id<1>)'
41-
// CHECK: {{.*}}use_constant 'void (__constant int *, range<1>, range<1>, id<1>)'
39+
// CHECK: {{.*}}use_local 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
40+
// CHECK: {{.*}}use_global 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
41+
// CHECK: {{.*}}use_constant 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'

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

+12-13
Original file line numberDiff line numberDiff line change
@@ -23,14 +23,14 @@ int main() {
2323

2424
// Check declaration of the kernel
2525

26-
// CHECK: FunctionDecl {{.*}}kernel_wrapper 'void (__global int *, range<1>, range<1>, id<1>)'
26+
// CHECK: FunctionDecl {{.*}}kernel_wrapper 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
2727

2828
// Check parameters of the kernel
2929

3030
// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *'
31-
// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>'
32-
// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>'
33-
// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'id<1>':'cl::sycl::id<1>'
31+
// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
32+
// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
33+
// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
3434

3535
// Check body of the kernel
3636

@@ -42,18 +42,17 @@ int main() {
4242
// Check accessor initialization
4343

4444
// CHECK: CXXMemberCallExpr {{.*}} 'void'
45-
// CHECK-NEXT: MemberExpr {{.*}} 'void (__global int *, range<1>, range<1>, id<1>)' lvalue .__init
45+
// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
4646
// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' lvalue .
4747
// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var
48-
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
49-
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' lvalue <NoOp>
48+
// CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
5049
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *'
51-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <NoOp>
52-
// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'range<1>':'cl::sycl::range<1>'
53-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <NoOp>
54-
// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'range<1>':'cl::sycl::range<1>'
55-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' <NoOp>
56-
// CHECK-NEXT: DeclRefExpr {{.*}} 'id<1>':'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'id<1>':'cl::sycl::id<1>'
50+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <LValueToRValue>
51+
// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'cl::sycl::range<1>'
52+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <LValueToRValue>
53+
// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'cl::sycl::range<1>'
54+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' <LValueToRValue>
55+
// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'cl::sycl::id<1>'
5756

5857
// Check that body of the kernel caller function is included into kernel
5958

‎clang/test/SemaSYCL/fake-accessors.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,6 @@ int main() {
5151
});
5252
return 0;
5353
}
54-
// CHECK: fake_accessors 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor)
55-
// CHECK: accessor_typedef 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor)
56-
// CHECK: accessor_alias 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor)
54+
// CHECK: fake_accessors 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
55+
// CHECK: accessor_typedef 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
56+
// CHECK: accessor_alias 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)

‎clang/test/SemaSYCL/sampler.cpp

+3-4
Original file line numberDiff line numberDiff line change
@@ -16,10 +16,10 @@ int main() {
1616
}
1717

1818
// Check declaration of the test kernel
19-
// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__ocl_sampler_t)'
19+
// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (sampler_t)'
2020
//
2121
// Check parameters of the test kernel
22-
// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] '__ocl_sampler_t'
22+
// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] 'sampler_t'
2323
//
2424
// Check that sampler field of the test kernel object is initialized using __init method
2525
// CHECK: CXXMemberCallExpr {{.*}} 'void'
@@ -29,5 +29,4 @@ int main() {
2929
//
3030
// Check the parameters of __init method
3131
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' <LValueToRValue>
32-
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue <NoOp>
33-
// CHECK-NEXT: DeclRefExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__ocl_sampler_t':'sampler_t'
32+
// CHECK-NEXT: DeclRefExpr {{.*}} 'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' 'sampler_t'

‎clang/test/SemaSYCL/wrapped-accessor.cpp

+12-13
Original file line numberDiff line numberDiff line change
@@ -23,14 +23,14 @@ int main() {
2323
}
2424

2525
// Check declaration of the kernel
26-
// CHECK: wrapped_access 'void (AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >, __global int *, range<1>, range<1>, id<1>)'
26+
// CHECK: wrapped_access 'void (AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
2727

2828
// Check parameters of the kernel
2929
// CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >'
3030
// CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *'
31-
// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>'
32-
// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>'
33-
// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'id<1>':'cl::sycl::id<1>'
31+
// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
32+
// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
33+
// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
3434

3535
// Check that wrapper object itself is initialized with corresponding kernel argument using operator=
3636
// CHECK: BinaryOperator {{.*}} 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >' lvalue '='
@@ -45,18 +45,17 @@ int main() {
4545

4646
// Check that accessor field of the wrapper object is initialized using __init method
4747
// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
48-
// CHECK-NEXT: MemberExpr {{.*}} 'void (__global int *, range<1>, range<1>, id<1>)' lvalue .__init
48+
// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
4949
// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' lvalue .accessor {{.*}}
5050
// CHECK-NEXT: MemberExpr {{.*}} 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >' lvalue .
5151
// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
5252

5353
// Parameters of the _init method
54-
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
55-
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' lvalue <NoOp>
54+
// CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
5655
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global int *'
57-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <NoOp>
58-
// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'range<1>':'cl::sycl::range<1>'
59-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <NoOp>
60-
// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'range<1>':'cl::sycl::range<1>'
61-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' <NoOp>
62-
// CHECK-NEXT: DeclRefExpr {{.*}} 'id<1>':'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'id<1>':'cl::sycl::id<1>'
56+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <LValueToRValue>
57+
// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'cl::sycl::range<1>'
58+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <LValueToRValue>
59+
// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'cl::sycl::range<1>'
60+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' <LValueToRValue>
61+
// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'cl::sycl::id<1>'

0 commit comments

Comments
 (0)