Skip to content

Commit 2bb7c21

Browse files
dklochkov-embdklochkov-emb
committed
[SYCL] add unit tests of free function namespace support
1 parent a5edf00 commit 2bb7c21

File tree

1 file changed

+248
-0
lines changed

1 file changed

+248
-0
lines changed

‎clang/test/CodeGenSYCL/free_function_int_header.cpp‎

Lines changed: 248 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,47 @@ __attribute__((sycl_device))
101101
void ff_8(sycl::work_group_memory<int>) {
102102
}
103103

104+
// function in namespace
105+
namespace free_functions {
106+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
107+
void ff_9(int start, int *ptr) {
108+
}
109+
}
110+
111+
// function in nested namespace
112+
namespace free_functions::tests {
113+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
114+
void ff_10(int start, int *ptr) {
115+
}
116+
}
117+
118+
// function in inline namespace
119+
namespace free_functions::tests {
120+
inline namespace V1 {
121+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
122+
void ff_11(int start, int *ptr) {
123+
}
124+
}
125+
}
126+
127+
//function in anonymous namespace
128+
namespace {
129+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
130+
void ff_12(int start, int *ptr) {
131+
}
132+
}
133+
134+
// functions with the same name but in different namespaces
135+
namespace free_functions {
136+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
137+
void ff_13(int start, int *ptr) {
138+
}
139+
}
140+
namespace free_functions::tests {
141+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
142+
void ff_13(int start, int *ptr) {
143+
}
144+
}
104145

105146
// CHECK: const char* const kernel_names[] = {
106147
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piii
@@ -112,6 +153,12 @@ void ff_8(sycl::work_group_memory<int>) {
112153
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
113154
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
114155
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
156+
// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions4ff_9EiPi
157+
// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests5ff_10EiPi
158+
// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests2V15ff_11EiPi
159+
// CHECK-NEXT: {{.*}}__sycl_kernel__GLOBAL__N_15ff_12EiPi
160+
// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5ff_13EiPi
161+
// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests5ff_13EiPi
115162
// CHECK-NEXT: ""
116163
// CHECK-NEXT: };
117164

@@ -158,6 +205,30 @@ void ff_8(sycl::work_group_memory<int>) {
158205
// CHECK: //--- _Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
159206
// CHECK-NEXT: { kernel_param_kind_t::kind_work_group_memory, 8, 0 },
160207

208+
// CHECK: //--- _ZN28__sycl_kernel_free_functions4ff_9EiPi
209+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
210+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 },
211+
212+
// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi
213+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
214+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 },
215+
216+
// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi
217+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
218+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 },
219+
220+
// CHECK: //--- _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi
221+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
222+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 },
223+
224+
// CHECK: //--- _ZN28__sycl_kernel_free_functions5ff_13EiPi
225+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
226+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 },
227+
228+
// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi
229+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
230+
// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 },
231+
161232
// CHECK: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
162233
// CHECK-NEXT: };
163234

