44
55// The name mangling for free function kernels currently does not work with PTX.
66// UNSUPPORTED: cuda
7+ // UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends.
78
89#include < sycl/detail/core.hpp>
910#include < sycl/ext/oneapi/free_function_queries.hpp>
@@ -26,13 +27,13 @@ namespace free_functions::tests {
2627SYCL_EXT_ONEAPI_FUNCTION_PROPERTY ((syclexp::nd_range_kernel<1 >))
2728void function_in_ns (float start, float *ptr) {
2829 size_t id = syclext::this_work_item::get_nd_item<1 >().get_global_linear_id ();
29- ptr[id] = start + static_cast <float >(id + 1 );
30+ ptr[id] = start + static_cast <float >(id);
3031}
3132
3233SYCL_EXT_ONEAPI_FUNCTION_PROPERTY ((syclexp::nd_range_kernel<1 >))
3334void func (float start, float *ptr) {
3435 size_t id = syclext::this_work_item::get_nd_item<1 >().get_global_linear_id ();
35- ptr[id] = start + static_cast <float >(id + 2 );
36+ ptr[id] = start + static_cast <float >(id);
3637}
3738} // namespace free_functions::tests
3839
@@ -41,7 +42,7 @@ inline namespace V1 {
4142SYCL_EXT_ONEAPI_FUNCTION_PROPERTY ((syclexp::nd_range_kernel<1 >))
4243void function_in_inline_ns (float start, float *ptr) {
4344 size_t id = syclext::this_work_item::get_nd_item<1 >().get_global_linear_id ();
44- ptr[id] = start + static_cast <float >(id + 2 );
45+ ptr[id] = start + static_cast <float >(id);
4546}
4647} // namespace V1
4748} // namespace free_functions::tests
@@ -50,7 +51,7 @@ namespace {
5051SYCL_EXT_ONEAPI_FUNCTION_PROPERTY ((syclexp::nd_range_kernel<1 >))
5152void function_in_anonymous_ns (float start, float *ptr) {
5253 size_t id = syclext::this_work_item::get_nd_item<1 >().get_global_linear_id ();
53- ptr[id] = start + static_cast <float >(id + 3 );
54+ ptr[id] = start + static_cast <float >(id);
5455}
5556} // namespace
5657
@@ -61,6 +62,12 @@ static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) {
6162 sycl::nd_range ndr{{NUM}, {WGSIZE}};
6263 cgh.parallel_for (ndr, kernel);
6364 }).wait ();
65+ // Check the result
66+ for (size_t i = 0 ; i < NUM; ++i) {
67+ const float expected = 3 .14f + static_cast <float >(i);
68+ assert (ptr[i] == expected &&
69+ " Kernel execution did not produce the expected result" );
70+ }
6471}
6572
6673void test_function_without_ns (sycl::queue &q, sycl::context &ctxt) {
0 commit comments