@@ -76,6 +76,26 @@ __attribute__((sycl_device))
7676
7777template void ff_6 (Agg S1, Derived S2, int );
7878
79+ constexpr int TestArrSize = 3 ;
80+
81+ template <int ArrSize>
82+ struct KArgWithPtrArray {
83+ int *data[ArrSize];
84+ int start[ArrSize];
85+ int end[ArrSize];
86+ constexpr int getArrSize () { return ArrSize; }
87+ };
88+
89+ template <int ArrSize>
90+ [[__sycl_detail__::add_ir_attributes_function(" sycl-single-task-kernel" , 0 )]]
91+ void ff_7 (KArgWithPtrArray<ArrSize> KArg) {
92+ for (int j = 0 ; j < ArrSize; j++)
93+ for (int i = KArg.start [j]; i <= KArg.end [j]; i++)
94+ KArg.data [j][i] = KArg.start [j] + KArg.end [j];
95+ }
96+
97+ template void ff_7 (KArgWithPtrArray<TestArrSize> KArg);
98+
7999// CHECK: const char* const kernel_names[] = {
80100// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piii
81101// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piiii
@@ -84,6 +104,7 @@ template void ff_6(Agg S1, Derived S2, int);
84104// CHECK-NEXT: {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
85105// CHECK-NEXT: {{.*}}__sycl_kernel_ff_410NoPointers8Pointers3Agg
86106// CHECK-NEXT: {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
107+ // CHECK-NEXT: {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
87108// CHECK-NEXT: ""
88109// CHECK-NEXT: };
89110
@@ -124,6 +145,9 @@ template void ff_6(Agg S1, Derived S2, int);
124145// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 40, 32 },
125146// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 72 },
126147
148+ // CHECK: //--- _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
149+ // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 48, 0 },
150+
127151// CHECK: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
128152// CHECK-NEXT: };
129153
@@ -249,6 +273,26 @@ template void ff_6(Agg S1, Derived S2, int);
249273// CHECK-NEXT: static constexpr bool value = true;
250274// CHECK-NEXT: };
251275// CHECK-NEXT: }
276+ //
277+ // CHECK: Definition of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE as a free function kernel
278+
279+ // CHECK: Forward declarations of kernel and its argument types:
280+ // CHECK: template <int ArrSize> struct KArgWithPtrArray;
281+ //
282+ // CHECK: template <int ArrSize> void ff_7(KArgWithPtrArray<ArrSize> KArg);
283+ // CHECK-NEXT: static constexpr auto __sycl_shim8() {
284+ // CHECK-NEXT: return (void (*)(struct KArgWithPtrArray<3>))ff_7<3>;
285+ // CHECK-NEXT: }
286+ // CHECK-NEXT: namespace sycl {
287+ // CHECK-NEXT: template <>
288+ // CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim8()> {
289+ // CHECK-NEXT: static constexpr bool value = true;
290+ // CHECK-NEXT: };
291+ // CHECK-NEXT: template <>
292+ // CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim8()> {
293+ // CHECK-NEXT: static constexpr bool value = true;
294+ // CHECK-NEXT: };
295+ // CHECK-NEXT: }
252296
253297// CHECK: #include <sycl/kernel_bundle.hpp>
254298
@@ -307,3 +351,11 @@ template void ff_6(Agg S1, Derived S2, int);
307351// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i"});
308352// CHECK-NEXT: }
309353// CHECK-NEXT: }
354+
355+ // CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
356+ // CHECK-NEXT: namespace sycl {
357+ // CHECK-NEXT: template <>
358+ // CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim8()>() {
359+ // CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE"});
360+ // CHECK-NEXT: }
361+ // CHECK-NEXT: }
0 commit comments