@@ -324,6 +395,135 @@ void ff_8(sycl::work_group_memory<int>) {
324395
// CHECK-NEXT: };
325396
// CHECK-NEXT: }
326397

398+
// CHECK: Definition of _ZN28__sycl_kernel_free_functions4ff_9EiPi as a free function kernel
399+
// CHECK: Forward declarations of kernel and its argument types:
400+
// CHECK: namespace free_functions {
401+
// CHECK-NEXT: void ff_9(int start, int * ptr);
402+
// CHECK-NEXT: } // namespace free_functions
403+
404+
// CHECK: static constexpr auto __sycl_shim10() {
405+
// CHECK-NEXT: return (void (*)(int, int *))free_functions::ff_9;
406+
// CHECK-NEXT: }
407+
408+
// CHECK: namespace sycl {
409+
// CHECK-NEXT: template <>
410+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim10()> {
411+
// CHECK-NEXT: static constexpr bool value = true;
412+
// CHECK-NEXT: };
413+
// CHECK-NEXT: template <>
414+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim10()> {
415+
// CHECK-NEXT: static constexpr bool value = true;
416+
// CHECK-NEXT: };
417+
// CHECK-NEXT: }
418+
419+
// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi as a free function kernel
420+
// CHECK: Forward declarations of kernel and its argument types:
421+
422+
// CHECK: namespace free_functions {
423+
// CHECK-NEXT: namespace tests {
424+
// CHECK-NEXT: void ff_10(int start, int * ptr);
425+
// CHECK-NEXT: } // namespace tests
426+
// CHECK-NEXT: } // namespace free_functions
427+
// CHECK: static constexpr auto __sycl_shim11() {
428+
// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::ff_10;
429+
// CHECK-NEXT: }
430+
// CHECK-NEXT: namespace sycl {
431+
// CHECK-NEXT: template <>
432+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim11()> {
433+
// CHECK-NEXT: static constexpr bool value = true;
434+
// CHECK-NEXT: };
435+
// CHECK-NEXT: template <>
436+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim11()> {
437+
// CHECK-NEXT: static constexpr bool value = true;
438+
// CHECK-NEXT: };
439+
// CHECK-NEXT: }
440+
441+
// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi as a free function kernel
442+
// CHECK: Forward declarations of kernel and its argument types:
443+
444+
// CHECK: namespace free_functions {
445+
// CHECK-NEXT: namespace tests {
446+
// CHECK-NEXT: inline namespace V1 {
447+
// CHECK-NEXT: void ff_11(int start, int * ptr);
448+
// CHECK-NEXT: } // inline namespace V1
449+
// CHECK-NEXT: } // namespace tests
450+
// CHECK-NEXT: } // namespace free_functions
451+
// CHECK: static constexpr auto __sycl_shim12() {
452+
// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::V1::ff_11;
453+
// CHECK-NEXT: }
454+
// CHECK-NEXT: namespace sycl {
455+
// CHECK-NEXT: template <>
456+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim12()> {
457+
// CHECK-NEXT: static constexpr bool value = true;
458+
// CHECK-NEXT: };
459+
// CHECK-NEXT: template <>
460+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim12()> {
461+
// CHECK-NEXT: static constexpr bool value = true;
462+
// CHECK-NEXT: };
463+
// CHECK-NEXT: }
464+
465+
// CHECK: Definition of _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi as a free function kernel
466+
// CHECK: Forward declarations of kernel and its argument types:
467+
468+
// CHECK: namespace {
469+
// CHECK-NEXT: void ff_12(int start, int * ptr);
470+
// CHECK-NEXT: } // namespace
471+
// CHECK: static constexpr auto __sycl_shim13() {
472+
// CHECK-NEXT: return (void (*)(int, int *))ff_12;
473+
// CHECK-NEXT: }
474+
// CHECK-NEXT: namespace sycl {
475+
// CHECK-NEXT: template <>
476+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim13()> {
477+
// CHECK-NEXT: static constexpr bool value = true;
478+
// CHECK-NEXT: };
479+
// CHECK-NEXT: template <>
480+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim13()> {
481+
// CHECK-NEXT: static constexpr bool value = true;
482+
// CHECK-NEXT: };
483+
// CHECK-NEXT: }
484+
485+
// CHECK: Definition of _ZN28__sycl_kernel_free_functions5ff_13EiPi as a free function kernel
486+
// CHECK: Forward declarations of kernel and its argument types:
487+
488+
// CHECK: namespace free_functions {
489+
// CHECK-NEXT: void ff_13(int start, int * ptr);
490+
// CHECK-NEXT: } // namespace free_functions
491+
// CHECK: static constexpr auto __sycl_shim14() {
492+
// CHECK-NEXT: return (void (*)(int, int *))free_functions::ff_13;
493+
// CHECK-NEXT: }
494+
// CHECK-NEXT: namespace sycl {
495+
// CHECK-NEXT: template <>
496+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim14()> {
497+
// CHECK-NEXT: static constexpr bool value = true;
498+
// CHECK-NEXT: };
499+
// CHECK-NEXT: template <>
500+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim14()> {
501+
// CHECK-NEXT: static constexpr bool value = true;
502+
// CHECK-NEXT: };
503+
// CHECK-NEXT: }
504+
505+
// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi as a free function kernel
506+
// CHECK: Forward declarations of kernel and its argument types:
507+
508+
// CHECK: namespace free_functions {
509+
// CHECK-NEXT: namespace tests {
510+
// CHECK-NEXT: void ff_13(int start, int * ptr);
511+
// CHECK-NEXT: } // namespace tests
512+
// CHECK-NEXT: } // namespace free_functions
513+
// CHECK: static constexpr auto __sycl_shim15() {
514+
// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::ff_13;
515+
// CHECK-NEXT: }
516+
// CHECK-NEXT: namespace sycl {
517+
// CHECK-NEXT: template <>
518+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim15()> {
519+
// CHECK-NEXT: static constexpr bool value = true;
520+
// CHECK-NEXT: };
521+
// CHECK-NEXT: template <>
522+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim15()> {
523+
// CHECK-NEXT: static constexpr bool value = true;
524+
// CHECK-NEXT: };
525+
// CHECK-NEXT: }
526+
327527
// CHECK: #include <sycl/kernel_bundle.hpp>
328528

329529
// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
@@ -397,3 +597,51 @@ void ff_8(sycl::work_group_memory<int>) {
397597
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE"});
398598
// CHECK-NEXT: }
399599
// CHECK-NEXT: }
600+
601+
// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions4ff_9EiPi
602+
// CHECK-NEXT: namespace sycl {
603+
// CHECK-NEXT: template <>
604+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim10()>() {
605+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions4ff_9EiPi"});
606+
// CHECK-NEXT: }
607+
// CHECK-NEXT: }
608+
609+
// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi
610+
// CHECK-NEXT: namespace sycl {
611+
// CHECK-NEXT: template <>
612+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim11()>() {
613+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_10EiPi"});
614+
// CHECK-NEXT: }
615+
// CHECK-NEXT: }
616+
617+
// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi
618+
// CHECK-NEXT: namespace sycl {
619+
// CHECK-NEXT: template <>
620+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim12()>() {
621+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi"});
622+
// CHECK-NEXT: }
623+
// CHECK-NEXT: }
624+
625+
// CHECK: Definition of kernel_id of _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi
626+
// CHECK-NEXT: namespace sycl {
627+
// CHECK-NEXT: template <>
628+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim13()>() {
629+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi"});
630+
// CHECK-NEXT: }
631+
// CHECK-NEXT: }
632+
633+
// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_13EiPi
634+
// CHECK-NEXT: namespace sycl {
635+
// CHECK-NEXT: template <>
636+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim14()>() {
637+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_13EiPi"});
638+
// CHECK-NEXT: }
639+
// CHECK-NEXT: }
640+
641+
// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi
642+
// CHECK-NEXT: namespace sycl {
643+
// CHECK-NEXT: template <>
644+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim15()>() {
645+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_13EiPi"});
646+
// CHECK-NEXT: }
647+
// CHECK-NEXT: }

0 commit comments

Comments
 (0)