From 08a0aafc96eea0b53e4caa51ee33d5f339c25f42 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Mon, 7 Apr 2025 16:31:16 +0200 Subject: [PATCH 01/24] [SYCL][Doc] Introduce test plan for free function kernels extension --- .../test-e2e/FreeFunctionKernels/test-plan.md | 150 ++++++++++++++++++ 1 file changed, 150 insertions(+) create mode 100644 sycl/test-e2e/FreeFunctionKernels/test-plan.md diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md new file mode 100644 index 0000000000000..2516fdbf729df --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -0,0 +1,150 @@ +# Test plan for [`sycl_ext_oneapi_free_function_kernels`][spec-link] extension + +## Testing scope + +### Device coverage + +The tests should be launched on every supported device configuration we have. + +### Type coverage +New APIs for new way to define a kernel as a simple C++ function, where the kernel arguments are parameters to this function described by the extension can take only allowed kernel parameter type as specified in section 4.12.4 "Rules for parameter passing to kernels" of the core SYCL specification with exceptions to parameters of type `reducer` or `kernel_handler`. + +Therefore those tests should be structured in a way that checks are performed on all allowed kernel parameter types to ensure that everything works correctly. + +## Tests + +### Unit tests + +Tests in this category may not fully exercise the extension functionality, but are instead they are focused on making sure that all APIs are consistent with respect to other APIs. + +Perform tests on new traits for free function kernels which should check the following: + - that `is_nd_range_kernel_v` trait returns true whose function declaration is decorated with `nd_range_kernel` property and when it is not then it returns false. + + - that `is_single_task_kernel_v` trait returns true function whose declaration is decorated with `single_task_kernel` and when it is not then it returns false. + +- that `is_kernel_v` trait returns true for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it returns false. + + +Perform tests on new kernel bundle member functions for free function kernels which should check the following: + +- that `get_kernel_id` member function returns valid kernel identifier that is associated with that kernel. + +- that `get_kernel_bundle(const context& ctxt)` member function returns kernel bundle in which associated free function kernel can be found. + +- that `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns kernel bundle in which associated free function kernel can be found. + +- that `has_kernel_bundle(const context& ctxt)` returns true when free function kernel can be represent in a device image of state and free function kernel is compatible with at least one of the devices in context. + +- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represent in a device image of state and free function kernel is compatible with at least one of the devices. + +- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represent in a device image of state and free function kernel is compatible with at least one of the devices. + +- that `is_compatible(const device& dev)` returns true when free function kernel is compatible with the device. + +- that `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the free function kernel. + +- that `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the free function kernel and if that kernel is compatible with the device. + +- that `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. + +- that `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the free function kernel does not reside in this kernel bundle. + + +Perform tests on new free functions to query kernel information descriptors which should check the following: + +- that `get_kernel_info(const context& ctxt)` produces the same result as would be computed by + ``` + auto bundle = sycl::get_kernel_bundle(ctxt); + auto ret = bundle.ext_oneapi_get_kernel().get_info()` + ``` + +- that `get_kernel_info(const context& ctxt, const device& dev)` produces the same result as would be computed by + ``` + auto bundle = + sycl::get_kernel_bundle(ctxt); + auto ret = bundle.ext_oneapi_get_kernel().get_info(dev); + ``` + +- that ` get_kernel_info(const queue& q)` produces the same result as would be computed by + ``` + sycl::context ctxt = q.get_context(); + sycl::device dev = q.get_device(); + auto bundle = + sycl::get_kernel_bundle(ctxt); + auto ret = bundle.ext_oneapi_get_kernel().get_info(dev); + ``` + +### E2E tests + +Tests in this category perform some meaningful actions with the extension to +see that the extension works in a scenarios which mimic real-life usage of the +extension. + +With the exception of the`single_task_kernel` free function kernels, all subsequent tests are executed with Dimensions = 1, 2, 3. + +#### Test `accessor` as kernel parameter to `single_task_kernel`: + +A series of tests should be performed that accessor in `single_task_kernel` is supported when +templated with `target::device`, `target::global_buffer`, `target::constant_buffer` and `target::local` inside free function kernel when passed as kernel parameter. + + + +#### Test `USM` pointer as kernel parameter to `single_task_kernel`: +A series of checks should be performed that USM memory in `single_task_kernel` with with three types of memory allocations `host`, `device` and `shared` inside free function kernel when passed as kernel parameter. + +#### Test `id` as kernel parameter to `single_task_kernel`: +A series of checks should be performed that we can pass `id` with different dimensions = 1, 2, 3 as kernel parameter to `single_task_kernel` and use it within kernel. + +#### Test `range` as kernel parameter to `single_task_kernel`: +A series of checks should be performed that we can pass `range` with different dimensions = 1, 2, 3 as kernel parameter to `single_task_kernel` and use it within kernel. + +#### Test `marray` when T is device copyable as kernel parameter to `single_task_kernel`: +A series of checks should be performed that we can pass `marray` as kernel parameter to `single_task_kernel` and use it within kernel. + +#### Test `vec` as kernel parameter to `single_task_kernel`: +A series of checks should be performed that we can pass `vec` as kernel parameter to `single_task_kernel` and use it within kernel. + +#### Test `sampled_image_accessor` as kernel parameter to `single_task_kernel`: +A series of checks should be performed that we can pass `sampled_image_accessor` as kernel parameter to `single_task_kernel` and use it within kernel. + +#### Test `unsampled_image_accessor` as kernel parameter to `single_task_kernel`: +A series of checks should be performed that we can pass `unsampled_image_accessor` as kernel parameter to `single_task_kernel` and use it within kernel. + +#### Test `accessor` as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that accessor in `nd_range_kernel` is supported when +templated with `target::device`, `target::global_buffer`, `target::constant_buffer` and `target::local` inside free function kernel when passed as kernel parameter. + +#### Test `local_accessor` as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that we can pass `local_accessor` as kernel parameter to `nd_range_kernel` and use it within kernel + +#### Test `USM` pointer as kernel parameter to `single_task_kernel`: +A series of checks should be performed that USM memory in `single_task_kernel` with with three types of memory allocations `host`, `device` and `shared` inside free function kernel when passed as kernel parameter. + +#### Test `id` as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that we can pass `id` with different dimensions = 1, 2, 3 as kernel parameter to `nd_range_kernel` and use it within kernel. + +#### Test `range` as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that we can pass `range` with different dimensions = 1, 2, 3 as kernel parameter to `nd_range_kernel` and use it within kernel. + +#### Test `marray` when T is device copyable as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that we can pass `marray` as kernel parameter to `nd_range_kernel` and use it within kernel. + +#### Test `vec` as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that we can pass `vec` as kernel parameter to `nd_range_kernel` and use it within kernel. + +#### Test `sampled_image_accessor` as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that we can pass `sampled_image_accessor` as kernel parameter to `nd_range_kernel` and use it within kernel. + +#### Test `unsampled_image_accessor` as kernel parameter to `nd_range_kernel`: +A series of checks should be performed that we can pass `unsampled_image_accessor` as kernel parameter to `nd_range_kernel` and use it within kernel. + +#### Interaction with additional kernel properties +A series of checks should be performed to check that to the free function kernels may also be decorated with any of the properties defined in `sycl_ext_oneapi_kernel_properties`. This test should perform simple checks verifying if applied kernel_properties work within defined kernels. + +#### Free function kernels compatibility with L0 backend +A series of checks should be performed to check compatibility of free function kernels with Level Zero Backend without going through the SYCL host runtime. + +#### Free function kernels compatibility with OpenCL backend +A series of checks should be performed to check compatibility of free function kernels with OpenCL Backend without going through the SYCL host runtime. + +[spec-link]: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc \ No newline at end of file From 8a2431686b63ea10b96695755c96c6b077392f24 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Mon, 7 Apr 2025 16:37:14 +0200 Subject: [PATCH 02/24] fixes --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 2516fdbf729df..98626da3c2664 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -18,9 +18,9 @@ Therefore those tests should be structured in a way that checks are performed on Tests in this category may not fully exercise the extension functionality, but are instead they are focused on making sure that all APIs are consistent with respect to other APIs. Perform tests on new traits for free function kernels which should check the following: - - that `is_nd_range_kernel_v` trait returns true whose function declaration is decorated with `nd_range_kernel` property and when it is not then it returns false. + - that `is_nd_range_kernel_v` trait returns true if function declaration is decorated with `nd_range_kernel` property and false if it is not. - - that `is_single_task_kernel_v` trait returns true function whose declaration is decorated with `single_task_kernel` and when it is not then it returns false. + - that `is_single_task_kernel_v` trait returns true function if declaration is decorated with `single_task_kernel` and false if it is not. - that `is_kernel_v` trait returns true for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it returns false. @@ -74,7 +74,7 @@ Perform tests on new free functions to query kernel information descriptors whic auto ret = bundle.ext_oneapi_get_kernel().get_info(dev); ``` -### E2E tests +### End-to-end tests Tests in this category perform some meaningful actions with the extension to see that the extension works in a scenarios which mimic real-life usage of the From 59cd1ff8c9263c267316dfa03cb51210f205f316 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 8 Apr 2025 15:58:44 +0200 Subject: [PATCH 03/24] fix --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 98626da3c2664..a69c372fcac90 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -85,7 +85,7 @@ With the exception of the`single_task_kernel` free function kernels, all subsequ #### Test `accessor` as kernel parameter to `single_task_kernel`: A series of tests should be performed that accessor in `single_task_kernel` is supported when -templated with `target::device`, `target::global_buffer`, `target::constant_buffer` and `target::local` inside free function kernel when passed as kernel parameter. +templated with `target::device`, inside free function kernel when passed as kernel parameter. @@ -112,7 +112,7 @@ A series of checks should be performed that we can pass `unsampled_image_accesso #### Test `accessor` as kernel parameter to `nd_range_kernel`: A series of checks should be performed that accessor in `nd_range_kernel` is supported when -templated with `target::device`, `target::global_buffer`, `target::constant_buffer` and `target::local` inside free function kernel when passed as kernel parameter. +templated with `target::device`, inside free function kernel when passed as kernel parameter. #### Test `local_accessor` as kernel parameter to `nd_range_kernel`: A series of checks should be performed that we can pass `local_accessor` as kernel parameter to `nd_range_kernel` and use it within kernel From b6203765bca2290f738d84e1fcbb1cac79ea40be Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 9 Apr 2025 19:11:01 +0200 Subject: [PATCH 04/24] applied changes based on feedback --- .../test-e2e/FreeFunctionKernels/test-plan.md | 79 ++++++++----------- 1 file changed, 33 insertions(+), 46 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index a69c372fcac90..6a69ae00ffbd5 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -17,6 +17,19 @@ Therefore those tests should be structured in a way that checks are performed on Tests in this category may not fully exercise the extension functionality, but are instead they are focused on making sure that all APIs are consistent with respect to other APIs. + +Perform tests on free function kernels requirments which should check the following: + - that compiler will emit diagnostic when free function kernel is declared with parameters of type `reducer` or `kernel_handler`. + - that compiler will emit diagnostic when free function kernel is declared with variadic arguments. + + - that compiler will emit diagnostic when free function kernel provides default parameter values. + + - that compiler will emit diagnostic when free function kernel return type is not `void`. + + - that compiler will emit diagnostic if free function kernel is decorated with more than one of the following properites (`nd_range_kernel` or `single_task_kernel `). + + - that compiler will emit diagnostic when free function kernel is decorated with multiple instances of the same `nd_range_kernel` property with different argument. + Perform tests on new traits for free function kernels which should check the following: - that `is_nd_range_kernel_v` trait returns true if function declaration is decorated with `nd_range_kernel` property and false if it is not. @@ -80,63 +93,37 @@ Tests in this category perform some meaningful actions with the extension to see that the extension works in a scenarios which mimic real-life usage of the extension. -With the exception of the`single_task_kernel` free function kernels, all subsequent tests are executed with Dimensions = 1, 2, 3. - -#### Test `accessor` as kernel parameter to `single_task_kernel`: - -A series of tests should be performed that accessor in `single_task_kernel` is supported when -templated with `target::device`, inside free function kernel when passed as kernel parameter. - - - -#### Test `USM` pointer as kernel parameter to `single_task_kernel`: -A series of checks should be performed that USM memory in `single_task_kernel` with with three types of memory allocations `host`, `device` and `shared` inside free function kernel when passed as kernel parameter. - -#### Test `id` as kernel parameter to `single_task_kernel`: -A series of checks should be performed that we can pass `id` with different dimensions = 1, 2, 3 as kernel parameter to `single_task_kernel` and use it within kernel. - -#### Test `range` as kernel parameter to `single_task_kernel`: -A series of checks should be performed that we can pass `range` with different dimensions = 1, 2, 3 as kernel parameter to `single_task_kernel` and use it within kernel. - -#### Test `marray` when T is device copyable as kernel parameter to `single_task_kernel`: -A series of checks should be performed that we can pass `marray` as kernel parameter to `single_task_kernel` and use it within kernel. - -#### Test `vec` as kernel parameter to `single_task_kernel`: -A series of checks should be performed that we can pass `vec` as kernel parameter to `single_task_kernel` and use it within kernel. - -#### Test `sampled_image_accessor` as kernel parameter to `single_task_kernel`: -A series of checks should be performed that we can pass `sampled_image_accessor` as kernel parameter to `single_task_kernel` and use it within kernel. +With the exception of the`single_task_kernel` free function kernels, all subsequent tests are executed with Dimensions = 1, 2, 3. -#### Test `unsampled_image_accessor` as kernel parameter to `single_task_kernel`: -A series of checks should be performed that we can pass `unsampled_image_accessor` as kernel parameter to `single_task_kernel` and use it within kernel. +In all subsequent tests, free function kernels should be declared within a namespace, as static member functions of a class, or in the global namespace. -#### Test `accessor` as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that accessor in `nd_range_kernel` is supported when +#### Test `accessor` as kernel parameter: +A series of tests should be performed that `accessor` is supported when templated with `target::device`, inside free function kernel when passed as kernel parameter. -#### Test `local_accessor` as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that we can pass `local_accessor` as kernel parameter to `nd_range_kernel` and use it within kernel +#### Test `USM` pointer as kernel parameter: +A series of checks should be performed that USM memory with three types of memory allocations `host`, `device` and `shared` is supported inside free function kernel when passed as kernel parameter. -#### Test `USM` pointer as kernel parameter to `single_task_kernel`: -A series of checks should be performed that USM memory in `single_task_kernel` with with three types of memory allocations `host`, `device` and `shared` inside free function kernel when passed as kernel parameter. +#### Test `id` as kernel parameter: +A series of checks should be performed that we can pass `id` with different dimensions = 1, 2, 3 as kernel parameter to free function kernel and use it within kernel. -#### Test `id` as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that we can pass `id` with different dimensions = 1, 2, 3 as kernel parameter to `nd_range_kernel` and use it within kernel. +#### Test `range` as kernel parameter: +A series of checks should be performed that we can pass `range` with different dimensions = 1, 2, 3 as kernel parameter to free function kernel and use it within kernel. -#### Test `range` as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that we can pass `range` with different dimensions = 1, 2, 3 as kernel parameter to `nd_range_kernel` and use it within kernel. +#### Test `marray` when `T` is device copyable as kernel parameter: +A series of checks should be performed that we can pass `marray` as kernel parameter to free function kernel and use it within kernel. -#### Test `marray` when T is device copyable as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that we can pass `marray` as kernel parameter to `nd_range_kernel` and use it within kernel. +#### Test `vec` when `T` is device copyable as kernel parameter: +A series of checks should be performed that we can pass `vec` as kernel parameter to free function kernel and use it within kernel. -#### Test `vec` as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that we can pass `vec` as kernel parameter to `nd_range_kernel` and use it within kernel. +#### Test `sampled_image_accessor` as kernel parameter: +A series of checks should be performed that we can pass `sampled_image_accessor` as kernel parameter to free function kernel and use it within kernel. -#### Test `sampled_image_accessor` as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that we can pass `sampled_image_accessor` as kernel parameter to `nd_range_kernel` and use it within kernel. +#### Test `unsampled_image_accessor` as kernel parameter: +A series of checks should be performed that we can pass `unsampled_image_accessor` as kernel parameter to free function kernel and use it within kernel. -#### Test `unsampled_image_accessor` as kernel parameter to `nd_range_kernel`: -A series of checks should be performed that we can pass `unsampled_image_accessor` as kernel parameter to `nd_range_kernel` and use it within kernel. +#### Test `local_accessor` as kernel parameter: +A series of checks should be performed that we can pass `local_accessor` as kernel parameter to free function kernel and use it within kernel #### Interaction with additional kernel properties A series of checks should be performed to check that to the free function kernels may also be decorated with any of the properties defined in `sycl_ext_oneapi_kernel_properties`. This test should perform simple checks verifying if applied kernel_properties work within defined kernels. From 7e8a99d73ffc43e4f3579dfae72bdb44b3bf6336 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 9 Apr 2025 19:18:54 +0200 Subject: [PATCH 05/24] fix --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 6a69ae00ffbd5..416d0e9769e4e 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -19,7 +19,10 @@ Tests in this category may not fully exercise the extension functionality, but a Perform tests on free function kernels requirments which should check the following: + - that compiler will emit diagnostic when free function kernel is declared with reference types as parameters. + - that compiler will emit diagnostic when free function kernel is declared with parameters of type `reducer` or `kernel_handler`. + - that compiler will emit diagnostic when free function kernel is declared with variadic arguments. - that compiler will emit diagnostic when free function kernel provides default parameter values. From 12176ce033ec42bdb652ca6165332cf15cde02f7 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Thu, 10 Apr 2025 15:17:28 +0200 Subject: [PATCH 06/24] fix --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 7 ------- 1 file changed, 7 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 416d0e9769e4e..01cfa18aa7164 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -21,18 +21,12 @@ Tests in this category may not fully exercise the extension functionality, but a Perform tests on free function kernels requirments which should check the following: - that compiler will emit diagnostic when free function kernel is declared with reference types as parameters. - - that compiler will emit diagnostic when free function kernel is declared with parameters of type `reducer` or `kernel_handler`. - - that compiler will emit diagnostic when free function kernel is declared with variadic arguments. - that compiler will emit diagnostic when free function kernel provides default parameter values. - that compiler will emit diagnostic when free function kernel return type is not `void`. - - that compiler will emit diagnostic if free function kernel is decorated with more than one of the following properites (`nd_range_kernel` or `single_task_kernel `). - - - that compiler will emit diagnostic when free function kernel is decorated with multiple instances of the same `nd_range_kernel` property with different argument. - Perform tests on new traits for free function kernels which should check the following: - that `is_nd_range_kernel_v` trait returns true if function declaration is decorated with `nd_range_kernel` property and false if it is not. @@ -40,7 +34,6 @@ Perform tests on new traits for free function kernels which should check the fol - that `is_kernel_v` trait returns true for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it returns false. - Perform tests on new kernel bundle member functions for free function kernels which should check the following: - that `get_kernel_id` member function returns valid kernel identifier that is associated with that kernel. From ac28fa9e4df707f103329d6d8395f782ef0ebaf1 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Mon, 14 Apr 2025 16:59:16 +0200 Subject: [PATCH 07/24] applied feedback and added addional tests to the test plan --- .../test-e2e/FreeFunctionKernels/test-plan.md | 44 ++++++++++++++++--- 1 file changed, 37 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 01cfa18aa7164..d416bad885549 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -27,6 +27,8 @@ Perform tests on free function kernels requirments which should check the follow - that compiler will emit diagnostic when free function kernel return type is not `void`. + - that compiler will emit diagnostic when a non-static member function is used as a kernel. Only static member function at class scope are allowed as free function kernel. + Perform tests on new traits for free function kernels which should check the following: - that `is_nd_range_kernel_v` trait returns true if function declaration is decorated with `nd_range_kernel` property and false if it is not. @@ -42,11 +44,11 @@ Perform tests on new kernel bundle member functions for free function kernels wh - that `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns kernel bundle in which associated free function kernel can be found. -- that `has_kernel_bundle(const context& ctxt)` returns true when free function kernel can be represent in a device image of state and free function kernel is compatible with at least one of the devices in context. +- that `has_kernel_bundle(const context& ctxt)` returns true when free function kernel can be represented in a device image of state and free function kernel is compatible with at least one of the devices in context. -- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represent in a device image of state and free function kernel is compatible with at least one of the devices. +- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image of state and free function kernel is compatible with at least one of the devices. -- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represent in a device image of state and free function kernel is compatible with at least one of the devices. +- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image of state and free function kernel is compatible with at least one of the devices. - that `is_compatible(const device& dev)` returns true when free function kernel is compatible with the device. @@ -58,6 +60,9 @@ Perform tests on new kernel bundle member functions for free function kernels wh - that `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the free function kernel does not reside in this kernel bundle. +- that `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. + +- that `info::kernel::num_args` returns the number of parameters in the function definition of free function kernel. Perform tests on new free functions to query kernel information descriptors which should check the following: @@ -83,6 +88,10 @@ Perform tests on new free functions to query kernel information descriptors whic auto ret = bundle.ext_oneapi_get_kernel().get_info(dev); ``` +Perform tests on use of illegal types for kernel paramters: +- that A class type S with a virtual base class of type T can not be used as free function kernel parameter type. +- that A class type S with a virtual member function can not be used as free function kernel parameter type. + ### End-to-end tests Tests in this category perform some meaningful actions with the extension to @@ -119,15 +128,36 @@ A series of checks should be performed that we can pass `sampled_image_accessor` A series of checks should be performed that we can pass `unsampled_image_accessor` as kernel parameter to free function kernel and use it within kernel. #### Test `local_accessor` as kernel parameter: -A series of checks should be performed that we can pass `local_accessor` as kernel parameter to free function kernel and use it within kernel +A series of checks should be performed that we can pass `local_accessor` as kernel parameter to free function kernel and use it within kernel. + +#### Test structs that contain one of the following `accessor`, `local_accessor`, `sampled_image_accessor` or `unsampled_image_accessor` types when used as kernel parameter: +A series of checks should be performed that we can pass struct that contain one of the following `accessor`, `local_accessor`, `sampled_image_accessor` or `unsampled_image_accessor` types as kernel parameter to free function kernel and use it within kernel. + +#### Test `struct` defined at namespace scope as kernel parameter: +A series of checks should be performed that we can pass `struct` as kernel parameter and use it within kernel. -#### Interaction with additional kernel properties +#### Test `class` defined at namespace scope as kernel parameter: +A series of checks should be performed that we can pass `class` as kernel parameter and use it within kernel. + +#### Test scoped enumeration defined at namespace scope as kernel parameter: +A series of checks should be performed that we can pass scoped enumeration as kernel parameter and use it within kernel. + +#### Test unscoped enumeration that has an explicit underlying type defined at namespace scope as kernel parameter: +A series of checks should be performed that we can pass unscoped enumeration that has an explicit underlying type as kernel parameter and use it within kernel. + +#### Test type aliases to allowed kernel paramater types as kernel parameter: +A series of checks should be performed that we can pass type aliases to allowed kernel paramater types as kernel parameter and use it within kernel. + +#### Interaction with additional kernel properties: A series of checks should be performed to check that to the free function kernels may also be decorated with any of the properties defined in `sycl_ext_oneapi_kernel_properties`. This test should perform simple checks verifying if applied kernel_properties work within defined kernels. -#### Free function kernels compatibility with L0 backend +#### Free function kernels compatibility with L0 backend: A series of checks should be performed to check compatibility of free function kernels with Level Zero Backend without going through the SYCL host runtime. -#### Free function kernels compatibility with OpenCL backend +#### Free function kernels compatibility with OpenCL backend: A series of checks should be performed to check compatibility of free function kernels with OpenCL Backend without going through the SYCL host runtime. +#### Test template support in free function kernels +A series of checks should be performed to check compatibility of free function kernels with templateed kernel parameters. + [spec-link]: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc \ No newline at end of file From fa943fb56635aac17008597ea4039ff200c234fa Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Mon, 14 Apr 2025 17:27:32 +0200 Subject: [PATCH 08/24] fixed typo --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index d416bad885549..311e0ebaa8fdc 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -157,7 +157,7 @@ A series of checks should be performed to check compatibility of free function k #### Free function kernels compatibility with OpenCL backend: A series of checks should be performed to check compatibility of free function kernels with OpenCL Backend without going through the SYCL host runtime. -#### Test template support in free function kernels +#### Test template support in free function kernels: A series of checks should be performed to check compatibility of free function kernels with templateed kernel parameters. [spec-link]: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc \ No newline at end of file From 6a341c9f20cdc7576044b607d0b2dd33fb7a2f30 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 15 Apr 2025 15:16:16 +0200 Subject: [PATCH 09/24] added new checks to test plan --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 311e0ebaa8fdc..2774dd6204fb1 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -29,6 +29,12 @@ Perform tests on free function kernels requirments which should check the follow - that compiler will emit diagnostic when a non-static member function is used as a kernel. Only static member function at class scope are allowed as free function kernel. +Perform tests on free function kernel declaration with properties `nd_range_kernel` and `single_task_kernel` which should check the following: + - that if the property does not appear on the first declaration of the function in the translation unit, it will result in a compilation error. + - that if the property appears on the first declaration of the function, and the following redeclarations do not have this property, it will not result in a compilation error. + - that if a function is decorated with more than one of these properties, it will result in a compilation error. + - that if a redeclaration of a function is decorated with the same property but with different arguments, the program should result in a compilation error. + Perform tests on new traits for free function kernels which should check the following: - that `is_nd_range_kernel_v` trait returns true if function declaration is decorated with `nd_range_kernel` property and false if it is not. @@ -64,6 +70,8 @@ Perform tests on new kernel bundle member functions for free function kernels wh - that `info::kernel::num_args` returns the number of parameters in the function definition of free function kernel. +- that performs all the checks mentioned above on a free function kernel that is declared in one translation unit and defined in a separate one. + Perform tests on new free functions to query kernel information descriptors which should check the following: - that `get_kernel_info(const context& ctxt)` produces the same result as would be computed by From 082a62991936b26c097ec023ce293b8978f9de29 Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:04:50 +0200 Subject: [PATCH 10/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 2774dd6204fb1..31988a2fafb21 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -7,7 +7,7 @@ The tests should be launched on every supported device configuration we have. ### Type coverage -New APIs for new way to define a kernel as a simple C++ function, where the kernel arguments are parameters to this function described by the extension can take only allowed kernel parameter type as specified in section 4.12.4 "Rules for parameter passing to kernels" of the core SYCL specification with exceptions to parameters of type `reducer` or `kernel_handler`. +New APIs for new way to define a kernel as a simple C++ function, where the kernel arguments are parameters to this function described by the extension can take only allowed kernel parameter type as specified in section [4.12.4 of the SYCL 2020 specification](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.parameter.passing) with exceptions to parameters of type `reducer` or `kernel_handler`. Therefore those tests should be structured in a way that checks are performed on all allowed kernel parameter types to ensure that everything works correctly. From 353a676f29800890a2831a8760b56f0888c16634 Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:05:04 +0200 Subject: [PATCH 11/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 31988a2fafb21..f53f805f61f36 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -15,7 +15,7 @@ Therefore those tests should be structured in a way that checks are performed on ### Unit tests -Tests in this category may not fully exercise the extension functionality, but are instead they are focused on making sure that all APIs are consistent with respect to other APIs. +Tests in this category may not fully exercise the extension functionality, but are instead they are focused on making sure that all APIs are consistent with respect to existing APIs. Perform tests on free function kernels requirments which should check the following: From b01a3f4cf0aaa322ce52e8c026128e02f49aa401 Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:05:14 +0200 Subject: [PATCH 12/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index f53f805f61f36..0829b6d2841d2 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -92,7 +92,7 @@ Perform tests on new free functions to query kernel information descriptors whic sycl::context ctxt = q.get_context(); sycl::device dev = q.get_device(); auto bundle = - sycl::get_kernel_bundle(ctxt); + sycl::get_kernel_bundle(ctxt); auto ret = bundle.ext_oneapi_get_kernel().get_info(dev); ``` From 429853f6932ab2732882eab22b0eecfbf5249abe Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:05:55 +0200 Subject: [PATCH 13/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 0829b6d2841d2..a3df9134b0ef9 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -118,7 +118,7 @@ templated with `target::device`, inside free function kernel when passed as kern A series of checks should be performed that USM memory with three types of memory allocations `host`, `device` and `shared` is supported inside free function kernel when passed as kernel parameter. #### Test `id` as kernel parameter: -A series of checks should be performed that we can pass `id` with different dimensions = 1, 2, 3 as kernel parameter to free function kernel and use it within kernel. +A series of checks should be performed that we can pass `id` where `Dimensions` is in {1, 2, 3} as kernel parameter to free function kernel and use it within kernel. #### Test `range` as kernel parameter: A series of checks should be performed that we can pass `range` with different dimensions = 1, 2, 3 as kernel parameter to free function kernel and use it within kernel. From 5f7770714561f473bd547b6b0fcebbe1e5eb5b6c Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:06:07 +0200 Subject: [PATCH 14/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index a3df9134b0ef9..1390182ff7138 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -106,7 +106,7 @@ Tests in this category perform some meaningful actions with the extension to see that the extension works in a scenarios which mimic real-life usage of the extension. -With the exception of the`single_task_kernel` free function kernels, all subsequent tests are executed with Dimensions = 1, 2, 3. +With the exception of the `single_task_kernel` free function kernels, all subsequent tests are executed with `Dimensions` $$\in {1, 2, 3}$$. In all subsequent tests, free function kernels should be declared within a namespace, as static member functions of a class, or in the global namespace. From e8d65653ca342a90bfb5a1adf4b52843ef49a9f9 Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:06:16 +0200 Subject: [PATCH 15/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 1390182ff7138..229364fbe4d23 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -97,8 +97,8 @@ Perform tests on new free functions to query kernel information descriptors whic ``` Perform tests on use of illegal types for kernel paramters: -- that A class type S with a virtual base class of type T can not be used as free function kernel parameter type. -- that A class type S with a virtual member function can not be used as free function kernel parameter type. +- that a class type `S` with a virtual base class of type `T` can not be used as free function kernel parameter type. +- that a class type `S` with a virtual member function can not be used as free function kernel parameter type. ### End-to-end tests From a74264c97991927e8c931acde46d1d8a19c41690 Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:06:36 +0200 Subject: [PATCH 16/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 229364fbe4d23..d0b361ab5f7a3 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -18,16 +18,16 @@ Therefore those tests should be structured in a way that checks are performed on Tests in this category may not fully exercise the extension functionality, but are instead they are focused on making sure that all APIs are consistent with respect to existing APIs. -Perform tests on free function kernels requirments which should check the following: - - that compiler will emit diagnostic when free function kernel is declared with reference types as parameters. +Perform tests on free function kernels requirements which should check that: + - the compiler will emit diagnostic when free function kernel is declared with reference types as parameters. - - that compiler will emit diagnostic when free function kernel is declared with variadic arguments. + - the compiler will emit diagnostic when free function kernel is declared with variadic arguments. - - that compiler will emit diagnostic when free function kernel provides default parameter values. + - the compiler will emit diagnostic when free function kernel provides default parameter values. - - that compiler will emit diagnostic when free function kernel return type is not `void`. + - the compiler will emit diagnostic when free function kernel return type is not `void`. - - that compiler will emit diagnostic when a non-static member function is used as a kernel. Only static member function at class scope are allowed as free function kernel. + - the compiler will emit diagnostic when a non-static member function is used as a kernel. Only static member function at class scope are allowed as free function kernel. Perform tests on free function kernel declaration with properties `nd_range_kernel` and `single_task_kernel` which should check the following: - that if the property does not appear on the first declaration of the function in the translation unit, it will result in a compilation error. From 9615e3260ec791816cb51fca8952254acc06bb7c Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:07:20 +0200 Subject: [PATCH 17/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- .../test-e2e/FreeFunctionKernels/test-plan.md | 54 +++++++++---------- 1 file changed, 27 insertions(+), 27 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index d0b361ab5f7a3..aaccb5bb22b7a 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -42,33 +42,33 @@ Perform tests on new traits for free function kernels which should check the fol - that `is_kernel_v` trait returns true for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it returns false. -Perform tests on new kernel bundle member functions for free function kernels which should check the following: - -- that `get_kernel_id` member function returns valid kernel identifier that is associated with that kernel. - -- that `get_kernel_bundle(const context& ctxt)` member function returns kernel bundle in which associated free function kernel can be found. - -- that `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns kernel bundle in which associated free function kernel can be found. - -- that `has_kernel_bundle(const context& ctxt)` returns true when free function kernel can be represented in a device image of state and free function kernel is compatible with at least one of the devices in context. - -- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image of state and free function kernel is compatible with at least one of the devices. - -- that `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image of state and free function kernel is compatible with at least one of the devices. - -- that `is_compatible(const device& dev)` returns true when free function kernel is compatible with the device. - -- that `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the free function kernel. - -- that `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the free function kernel and if that kernel is compatible with the device. - -- that `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. - -- that `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the free function kernel does not reside in this kernel bundle. - -- that `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. - -- that `info::kernel::num_args` returns the number of parameters in the function definition of free function kernel. +Perform tests on new `kernel_bundle` member functions for free function kernels which should check that: + + - the `get_kernel_id` member function returns a valid kernel identifier that is associated with that kernel. + + - the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. + + - the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. + + - the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. + + - the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. + + - the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. + + - the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. + + - the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. + + - the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. + + - the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. + + - the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. + + - the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. + + - the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. - that performs all the checks mentioned above on a free function kernel that is declared in one translation unit and defined in a separate one. From b745598b455dfbd0021961f2e1644f7c81b8117e Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 15 Apr 2025 17:08:03 +0200 Subject: [PATCH 18/24] Update sycl/test-e2e/FreeFunctionKernels/test-plan.md Co-authored-by: Steffen Larsen --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index aaccb5bb22b7a..1cf5bc8d70c2c 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -77,7 +77,7 @@ Perform tests on new free functions to query kernel information descriptors whic - that `get_kernel_info(const context& ctxt)` produces the same result as would be computed by ``` auto bundle = sycl::get_kernel_bundle(ctxt); - auto ret = bundle.ext_oneapi_get_kernel().get_info()` + auto ret = bundle.ext_oneapi_get_kernel().get_info(); ``` - that `get_kernel_info(const context& ctxt, const device& dev)` produces the same result as would be computed by From 34915254769e63cc28fc2b61b82482d4203cc36a Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 15 Apr 2025 18:25:32 +0200 Subject: [PATCH 19/24] reworked formatting and applied changes based on feedback --- .../test-e2e/FreeFunctionKernels/test-plan.md | 38 +++++++++---------- 1 file changed, 19 insertions(+), 19 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 1cf5bc8d70c2c..177a2148fd5bd 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -9,13 +9,13 @@ The tests should be launched on every supported device configuration we have. ### Type coverage New APIs for new way to define a kernel as a simple C++ function, where the kernel arguments are parameters to this function described by the extension can take only allowed kernel parameter type as specified in section [4.12.4 of the SYCL 2020 specification](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.parameter.passing) with exceptions to parameters of type `reducer` or `kernel_handler`. -Therefore those tests should be structured in a way that checks are performed on all allowed kernel parameter types to ensure that everything works correctly. +Therefore those tests should be structured in a way that checks are performed on set of types that satisfy the various requirements for kernel arguments to ensure that everything works correctly. ## Tests ### Unit tests -Tests in this category may not fully exercise the extension functionality, but are instead they are focused on making sure that all APIs are consistent with respect to existing APIs. +Tests in this category may not fully exercise the extension functionality, but are instead focused on making sure that all APIs are consistent with respect to other APIs. Perform tests on free function kernels requirements which should check that: @@ -36,41 +36,41 @@ Perform tests on free function kernel declaration with properties `nd_range_kern - that if a redeclaration of a function is decorated with the same property but with different arguments, the program should result in a compilation error. Perform tests on new traits for free function kernels which should check the following: - - that `is_nd_range_kernel_v` trait returns true if function declaration is decorated with `nd_range_kernel` property and false if it is not. + - that `is_nd_range_kernel_v` trait should be a subclass of `true_type` if function declaration is decorated with `nd_range_kernel` property and false if it is not. - - that `is_single_task_kernel_v` trait returns true function if declaration is decorated with `single_task_kernel` and false if it is not. + - that `is_single_task_kernel_v` trait should be a subclass of `true_type` if declaration is decorated with `single_task_kernel` and false if it is not. - that `is_kernel_v` trait returns true for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it returns false. Perform tests on new `kernel_bundle` member functions for free function kernels which should check that: - - the `get_kernel_id` member function returns a valid kernel identifier that is associated with that kernel. +- the `get_kernel_id` member function returns a valid kernel identifier that is associated with that kernel. - - the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. +- the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. - - the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. +- the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. - - the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. +- the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. - - the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. +- the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. - - the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. +- the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. - - the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. +- the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. - - the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. +- the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. - - the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. +- the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. - - the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. +- the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. - - the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. +- the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. - - the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. +- the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. - - the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. +- the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. -- that performs all the checks mentioned above on a free function kernel that is declared in one translation unit and defined in a separate one. +- that performs all the checks mentioned above on a free function kernel that is declared in one translation unit and defined in another. Perform tests on new free functions to query kernel information descriptors which should check the following: @@ -121,7 +121,7 @@ A series of checks should be performed that USM memory with three types of memor A series of checks should be performed that we can pass `id` where `Dimensions` is in {1, 2, 3} as kernel parameter to free function kernel and use it within kernel. #### Test `range` as kernel parameter: -A series of checks should be performed that we can pass `range` with different dimensions = 1, 2, 3 as kernel parameter to free function kernel and use it within kernel. +A series of checks should be performed that we can pass `range` where `Dimensions` is in {1, 2, 3} as kernel parameter to free function kernel and use it within kernel. #### Test `marray` when `T` is device copyable as kernel parameter: A series of checks should be performed that we can pass `marray` as kernel parameter to free function kernel and use it within kernel. From 1af15cc49780bcc964c6926cb8e23e534d486eb9 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 15 Apr 2025 21:37:00 +0200 Subject: [PATCH 20/24] fix --- .../test-e2e/FreeFunctionKernels/test-plan.md | 34 +++++++++---------- 1 file changed, 16 insertions(+), 18 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 177a2148fd5bd..2a6feb79ed23d 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -15,7 +15,7 @@ Therefore those tests should be structured in a way that checks are performed on ### Unit tests -Tests in this category may not fully exercise the extension functionality, but are instead focused on making sure that all APIs are consistent with respect to other APIs. +Tests in this category may not fully exercise the extension functionality, but are instead focused on making sure that all APIs are consistent with respect to existing APIs. Perform tests on free function kernels requirements which should check that: @@ -40,37 +40,35 @@ Perform tests on new traits for free function kernels which should check the fol - that `is_single_task_kernel_v` trait should be a subclass of `true_type` if declaration is decorated with `single_task_kernel` and false if it is not. -- that `is_kernel_v` trait returns true for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it returns false. +- that `is_kernel_v` trait should be a subclass of `true_type` for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it should be a subclass of `false_type`. -Perform tests on new `kernel_bundle` member functions for free function kernels which should check that: +Perform tests on new `kernel_bundle` member functions for free function kernel. These tests should declare `nd_range_kernel` and `single_task_kernel` free function kernels then check that: -- the `get_kernel_id` member function returns a valid kernel identifier that is associated with that kernel. + - the `get_kernel_id` member function returns a valid kernel identifier that is associated with free function kernel. -- the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. + - the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. -- the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. + - the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. -- the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. + - the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. -- the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. + - the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when declared free function kernel can be represented in a device image in the corresponding state and that free function kernel is compatible with at least one of the devices in `devs`. -- the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when free function kernel can be represented in a device image in the corresponding state and free function kernel is compatible with at least one of the devices in `devs`. + - the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. -- the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. + - the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. -- the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. + - the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. -- the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. + - the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. -- the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. + - the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. -- the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. + - the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. -- the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. - -- the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. + - the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. -- that performs all the checks mentioned above on a free function kernel that is declared in one translation unit and defined in another. +- Write tests that perform all the checks mentioned above on `nd_range_kernel` and `single_task_kernel` free functions, which are declared in one translation unit and defined in another. Perform tests on new free functions to query kernel information descriptors which should check the following: From b5aca8be5fa3828edc07d58462bf46cf64fb73ec Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 15 Apr 2025 21:52:13 +0200 Subject: [PATCH 21/24] fixed formatting --- .../test-e2e/FreeFunctionKernels/test-plan.md | 38 +++++++++---------- 1 file changed, 19 insertions(+), 19 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 2a6feb79ed23d..51c02244d86fa 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -18,7 +18,7 @@ Therefore those tests should be structured in a way that checks are performed on Tests in this category may not fully exercise the extension functionality, but are instead focused on making sure that all APIs are consistent with respect to existing APIs. -Perform tests on free function kernels requirements which should check that: +#### Perform tests on free function kernels requirements which should check that: - the compiler will emit diagnostic when free function kernel is declared with reference types as parameters. - the compiler will emit diagnostic when free function kernel is declared with variadic arguments. @@ -29,48 +29,48 @@ Perform tests on free function kernels requirements which should check that: - the compiler will emit diagnostic when a non-static member function is used as a kernel. Only static member function at class scope are allowed as free function kernel. -Perform tests on free function kernel declaration with properties `nd_range_kernel` and `single_task_kernel` which should check the following: +#### Perform tests on free function kernel declaration with properties `nd_range_kernel` and `single_task_kernel` which should check the following: - that if the property does not appear on the first declaration of the function in the translation unit, it will result in a compilation error. - that if the property appears on the first declaration of the function, and the following redeclarations do not have this property, it will not result in a compilation error. - that if a function is decorated with more than one of these properties, it will result in a compilation error. - that if a redeclaration of a function is decorated with the same property but with different arguments, the program should result in a compilation error. -Perform tests on new traits for free function kernels which should check the following: +#### Perform tests on new traits for free function kernels which should check the following: - that `is_nd_range_kernel_v` trait should be a subclass of `true_type` if function declaration is decorated with `nd_range_kernel` property and false if it is not. - that `is_single_task_kernel_v` trait should be a subclass of `true_type` if declaration is decorated with `single_task_kernel` and false if it is not. - that `is_kernel_v` trait should be a subclass of `true_type` for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it should be a subclass of `false_type`. -Perform tests on new `kernel_bundle` member functions for free function kernel. These tests should declare `nd_range_kernel` and `single_task_kernel` free function kernels then check that: +#### Perform tests on new `kernel_bundle` member functions for free function kernel. These tests should declare `nd_range_kernel` and `single_task_kernel` free function kernels then check that: - - the `get_kernel_id` member function returns a valid kernel identifier that is associated with free function kernel. +- the `get_kernel_id` member function returns a valid kernel identifier that is associated with free function kernel. - - the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. +- the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. - - the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. +- the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. - - the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. +- the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. - - the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when declared free function kernel can be represented in a device image in the corresponding state and that free function kernel is compatible with at least one of the devices in `devs`. +- the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when declared free function kernel can be represented in a device image in the corresponding state and that free function kernel is compatible with at least one of the devices in `devs`. - - the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. +- the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. - - the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. +- the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. - - the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. +- the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. - - the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. +- the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. - - the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. +- the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. - - the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. +- the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. - - the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. +- the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. -- Write tests that perform all the checks mentioned above on `nd_range_kernel` and `single_task_kernel` free functions, which are declared in one translation unit and defined in another. + Write test that perform all the checks mentioned above on `nd_range_kernel` and `single_task_kernel` free functions, which are declared in one translation unit and defined in another. -Perform tests on new free functions to query kernel information descriptors which should check the following: +#### Perform tests on new free functions to query kernel information descriptors which should check the following: - that `get_kernel_info(const context& ctxt)` produces the same result as would be computed by ``` @@ -94,7 +94,7 @@ Perform tests on new free functions to query kernel information descriptors whic auto ret = bundle.ext_oneapi_get_kernel().get_info(dev); ``` -Perform tests on use of illegal types for kernel paramters: +#### Perform tests on use of illegal types for kernel paramters: - that a class type `S` with a virtual base class of type `T` can not be used as free function kernel parameter type. - that a class type `S` with a virtual member function can not be used as free function kernel parameter type. From 49049f4ba20e22083ded15ce8465643f7fe090ef Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 16 Apr 2025 10:11:53 +0200 Subject: [PATCH 22/24] fix --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 51c02244d86fa..29b2143b6bf2f 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -42,13 +42,13 @@ Tests in this category may not fully exercise the extension functionality, but a - that `is_kernel_v` trait should be a subclass of `true_type` for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it should be a subclass of `false_type`. -#### Perform tests on new `kernel_bundle` member functions for free function kernel. These tests should declare `nd_range_kernel` and `single_task_kernel` free function kernels then check that: +#### Perform tests on new `kernel_bundle` member functions for free function kernels by declaring `nd_range_kernel` and `single_task_kernel` and verifying that: -- the `get_kernel_id` member function returns a valid kernel identifier that is associated with free function kernel. +- the `get_kernel_id` member function returns a valid kernel identifier associated with free function kernel. -- the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle in which the associated free function kernel can be found. - -- the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle in which the associated free function kernel can be found. +- the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle that contains the corresponding free function kernel. + +- the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle that contains the corresponding free function kernel. - the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. @@ -66,7 +66,7 @@ Tests in this category may not fully exercise the extension functionality, but a - the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. -- the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. +- the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. Write test that perform all the checks mentioned above on `nd_range_kernel` and `single_task_kernel` free functions, which are declared in one translation unit and defined in another. @@ -155,7 +155,7 @@ A series of checks should be performed that we can pass unscoped enumeration tha A series of checks should be performed that we can pass type aliases to allowed kernel paramater types as kernel parameter and use it within kernel. #### Interaction with additional kernel properties: -A series of checks should be performed to check that to the free function kernels may also be decorated with any of the properties defined in `sycl_ext_oneapi_kernel_properties`. This test should perform simple checks verifying if applied kernel_properties work within defined kernels. +A series of checks should be performed to check that to the free function kernels may also be decorated with the properties defined in `sycl_ext_oneapi_kernel_properties`. This test should perform simple checks verifying if applied kernel_properties work within defined kernels. #### Free function kernels compatibility with L0 backend: A series of checks should be performed to check compatibility of free function kernels with Level Zero Backend without going through the SYCL host runtime. From a121cb0d168d70940755e70469105bc6548694ac Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Apr 2025 16:38:54 +0200 Subject: [PATCH 23/24] Applied changes based on feedback --- .../test-e2e/FreeFunctionKernels/test-plan.md | 23 +++---------------- 1 file changed, 3 insertions(+), 20 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 29b2143b6bf2f..6920da69d50fa 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -20,13 +20,9 @@ Tests in this category may not fully exercise the extension functionality, but a #### Perform tests on free function kernels requirements which should check that: - the compiler will emit diagnostic when free function kernel is declared with reference types as parameters. - - the compiler will emit diagnostic when free function kernel is declared with variadic arguments. - - the compiler will emit diagnostic when free function kernel provides default parameter values. - - the compiler will emit diagnostic when free function kernel return type is not `void`. - - the compiler will emit diagnostic when a non-static member function is used as a kernel. Only static member function at class scope are allowed as free function kernel. #### Perform tests on free function kernel declaration with properties `nd_range_kernel` and `single_task_kernel` which should check the following: @@ -36,36 +32,23 @@ Tests in this category may not fully exercise the extension functionality, but a - that if a redeclaration of a function is decorated with the same property but with different arguments, the program should result in a compilation error. #### Perform tests on new traits for free function kernels which should check the following: - - that `is_nd_range_kernel_v` trait should be a subclass of `true_type` if function declaration is decorated with `nd_range_kernel` property and false if it is not. - - - that `is_single_task_kernel_v` trait should be a subclass of `true_type` if declaration is decorated with `single_task_kernel` and false if it is not. - +- that `is_nd_range_kernel_v` trait should be a subclass of `true_type` if function declaration is decorated with `nd_range_kernel` property or a subclass of `false_type` if it is not. +- that `is_single_task_kernel_v` trait should be a subclass of `true_type` if declaration is decorated with `single_task_kernel` or a subclass of `false_type` if it is not. - that `is_kernel_v` trait should be a subclass of `true_type` for function whose declaration is decorated with either the `nd_range_kernel` property or the `single_task_kernel` property when it is not then it should be a subclass of `false_type`. #### Perform tests on new `kernel_bundle` member functions for free function kernels by declaring `nd_range_kernel` and `single_task_kernel` and verifying that: - the `get_kernel_id` member function returns a valid kernel identifier associated with free function kernel. - - the `get_kernel_bundle(const context& ctxt)` member function returns a kernel bundle that contains the corresponding free function kernel. - - the `get_kernel_bundle(const context& ctxt, const std::vector& devs)` member function returns a kernel bundle that contains the corresponding free function kernel. - - the `has_kernel_bundle(const context& ctxt)` returns true when a free function kernel can be represented in a device image in the corresponding state and the associated free function kernel is compatible with at least one of the devices in `ctxt`. - - the `has_kernel_bundle(const context& ctxt, const std::vector& devs)` returns true when declared free function kernel can be represented in a device image in the corresponding state and that free function kernel is compatible with at least one of the devices in `devs`. - - the `is_compatible(const device& dev)` returns true when the associated free function kernel is compatible with `dev`. - - the `ext_oneapi_has_kernel()` returns true only if the kernel bundle contains the associated free function kernel. - - the `ext_oneapi_has_kernel(const device &dev)` returns true when kernel bundle contains the associated free function kernel and if that kernel is compatible with `dev`. - - the `ext_oneapi_get_kernel` returns the kernel object representing that kernel if the free function kernel resides in this kernel bundle. - - the `ext_oneapi_get_kernel` throws exception with the error code `errc::invalid` if the associated free function kernel does not reside in this kernel bundle. - - the `get_kernel_ids()` returns all of the kernels defined in the source, whether they were defined as free function kernels, lambda expressions or named kernel objects. - - the `info::kernel::num_args` returns the number of parameters in the function definition of the associated free function kernel. Write test that perform all the checks mentioned above on `nd_range_kernel` and `single_task_kernel` free functions, which are declared in one translation unit and defined in another. @@ -104,7 +87,7 @@ Tests in this category perform some meaningful actions with the extension to see that the extension works in a scenarios which mimic real-life usage of the extension. -With the exception of the `single_task_kernel` free function kernels, all subsequent tests are executed with `Dimensions` $$\in {1, 2, 3}$$. +With the exception of the `single_task_kernel` free function kernels, all subsequent tests are executed with `Dimensions` $$\in \{1, 2, 3\}$$. In all subsequent tests, free function kernels should be declared within a namespace, as static member functions of a class, or in the global namespace. From 6e6f3452b125fd2b1d88cdfb8b5d53e125b9c4f7 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Apr 2025 16:42:19 +0200 Subject: [PATCH 24/24] Fix --- sycl/test-e2e/FreeFunctionKernels/test-plan.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/test-plan.md b/sycl/test-e2e/FreeFunctionKernels/test-plan.md index 6920da69d50fa..1dc2aa079768d 100644 --- a/sycl/test-e2e/FreeFunctionKernels/test-plan.md +++ b/sycl/test-e2e/FreeFunctionKernels/test-plan.md @@ -99,10 +99,10 @@ templated with `target::device`, inside free function kernel when passed as kern A series of checks should be performed that USM memory with three types of memory allocations `host`, `device` and `shared` is supported inside free function kernel when passed as kernel parameter. #### Test `id` as kernel parameter: -A series of checks should be performed that we can pass `id` where `Dimensions` is in {1, 2, 3} as kernel parameter to free function kernel and use it within kernel. +A series of checks should be performed that we can pass `id` where `Dimensions` is in $$\in \{1, 2, 3\}$$ as kernel parameter to free function kernel and use it within kernel. #### Test `range` as kernel parameter: -A series of checks should be performed that we can pass `range` where `Dimensions` is in {1, 2, 3} as kernel parameter to free function kernel and use it within kernel. +A series of checks should be performed that we can pass `range` where `Dimensions` is in $$\in \{1, 2, 3\}$$ as kernel parameter to free function kernel and use it within kernel. #### Test `marray` when `T` is device copyable as kernel parameter: A series of checks should be performed that we can pass `marray` as kernel parameter to free function kernel and use it within kernel.