Skip to content

Commit b23734f

Browse files
lbushi25lbushi25
committed
Add tests for free function kernel special type arguments
1 parent b4f18e1 commit b23734f

File tree

1 file changed

+144
-1
lines changed

1 file changed

+144
-1
lines changed

‎clang/test/CodeGenSYCL/free_function_kernel_params.cpp‎

Lines changed: 144 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -83,4 +83,147 @@ void ff_7(sycl::dynamic_work_group_memory<int> DynMem) {
8383
// CHECK-NEXT: %DynMem.ascast = addrspacecast ptr %DynMem to ptr addrspace(4)
8484
// CHECK: store ptr addrspace(3) %__arg_Ptr, ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
8585
// CHECK-NEXT: [[REGISTER:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
86-
// CHECK-NEXT: call spir_func void @{{.*}}dynamic_work_group_memory{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %DynMem.ascast, ptr addrspace(3) noundef [[REGISTER]])
86+
// CHECK-NEXT: call spir_func void @{{.*}}dynamic_work_group_memory{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %DynMem.ascast, ptr addrspace(3) noundef [[REGISTER]])
87+
88+
__attribute__((sycl_device))
89+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
90+
void ff_8(sycl::accessor<int, 1, sycl::access::mode::read> acc) {
91+
}
92+
93+
// CHECK: define dso_local spir_kernel void @_Z18__sycl_kernel_ff_8N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE(ptr addrspace(1) noundef align 4 %__arg_Ptr, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %__arg_AccessRange, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %__arg_MemRange, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %__arg_Offset) #5 !srcloc !9 !kernel_arg_buffer_location !23 !sycl_fixed_targets !6 {
94+
// CHECK-NEXT: %__arg_Ptr.addr = alloca ptr addrspace(1), align 8
95+
// CHECK-NEXT: %acc = alloca %"class.sycl::_V1::accessor", align 4
96+
// CHECK-NEXT: %agg.tmp = alloca %"struct.sycl::_V1::range", align 4
97+
// CHECK-NEXT: %agg.tmp1 = alloca %"struct.sycl::_V1::range", align 4
98+
// CHECK-NEXT: %agg.tmp2 = alloca %"struct.sycl::_V1::id", align 4
99+
// CHECK-NEXT: %agg.tmp3 = alloca %"class.sycl::_V1::accessor", align 4
100+
// CHECK-NEXT: %__arg_Ptr.addr.ascast = addrspacecast ptr %__arg_Ptr.addr to ptr addrspace(4)
101+
// CHECK-NEXT: %acc.ascast = addrspacecast ptr %acc to ptr addrspace(4)
102+
// CHECK-NEXT: %agg.tmp.ascast = addrspacecast ptr %agg.tmp to ptr addrspace(4)
103+
// CHECK-NEXT: %agg.tmp1.ascast = addrspacecast ptr %agg.tmp1 to ptr addrspace(4)
104+
// CHECK-NEXT: %agg.tmp2.ascast = addrspacecast ptr %agg.tmp2 to ptr addrspace(4)
105+
// CHECK-NEXT: %agg.tmp3.ascast = addrspacecast ptr %agg.tmp3 to ptr addrspace(4)
106+
// CHECK-NEXT: store ptr addrspace(1) %__arg_Ptr, ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
107+
// CHECK-NEXT: %__arg_AccessRange.ascast = addrspacecast ptr %__arg_AccessRange to ptr addrspace(4)
108+
// CHECK-NEXT: %__arg_MemRange.ascast = addrspacecast ptr %__arg_MemRange to ptr addrspace(4)
109+
// CHECK-NEXT: %__arg_Offset.ascast = addrspacecast ptr %__arg_Offset to ptr addrspace(4)
110+
111+
112+
__attribute__((sycl_device))
113+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
114+
void ff_8(sycl::local_accessor<int, 1> lacc) {
115+
}
116+
117+
// CHECK : define dso_local spir_kernel void @_Z18__sycl_kernel_ff_8N4sycl3_V114local_accessorIiLi1EEE(ptr addrspace(3) noundef align 4 %__arg_Ptr, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %__arg_AccessRange, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %__arg_MemRange, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %__arg_Offset) #5 !srcloc !9 !kernel_arg_buffer_location !22 !sycl_fixed_targets !6 {
118+
// CHECK: %__arg_Ptr.addr = alloca ptr addrspace(3), align 8
119+
// CHECK-NEXT: %lacc = alloca %"class.sycl::_V1::local_accessor", align 4
120+
// CHECK-NEXT: %agg.tmp = alloca %"struct.sycl::_V1::range", align 4
121+
// CHECK-NEXT: %agg.tmp1 = alloca %"struct.sycl::_V1::range", align 4
122+
// CHECK-NEXT: %agg.tmp2 = alloca %"struct.sycl::_V1::id", align 4
123+
// CHECK-NEXT: %agg.tmp3 = alloca %"class.sycl::_V1::local_accessor", align 4
124+
// CHECK-NEXT: %__arg_Ptr.addr.ascast = addrspacecast ptr %__arg_Ptr.addr to ptr addrspace(4)
125+
// CHECK-NEXT: %lacc.ascast = addrspacecast ptr %lacc to ptr addrspace(4)
126+
// CHECK-NEXT: %agg.tmp.ascast = addrspacecast ptr %agg.tmp to ptr addrspace(4)
127+
// CHECK-NEXT: %agg.tmp1.ascast = addrspacecast ptr %agg.tmp1 to ptr addrspace(4)
128+
// CHECK-NEXT: %agg.tmp2.ascast = addrspacecast ptr %agg.tmp2 to ptr addrspace(4)
129+
// CHECK-NEXT: %agg.tmp3.ascast = addrspacecast ptr %agg.tmp3 to ptr addrspace(4)
130+
// CHECK-NEXT: store ptr addrspace(3) %__arg_Ptr, ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
131+
// CHECK-NEXT: %__arg_AccessRange.ascast = addrspacecast ptr %__arg_AccessRange to ptr addrspace(4)
132+
// CHECK-NEXT: %__arg_MemRange.ascast = addrspacecast ptr %__arg_MemRange to ptr addrspace(4)
133+
// CHECK-NEXT: %__arg_Offset.ascast = addrspacecast ptr %__arg_Offset to ptr addrspace(4)
134+
// CHECK-NEXT: call spir_func void @_ZN4sycl3_V114local_accessorIiLi1EEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) %lacc.ascast) #8
135+
// CHECK: call spir_func void @_ZN4sycl3_V114local_accessorIiLi1EE6__initEPU3AS3iNS0_5rangeILi1EEES6_NS0_2idILi1EEE(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) %lacc.ascast, ptr addrspace(3) noundef %0, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.tmp.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.tmp1.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %agg.tmp2.ascast.ascast) #8
136+
// CHECK: call spir_func void @_Z4ff_8N4sycl3_V114local_accessorIiLi1EEE(ptr noundef byval(%"class.sycl::_V1::local_accessor") align 4 %agg.tmp3.ascast.ascast) #8
137+
138+
// CHECK: declare spir_func void @_ZN4sycl3_V114local_accessorIiLi1EEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24)) unnamed_addr #7
139+
140+
// CHECK: define linkonce_odr spir_func void @_ZN4sycl3_V114local_accessorIiLi1EE6__initEPU3AS3iNS0_5rangeILi1EEES6_NS0_2idILi1EEE(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) %this, ptr addrspace(3) noundef %Ptr, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %AccessRange, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %MemRange, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %Offset) #6 comdat align 2 !srcloc !23 {
141+
142+
__attribute__((sycl_device))
143+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
144+
void ff_8(sycl::sampler S) {
145+
}
146+
147+
// CHECK: define dso_local spir_kernel void @_Z18__sycl_kernel_ff_8N4sycl3_V17samplerE(target("spirv.Sampler") %__arg_Sampler) #5 !srcloc !10 !kernel_arg_buffer_location !16 !sycl_fixed_targets !6 {
148+
// CHECK-NEXT: %__arg_Sampler.addr = alloca target("spirv.Sampler"), align 8
149+
// CHECK-NEXT: %S = alloca %"class.sycl::_V1::sampler", align 8
150+
// CHECK-NEXT: %agg.tmp = alloca %"class.sycl::_V1::sampler", align 8
151+
// CHECK-NEXT: %__arg_Sampler.addr.ascast = addrspacecast ptr %__arg_Sampler.addr to ptr addrspace(4)
152+
// CHECK-NEXT: %S.ascast = addrspacecast ptr %S to ptr addrspace(4)
153+
// CHECK-NEXT: %agg.tmp.ascast = addrspacecast ptr %agg.tmp to ptr addrspace(4)
154+
// CHECK-NEXT: store target("spirv.Sampler") %__arg_Sampler, ptr addrspace(4) %__arg_Sampler.addr.ascast, align 8
155+
// CHECK-NEXT: %0 = load target("spirv.Sampler"), ptr addrspace(4) %__arg_Sampler.addr.ascast, align 8
156+
// CHECK-NEXT: call spir_func void @_ZN4sycl3_V17sampler6__initE11ocl_sampler(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %S.ascast, target("spirv.Sampler") %0) #8
157+
// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %agg.tmp.ascast, ptr addrspace(4) align 8 %S.ascast, i64 8, i1 false)
158+
// CHECK-NEXT: %agg.tmp.ascast.ascast = addrspacecast ptr addrspace(4) %agg.tmp.ascast to ptr
159+
// CHECK-NEXT: call spir_func void @_Z4ff_8N4sycl3_V17samplerE(ptr noundef byval(%"class.sycl::_V1::sampler") align 8 %agg.tmp.ascast.ascast) #8
160+
161+
// CHECK: define linkonce_odr spir_func void @_ZN4sycl3_V17sampler6__initE11ocl_sampler(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %this, target("spirv.Sampler") %Sampler) #6 comdat align 2 !srcloc !24 {
162+
163+
__attribute__((sycl_device))
164+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
165+
void ff_8(sycl::stream str) {
166+
}
167+
168+
// CHECK: define dso_local spir_kernel void @_Z18__sycl_kernel_ff_8N4sycl3_V16streamE(ptr addrspace(1) noundef align 1 %__arg_Ptr, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %__arg_AccessRange, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %__arg_MemRange, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %__arg_Offset, i32 noundef %__arg__FlushBufferSize) #5 !srcloc !11 !kernel_arg_buffer_location !25 !sycl_fixed_targets !6 {
169+
// CHECK-NEXT: %__arg_Ptr.addr = alloca ptr addrspace(1), align 8
170+
// CHECK-NEXT: %__arg__FlushBufferSize.addr = alloca i32, align 4
171+
// CHECK-NEXT: %str = alloca %"class.sycl::_V1::stream", align 4
172+
// CHECK-NEXT: %agg.tmp = alloca %"struct.sycl::_V1::range", align 4
173+
// CHECK-NEXT: %agg.tmp1 = alloca %"struct.sycl::_V1::range", align 4
174+
// CHECK-NEXT: %agg.tmp2 = alloca %"struct.sycl::_V1::id", align 4
175+
// CHECK-NEXT: %agg.tmp3 = alloca %"class.sycl::_V1::stream", align 4
176+
// CHECK-NEXT: %__arg_Ptr.addr.ascast = addrspacecast ptr %__arg_Ptr.addr to ptr addrspace(4)
177+
// CHECK-NEXT: %__arg__FlushBufferSize.addr.ascast = addrspacecast ptr %__arg__FlushBufferSize.addr to ptr addrspace(4)
178+
// CHECK-NEXT: %str.ascast = addrspacecast ptr %str to ptr addrspace(4)
179+
// CHECK-NEXT: %agg.tmp.ascast = addrspacecast ptr %agg.tmp to ptr addrspace(4)
180+
// CHECK-NEXT: %agg.tmp1.ascast = addrspacecast ptr %agg.tmp1 to ptr addrspace(4)
181+
// CHECK-NEXT: %agg.tmp2.ascast = addrspacecast ptr %agg.tmp2 to ptr addrspace(4)
182+
// CHECK-NEXT: %agg.tmp3.ascast = addrspacecast ptr %agg.tmp3 to ptr addrspace(4)
183+
// CHECK-NEXT: store ptr addrspace(1) %__arg_Ptr, ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
184+
// CHECK-NEXT: %__arg_AccessRange.ascast = addrspacecast ptr %__arg_AccessRange to ptr addrspace(4)
185+
// CHECK-NEXT: %__arg_MemRange.ascast = addrspacecast ptr %__arg_MemRange to ptr addrspace(4)
186+
// CHECK-NEXT: %__arg_Offset.ascast = addrspacecast ptr %__arg_Offset to ptr addrspace(4)
187+
// CHECK-NEXT: store i32 %__arg__FlushBufferSize, ptr addrspace(4) %__arg__FlushBufferSize.addr.ascast, align 4
188+
// CHECK-NEXT: call spir_func void @_ZN4sycl3_V16streamC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(16) %str.ascast) #8
189+
// CHECK-NEXT: %0 = load ptr addrspace(1), ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
190+
// CHECK: call spir_func void @_ZN4sycl3_V16stream6__initEPU3AS1cNS0_5rangeILi1EEES5_NS0_2idILi1EEEi(ptr addrspace(4) noundef align 4 dereferenceable_or_null(16) %str.ascast, ptr addrspace(1) noundef %0, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.tmp.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.tmp1.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %agg.tmp2.ascast.ascast, i32 noundef %1) #8
191+
// CHECK:define linkonce_odr spir_func void @_ZN4sycl3_V16streamC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(16) %this) unnamed_addr #6 comdat align 2 !srcloc !26 {
192+
// CHECK: define linkonce_odr spir_func void @_ZN4sycl3_V16stream6__initEPU3AS1cNS0_5rangeILi1EEES5_NS0_2idILi1EEEi(ptr addrspace(4) noundef align 4 dereferenceable_or_null(16) %this, ptr addrspace(1) noundef %Ptr, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %AccessRange, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %MemRange, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %Offset, i32 noundef %_FlushBufferSize) #6 comdat align 2 !srcloc !27 {
193+
194+
__attribute__((sycl_device))
195+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
196+
void ff_8(sycl::ext::oneapi::experimental::annotated_arg<int> arg) {
197+
}
198+
199+
// CHECK: define dso_local spir_kernel void @_Z18__sycl_kernel_ff_8N4sycl3_V13ext6oneapi12experimental13annotated_argIiJEEE(i32 noundef %__arg__obj) #5 !srcloc !12 !kernel_arg_buffer_location !16 !sycl_fixed_targets !6 {
200+
// CHECK-NEXT: %__arg__obj.addr = alloca i32, align 4
201+
// CHECK-NEXT: %arg = alloca %"class.sycl::_V1::ext::oneapi::experimental::annotated_arg", align 4
202+
// CHECK-NEXT: %agg.tmp = alloca %"class.sycl::_V1::ext::oneapi::experimental::annotated_arg", align 4
203+
// CHECK-NEXT: %__arg__obj.addr.ascast = addrspacecast ptr %__arg__obj.addr to ptr addrspace(4)
204+
// CHECK-NEXT: %arg.ascast = addrspacecast ptr %arg to ptr addrspace(4)
205+
// CHECK-NEXT: %agg.tmp.ascast = addrspacecast ptr %agg.tmp to ptr addrspace(4)
206+
// CHECK-NEXT: store i32 %__arg__obj, ptr addrspace(4) %__arg__obj.addr.ascast, align 4
207+
// CHECK-NEXT: %0 = load i32, ptr addrspace(4) %__arg__obj.addr.ascast, align 4
208+
// CHECK-NEXT: call spir_func void @_ZN4sycl3_V13ext6oneapi12experimental13annotated_argIiJEE6__initEi(ptr addrspace(4) noundef align 4 dereferenceable_or_null(4) %arg.ascast, i32 noundef %0) #8
209+
// CHECK: define linkonce_odr spir_func void @_ZN4sycl3_V13ext6oneapi12experimental13annotated_argIiJEE6__initEi(ptr addrspace(4) noundef align 4 dereferenceable_or_null(4) %this, i32 noundef %_obj) #6 comdat align 2 !srcloc !33 {
210+
211+
__attribute__((sycl_device))
212+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
213+
void ff_8(sycl::ext::oneapi::experimental::annotated_ptr<int> ptr) {
214+
}
215+
216+
// CHECK: define dso_local spir_kernel void @_Z18__sycl_kernel_ff_8N4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEEE(ptr addrspace(4) noundef align 4 %__arg__obj) #5 !srcloc !13 !kernel_arg_buffer_location !16 !sycl_fixed_targets !6 {
217+
// CHECK-NEXT: %__arg__obj.addr = alloca ptr addrspace(4), align 8
218+
// CHECK-NEXT: %ptr = alloca %"class.sycl::_V1::ext::oneapi::experimental::annotated_ptr", align 8
219+
// CHECK-NEXT: %agg.tmp = alloca %"class.sycl::_V1::ext::oneapi::experimental::annotated_ptr", align 8
220+
// CHECK-NEXT: %__arg__obj.addr.ascast = addrspacecast ptr %__arg__obj.addr to ptr addrspace(4)
221+
// CHECK-NEXT: %ptr.ascast = addrspacecast ptr %ptr to ptr addrspace(4)
222+
// CHECK-NEXT: %agg.tmp.ascast = addrspacecast ptr %agg.tmp to ptr addrspace(4)
223+
// CHECK-NEXT: store ptr addrspace(4) %__arg__obj, ptr addrspace(4) %__arg__obj.addr.ascast, align 8
224+
// CHECK-NEXT: %0 = load ptr addrspace(4), ptr addrspace(4) %__arg__obj.addr.ascast, align 8
225+
// CHECK-NEXT: call spir_func void @_ZN4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEE6__initEPi(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %ptr.ascast, ptr addrspace(4) noundef %0) #8
226+
// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %agg.tmp.ascast, ptr addrspace(4) align 8 %ptr.ascast, i64 8, i1 false)
227+
// CHECK-NEXT: %agg.tmp.ascast.ascast = addrspacecast ptr addrspace(4) %agg.tmp.ascast to ptr
228+
// CHECK-NEXT: call spir_func void @_Z4ff_8N4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEEE(ptr noundef byval(%"class.sycl::_V1::ext::oneapi::experimental::annotated_ptr") align 8 %agg.tmp.ascast.ascast) #8
229+
// CHECK: define linkonce_odr spir_func void @_ZN4sycl3_V13ext6oneapi12experimental13annotated_ptrIiJEE6__initEPi(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %this, ptr addrspace(4) noundef %_obj) #6 comdat align 2 !srcloc !34 {

0 commit comments

Comments
 (0)