diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp index e2ed6126f9bab..841080cb22661 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp @@ -1,4 +1,6 @@ // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 5 +// And edited to fix some failures. // This test checks a kernel argument that is an Accessor array @@ -22,69 +24,58 @@ int main() { acc[1].use(); }); } - -// Check kernel_A parameters -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A -// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], -// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], -// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], -// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]], -// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]], -// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]], -// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]], -// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]]) - -// CHECK alloca for pointer arguments -// CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca ptr addrspace(1), align 8 -// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca ptr addrspace(1), align 8 - -// CHECK lambda object alloca -// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4 - -// CHECK allocas for ranges -// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id" -// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id" - -// CHECK lambda object addrspacecast -// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr [[LOCAL_OBJECTA]] to ptr addrspace(4) - -// CHECK addrspacecasts for ranges -// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast ptr [[ACC_RANGE1A]] to ptr addrspace(4) -// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast ptr [[MEM_RANGE1A]] to ptr addrspace(4) -// CHECK: [[OFFSET1AS:%.*]] = addrspacecast ptr [[OFFSET1A]] to ptr addrspace(4) -// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast ptr [[ACC_RANGE2A]] to ptr addrspace(4) -// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast ptr [[MEM_RANGE2A]] to ptr addrspace(4) -// CHECK: [[OFFSET2AS:%.*]] = addrspacecast ptr [[OFFSET2A]] to ptr addrspace(4) -// CHECK accessor array default inits -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], ptr addrspace(4) [[ACCESSOR_ARRAY1]], i64 0, i64 0 -// Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP. -// CTOR Call #1 -// CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[BEGIN]]) -// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], ptr addrspace(4) [[BEGIN]], i64 1 -// CTOR Call #2 -// CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[ELEM2_GEP]]) - -// CHECK acc[0] __init method call -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[ACCESSOR_ARRAY1]], i64 0, i64 0 -// CHECK load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG1]] -// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast ptr addrspace(4) [[ACC_RANGE1AS]] to ptr -// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast ptr addrspace(4) [[MEM_RANGE1AS]] to ptr -// CHECK: [[OFFSET1:%.*]] = addrspacecast ptr addrspace(4) [[OFFSET1AS]] to ptr -// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[INDEX1]], ptr addrspace(1) noundef [[MEM_LOAD1]], ptr noundef byval({{.*}}) align 4 [[ACC_RANGE1]], ptr noundef byval({{.*}}) align 4 [[MEM_RANGE1]], ptr noundef byval({{.*}}) align 4 [[OFFSET1]]) - -// CHECK acc[1] __init method call -// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[ACCESSOR_ARRAY2]], i64 0, i64 1 -// CHECK load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG2]] -// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast ptr addrspace(4) [[ACC_RANGE2AS]] to ptr -// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast ptr addrspace(4) [[MEM_RANGE2AS]] to ptr -// CHECK: [[OFFSET2:%.*]] = addrspacecast ptr addrspace(4) [[OFFSET2AS]] to ptr -// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[INDEX2]], ptr addrspace(1) noundef [[MEM_LOAD2]], ptr noundef byval({{.*}}) align 4 [[ACC_RANGE2]], ptr noundef byval({{.*}}) align 4 [[MEM_RANGE2]], ptr noundef byval({{.*}}) align 4 [[OFFSET2]]) +// CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_A( +// CHECK-SAME: ptr addrspace(1) noundef align 4 [[_ARG_ACC:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC1:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC2:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC3:%.*]], ptr addrspace(1) noundef align 4 [[_ARG_ACC4:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC6:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC7:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC8:%.*]]) #[[ATTR0:[0-9]+]] +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[_ARG_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8 +// CHECK-NEXT: [[_ARG_ACC_ADDR5:%.*]] = alloca ptr addrspace(1), align 8 +// CHECK-NEXT: [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON:%.*]], align 4 +// CHECK-NEXT: [[AGG_TMP:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP10:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP11:%.*]] = alloca %"struct.sycl::_V1::id", align 4 +// CHECK-NEXT: [[AGG_TMP14:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP15:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP16:%.*]] = alloca %"struct.sycl::_V1::id", align 4 +// CHECK-NEXT: [[_ARG_ACC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC_ADDR]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_ACC_ADDR5_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC_ADDR5]] to ptr addrspace(4) +// CHECK-NEXT: [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP10_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP10]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP11_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP11]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP14_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP14]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP15_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP15]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP16_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP16]] to ptr addrspace(4) +// CHECK-NEXT: store ptr addrspace(1) [[_ARG_ACC]], ptr addrspace(4) [[_ARG_ACC_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[_ARG_ACC1_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC1]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_ACC2_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC2]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_ACC3_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC3]] to ptr addrspace(4) +// CHECK-NEXT: store ptr addrspace(1) [[_ARG_ACC4]], ptr addrspace(4) [[_ARG_ACC_ADDR5_ASCAST]], align 8 +// CHECK-NEXT: [[_ARG_ACC6_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC6]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_ACC7_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC7]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_ACC8_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC8]] to ptr addrspace(4) +// CHECK-NEXT: [[ACC:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ACC]]) #[[ATTR4:[0-9]+]] +// CHECK-NEXT: [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds %"class.sycl::_V1::accessor", ptr addrspace(4) [[ACC]], i64 1 +// CHECK-NEXT: call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYINIT_ELEMENT]]) #[[ATTR4]] +// CHECK-NEXT: [[ACC9:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[ACC9]], i64 0, i64 0 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_ACC_ADDR_ASCAST]], align 8 +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC1_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP10_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC2_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP11_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC3_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: [[AGG_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP10_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP10_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP11_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP11_ASCAST]] to ptr +// CHECK-NEXT: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX]], ptr addrspace(1) noundef [[TMP0]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP10_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP11_ASCAST_ASCAST]]) #[[ATTR4]] +// CHECK-NEXT: [[ACC12:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[ACC12]], i64 0, i64 1 +// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_ACC_ADDR5_ASCAST]], align 8 +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP14_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC6_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP15_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC7_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP16_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC8_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: [[AGG_TMP14_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP14_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP15_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP15_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP16_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP16_ASCAST]] to ptr +// CHECK-NEXT: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX13]], ptr addrspace(1) noundef [[TMP1]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP14_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP15_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP16_ASCAST_ASCAST]]) #[[ATTR4]] +// CHECK-NEXT: call spir_func void @_ZZ4mainENKUlvE_clEv(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) [[__SYCLKERNEL_ASCAST]]) #[[ATTR4]] +// CHECK-NEXT: ret void diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp index b569e152899f6..41d18645ff3eb 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp @@ -1,4 +1,6 @@ // RUN: %clang_cc1 -fsycl-is-device -fsycl-int-header=%t.h -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 5 +// And edited later to fix failures. // This test checks a kernel with struct parameter that contains an Accessor array. @@ -26,69 +28,61 @@ int main() { }); } -// CHECK kernel_C parameters -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_C -// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]], -// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]], -// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]], -// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]], -// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]], -// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]], -// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]], -// CHECK-SAME: ptr noundef byval({{.*}}) align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]]) - -// Check alloca for pointer arguments -// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca ptr addrspace(1), align 8 -// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca ptr addrspace(1), align 8 - -// Check lambda object alloca -// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class{{.*}}.anon, align 4 - -// Check allocas for ranges -// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id" -// CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range" -// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id" - -// Check lambda object addrspacecast -// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr %__SYCLKernel to ptr addrspace(4) - -// Check addrspacecast for ranges -// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast ptr [[ACC_RANGE1A]] to ptr addrspace(4) -// CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast ptr [[MEM_RANGE1A]] to ptr addrspace(4) -// CHECK: [[OFFSET1AS:%.*]] = addrspacecast ptr [[OFFSET1A]] to ptr addrspace(4) -// CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast ptr [[ACC_RANGE2A]] to ptr addrspace(4) -// CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast ptr [[MEM_RANGE2A]] to ptr addrspace(4) -// CHECK: [[OFFSET2AS:%.*]] = addrspacecast ptr [[OFFSET2A]] to ptr addrspace(4) - -// CHECK accessor array default inits -// CHECK: [[ACCESSOR_WRAPPER:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, ptr addrspace(4) [[ACCESSOR_WRAPPER]], i32 0, i32 0 -// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], ptr addrspace(4) [[ACCESSOR_ARRAY1]], i64 0, i64 0 -// CTOR Call #1 -// CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[BEGIN]]) -// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], ptr addrspace(4) [[BEGIN]], i64 1 -// CTOR Call #2 -// CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[ELEM2_GEP]]) - -// Check acc[0] __init method call -// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, ptr addrspace(4) [[GEP_LAMBDA1]], i32 0, i32 0 -// CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[GEP_MEMBER_ACC1]], i64 0, i64 0 -// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG1]].addr -// CHECK: [[ACC_RANGE1:%.*]] = addrspacecast ptr addrspace(4) [[ACC_RANGE1AS]] to ptr -// CHECK: [[MEM_RANGE1:%.*]] = addrspacecast ptr addrspace(4) [[MEM_RANGE1AS]] to ptr -// CHECK: [[OFFSET1:%.*]] = addrspacecast ptr addrspace(4) [[OFFSET1AS]] to ptr -// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[ARRAY_IDX1]], ptr addrspace(1) noundef [[MEM_LOAD1]], ptr noundef byval({{.*}}) align 4 [[ACC_RANGE1]], ptr noundef byval({{.*}}) align 4 [[MEM_RANGE1]], ptr noundef byval({{.*}}) align 4 [[OFFSET1]]) - -// Check acc[1] __init method call -// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, ptr addrspace(4) [[GEP_LAMBDA2]], i32 0, i32 0 -// CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[GEP_MEMBER_ACC2]], i64 0, i64 1 -// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG1]].addr -// CHECK: [[ACC_RANGE2:%.*]] = addrspacecast ptr addrspace(4) [[ACC_RANGE2AS]] to ptr -// CHECK: [[MEM_RANGE2:%.*]] = addrspacecast ptr addrspace(4) [[MEM_RANGE2AS]] to ptr -// CHECK: [[OFFSET2:%.*]] = addrspacecast ptr addrspace(4) [[OFFSET2AS]] to ptr -// CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[ARRAY_IDX2]], ptr addrspace(1) noundef [[MEM_LOAD2]], ptr noundef byval({{.*}}) align 4 [[ACC_RANGE2]], ptr noundef byval({{.*}}) align 4 [[MEM_RANGE2]], ptr noundef byval({{.*}}) align 4 [[OFFSET2]]) +// CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_C( +// CHECK-SAME: ptr addrspace(1) noundef align 4 [[_ARG_MEMBER_ACC:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC1:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC2:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_MEMBER_ACC3:%.*]], ptr addrspace(1) noundef align 4 [[_ARG_MEMBER_ACC4:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC6:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_MEMBER_ACC7:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_MEMBER_ACC8:%.*]]) #[[ATTR0:[0-9]+]] +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[_ARG_MEMBER_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8 +// CHECK-NEXT: [[_ARG_MEMBER_ACC_ADDR5:%.*]] = alloca ptr addrspace(1), align 8 +// CHECK-NEXT: [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON:%.*]], align 4 +// CHECK-NEXT: [[AGG_TMP:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP11:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP12:%.*]] = alloca %"struct.sycl::_V1::id", align 4 +// CHECK-NEXT: [[AGG_TMP16:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP17:%.*]] = alloca %"struct.sycl::_V1::range", align 4 +// CHECK-NEXT: [[AGG_TMP18:%.*]] = alloca %"struct.sycl::_V1::id", align 4 +// CHECK-NEXT: [[_ARG_MEMBER_ACC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC_ADDR]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_MEMBER_ACC_ADDR5_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC_ADDR5]] to ptr addrspace(4) +// CHECK-NEXT: [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP11_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP11]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP12_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP12]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP16_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP16]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP17_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP17]] to ptr addrspace(4) +// CHECK-NEXT: [[AGG_TMP18_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP18]] to ptr addrspace(4) +// CHECK-NEXT: store ptr addrspace(1) [[_ARG_MEMBER_ACC]], ptr addrspace(4) [[_ARG_MEMBER_ACC_ADDR_ASCAST]], align 8 +// CHECK-NEXT: [[_ARG_MEMBER_ACC1_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC1]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_MEMBER_ACC2_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC2]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_MEMBER_ACC3_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC3]] to ptr addrspace(4) +// CHECK-NEXT: store ptr addrspace(1) [[_ARG_MEMBER_ACC4]], ptr addrspace(4) [[_ARG_MEMBER_ACC_ADDR5_ASCAST]], align 8 +// CHECK-NEXT: [[_ARG_MEMBER_ACC6_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC6]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_MEMBER_ACC7_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC7]] to ptr addrspace(4) +// CHECK-NEXT: [[_ARG_MEMBER_ACC8_ASCAST:%.*]] = addrspacecast ptr [[_ARG_MEMBER_ACC8]] to ptr addrspace(4) +// CHECK-NEXT: [[STRUCT_ACC:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[MEMBER_ACC:%.*]] = getelementptr inbounds [[STRUCT_STRUCT_ACC_T:%.*]], ptr addrspace(4) [[STRUCT_ACC]], i32 0, i32 0 +// CHECK-NEXT: call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[MEMBER_ACC]]) #[[ATTR4:[0-9]+]] +// CHECK-NEXT: [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds %"class.sycl::_V1::accessor", ptr addrspace(4) [[MEMBER_ACC]], i64 1 +// CHECK-NEXT: call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYINIT_ELEMENT]]) #[[ATTR4]] +// CHECK-NEXT: [[STRUCT_ACC9:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[MEMBER_ACC10:%.*]] = getelementptr inbounds [[STRUCT_STRUCT_ACC_T]], ptr addrspace(4) [[STRUCT_ACC9]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[MEMBER_ACC10]], i64 0, i64 0 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_MEMBER_ACC_ADDR_ASCAST]], align 8 +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC1_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP11_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC2_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP12_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC3_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: [[AGG_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP11_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP11_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP12_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP12_ASCAST]] to ptr +// CHECK-NEXT: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX]], ptr addrspace(1) noundef [[TMP0]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP11_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP12_ASCAST_ASCAST]]) #[[ATTR4]] +// CHECK-NEXT: [[STRUCT_ACC13:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0 +// CHECK-NEXT: [[MEMBER_ACC14:%.*]] = getelementptr inbounds [[STRUCT_STRUCT_ACC_T]], ptr addrspace(4) [[STRUCT_ACC13]], i32 0, i32 0 +// CHECK-NEXT: [[ARRAYIDX15:%.*]] = getelementptr inbounds [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[MEMBER_ACC14]], i64 0, i64 1 +// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_MEMBER_ACC_ADDR5_ASCAST]], align 8 +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP16_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC6_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP17_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC7_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP18_ASCAST]], ptr addrspace(4) align 4 [[_ARG_MEMBER_ACC8_ASCAST]], i64 4, i1 false) +// CHECK-NEXT: [[AGG_TMP16_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP16_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP17_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP17_ASCAST]] to ptr +// CHECK-NEXT: [[AGG_TMP18_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP18_ASCAST]] to ptr +// CHECK-NEXT: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX15]], ptr addrspace(1) noundef [[TMP1]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP16_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP17_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP18_ASCAST_ASCAST]]) #[[ATTR4]] +// CHECK-NEXT: call spir_func void @_ZZ4mainENKUlvE_clEv(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) [[__SYCLKERNEL_ASCAST]]) #[[ATTR4]] +// CHECK-NEXT: ret void