Skip to content

Commit

Permalink
[SYCL][Test] Generate kernel-param test using script (#22546)
Browse files Browse the repository at this point in the history 
Will also be in sycl branch through pulldown in
#14327

---------

Co-authored-by: sys_ce_bb <sys_ce_bb@intel.com>
Co-authored-by: premanandrao <premanand.m.rao@intel.com>
  • Loading branch information
@jsji @sys-ce-bb @premanandrao
3 people committed Jun 28, 2024
1 parent 0378786 commit 5465207
Show file tree
Hide file tree
Showing 2 changed files with 117 additions and 132 deletions.
123 changes: 57 additions & 66 deletions clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
View file
Original file line number Diff line number Diff line change
@@ -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

Expand All @@ -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
Loading

0 comments on commit 5465207

Please sign in to comment.