diff --git a/sycl/test/extensions/free_function_errors.cpp b/sycl/test/extensions/free_function_errors.cpp index 2e0641fbd9023..d196ad00c4586 100644 --- a/sycl/test/extensions/free_function_errors.cpp +++ b/sycl/test/extensions/free_function_errors.cpp @@ -23,8 +23,6 @@ union U { float f; }; -using accType = accessor; - SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( (ext::oneapi::experimental::single_task_kernel)) void ff(struct S s) {} diff --git a/sycl/test/extensions/free_function_kernels/free_function_kernels_type_aliases.cpp b/sycl/test/extensions/free_function_kernels/free_function_kernels_type_aliases.cpp new file mode 100644 index 0000000000000..c92af27377881 --- /dev/null +++ b/sycl/test/extensions/free_function_kernels/free_function_kernels_type_aliases.cpp @@ -0,0 +1,176 @@ +// RUN: %clangxx -fsyntax-only -fsycl %s +// This test verifies whether type aliases are allowed as kernel parameter to +// free function kernel. +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +struct TestStruct {}; +using StructAlias = TestStruct; +typedef TestStruct StructAliasTypedef; + +class TestClass {}; +using ClassAlias = TestClass; +typedef TestClass ClassAliasTypedef; +namespace ns1 { +using WriteOnlyAcc = sycl::accessor; +typedef sycl::accessor + WriteOnlyAccAliasTypedef; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskKernelStruct(StructAlias Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelStruct(StructAlias Type) {} + +namespace ns2 { + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskKernelStructAliasTypedef(StructAliasTypedef Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelStructAliasTypedef(StructAliasTypedef Type) {} + +template struct TestStructY {}; + +template +using AliasTypeY = TestStructY; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskKernelClass(ClassAlias Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelClass(ClassAlias Type) {} + +namespace ns3 { + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskKernelClassAliasTypedef(ClassAliasTypedef Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelClassAliasTypedef(ClassAliasTypedef Type) {} + +using LocalAcc = sycl::local_accessor; +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskWriteOnlyAcc(WriteOnlyAcc Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelWriteOnlyAcc(WriteOnlyAcc Type) {} + +namespace ns4 { +template struct TestStructX {}; +using AliasTypeX = TestStructX; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskWriteOnlyAccAliasTypedef(WriteOnlyAccAliasTypedef Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelWriteOnlyAccAliasTypedef(WriteOnlyAccAliasTypedef Type) {} + +} // namespace ns4 +} // namespace ns3 +} // namespace ns2 +} // namespace ns1 + +namespace ns5 { +using ReadOnlyAcc = sycl::accessor; +typedef sycl::accessor + ReadOnlyAccAliasTypedef; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskReadOnlyAcc(ReadOnlyAcc Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelReadOnlyAcc(ReadOnlyAcc Type) {} + +} // namespace ns5 + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskReadOnlyAccAliasTypedef(ns5::ReadOnlyAccAliasTypedef Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelReadOnlyAccAliasTypedef(ns5::ReadOnlyAccAliasTypedef Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskLocalAcc(ns1::ns2::ns3::LocalAcc Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelLocalAcc(ns1::ns2::ns3::LocalAcc Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskTemplatedType(ns1::ns2::ns3::ns4::AliasTypeX Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelTemplatedType(ns1::ns2::ns3::ns4::AliasTypeX Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel) +void singleTaskTemplatedTypeAliasTemplatedType( + ns1::ns2::AliasTypeY Type) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>) +void ndRangeKernelTemplatedTypeAliasTemplatedType( + ns1::ns2::AliasTypeY Type) {} + +template void runNdRangeCheck(T Type) { + sycl::queue Queue; + sycl::kernel_bundle Bundle = + get_kernel_bundle(Queue.get_context()); + sycl::kernel_id Id = syclexp::get_kernel_id(); + sycl::kernel Kernel = Bundle.get_kernel(Id); + Queue.submit([&](sycl::handler &Handler) { + Handler.set_args(Type); + Handler.parallel_for(sycl::nd_range{{1}, {1}}, Kernel); + }); +} + +template void runSingleTaskTest(T Type) { + sycl::queue Queue; + sycl::kernel_bundle Bundle = + get_kernel_bundle(Queue.get_context()); + sycl::kernel_id Id = syclexp::get_kernel_id(); + sycl::kernel Kernel = Bundle.get_kernel(Id); + Queue.submit([&](sycl::handler &Handler) { + Handler.set_args(Type); + Handler.single_task(Kernel); + }); +} + +int main() { + runSingleTaskTest(StructAlias()); + runSingleTaskTest(ClassAlias()); + runSingleTaskTest(ns5::ReadOnlyAcc()); + runSingleTaskTest(ns1::WriteOnlyAcc()); + runSingleTaskTest(ns1::ns2::ns3::LocalAcc()); + runSingleTaskTest(ns1::ns2::ns3::ns4::AliasTypeX()); + runSingleTaskTest( + ns1::ns2::AliasTypeY()); + runSingleTaskTest( + StructAliasTypedef()); + runSingleTaskTest( + ClassAliasTypedef()); + runSingleTaskTest( + ns1::WriteOnlyAccAliasTypedef()); + runSingleTaskTest( + ns5::ReadOnlyAccAliasTypedef()); + + runNdRangeCheck(StructAlias()); + runNdRangeCheck(ClassAlias()); + runNdRangeCheck(ns5::ReadOnlyAcc()); + runNdRangeCheck( + ns1::WriteOnlyAcc()); + runNdRangeCheck(ns1::ns2::ns3::LocalAcc()); + runNdRangeCheck(ns1::ns2::ns3::ns4::AliasTypeX()); + runNdRangeCheck( + ns1::ns2::AliasTypeY()); + + runNdRangeCheck( + StructAliasTypedef()); + runNdRangeCheck( + ClassAliasTypedef()); + runNdRangeCheck( + ns1::WriteOnlyAccAliasTypedef()); + runNdRangeCheck( + ns5::ReadOnlyAccAliasTypedef()); + + return 0; +}