diff --git a/sycl/test-e2e/FreeFunctionKernels/properties.cpp b/sycl/test-e2e/FreeFunctionKernels/properties.cpp new file mode 100644 index 0000000000000..1b4a05a7c8253 --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/properties.cpp @@ -0,0 +1,180 @@ +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// UNSUPPORTED: cuda, hip +// UNSUPPORTED-INTENDED: Device incompatible error + +#include +#include +#include +#include + +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 32; +static constexpr size_t SGSIZE = 16; + +inline void kernel_code(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_wg_1dsize_before(float start, float *ptr) { + kernel_code(start, ptr); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +void range_wg_1dsize_after(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size_hint)) +void range_wg_1dsize_hint_after(float start, float *ptr) { + kernel_code(start, ptr); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size_hint)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_wg_1dsize_hint_before(float start, float *ptr) { + kernel_code(start, ptr); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_sg_1dsize_before(float start, float *ptr) { + kernel_code(start, ptr); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +void range_sg_1dsize_after(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::device_has)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_has_before(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::device_has)) +void range_has_after(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +void range_several_after(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_several_before(float start, float *ptr) { kernel_code(start, ptr); } + +template bool check_result(T *ptr) { + for (size_t i = 0; i < NUM; ++i) { + const T expected = static_cast(3.14f) + static_cast(i); + if (ptr[i] != expected) + return true; + } + return false; +} + +template +static bool call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { + T *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + cgh.set_args(3.14f, ptr); + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, kernel); + }).wait(); + const bool ret = check_result(ptr); + sycl::free(ptr, q); + return ret; +} + +template +bool test(sycl::queue &q, sycl::context &ctxt, std::string_view name) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + bool ret = call_kernel_code(q, k_func); + auto value = syclexp::get_kernel_info(q); + const auto kernel_ids = exe_bndl.get_kernel_ids(); + if (kernel_ids.empty()) + return true; + sycl::kernel k = exe_bndl.get_kernel(kernel_ids[0]); + const size_t kernel_value = k.get_info(q.get_device()); + ret |= (value != kernel_value); + if (ret) + std::cout << "Test " << name + << " did not pass: value got from get_kernel_info " << value + << ", value got from kernel get_info " << kernel_value + << std::endl; + return ret; +} + +template +bool test_has_desc(sycl::queue &q, sycl::context &ctxt) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + return call_kernel_code(q, k_func); +} + +using wg_size_desc = sycl::info::kernel_device_specific::work_group_size; +using sg_size_desc = sycl::info::kernel_device_specific::compile_sub_group_size; + +template +bool test_several_properties(sycl::queue &q, sycl::context &ctxt, + std::string_view name) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + bool ret = call_kernel_code(q, k_func); + const size_t value_wg_size = syclexp::get_kernel_info(q); + const size_t value_sg_size = syclexp::get_kernel_info(q); + const auto kernel_ids = exe_bndl.get_kernel_ids(); + if (kernel_ids.empty()) + return true; + sycl::kernel k = exe_bndl.get_kernel(kernel_ids[0]); + const size_t kernel_value_wg = k.get_info(q.get_device()); + const size_t kernel_value_sg = k.get_info(q.get_device()); + ret |= (value_wg_size != kernel_value_wg); + ret |= (value_sg_size != kernel_value_sg); + if (ret) + std::cout << "Test " << name << " did not pass: value_wg_size " + << value_wg_size << ", value_sg_size " << value_sg_size + << ", kernel_value_wg " << kernel_value_wg << ", kernel_value_sg " + << kernel_value_sg << std::endl; + return ret; +} + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); + + int ret = 0; + ret |= test( + q, ctxt, "range_wg_1dsize_before"); + ret |= test( + q, ctxt, "range_wg_1dsize_after"); + ret |= test( + q, ctxt, "range_wg_1dsize_hint_before"); + ret |= test( + q, ctxt, "range_wg_1dsize_hint_after"); + ret |= test( + q, ctxt, "range_sg_1dsize_before"); + ret |= test( + q, ctxt, "range_sg_1dsize_after"); + ret |= test_has_desc(q, ctxt); + ret |= test_has_desc(q, ctxt); + ret |= test_several_properties( + q, ctxt, "range_several_before"); + ret |= test_several_properties( + q, ctxt, "range_several_after"); + return ret; +}