From 5581ece98d09d0396e58ec4ee2621e1050d1d12f Mon Sep 17 00:00:00 2001 From: "Grau, Kacper" Date: Thu, 10 Oct 2024 18:50:56 +0200 Subject: [PATCH] Fix test_sampler * Fixed failing test case * Added 2d version of create_and_execute_function Signed-off-by: Kacper Grau --- .../core/test_sampler/src/test_sampler.cpp | 12 +-- .../test_harness/test_harness_image.hpp | 2 +- .../test_harness/test_harness_module.hpp | 16 +++- utils/test_harness/src/test_harness_image.cpp | 3 +- .../test_harness/src/test_harness_module.cpp | 76 +++++++++++++++++-- 5 files changed, 93 insertions(+), 16 deletions(-) diff --git a/conformance_tests/core/test_sampler/src/test_sampler.cpp b/conformance_tests/core/test_sampler/src/test_sampler.cpp index 99f452fe..4892a16e 100644 --- a/conformance_tests/core/test_sampler/src/test_sampler.cpp +++ b/conformance_tests/core/test_sampler/src/test_sampler.cpp @@ -118,12 +118,12 @@ static ze_image_handle_t create_sampler_image(lzt::ImagePNG32Bit png_image, int height, int width) { ze_image_desc_t image_description = {}; image_description.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; - image_description.format.layout = ZE_IMAGE_FORMAT_LAYOUT_32; + image_description.format.layout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; image_description.pNext = nullptr; image_description.flags = ZE_IMAGE_FLAG_KERNEL_WRITE; image_description.type = ZE_IMAGE_TYPE_2D; - image_description.format.type = ZE_IMAGE_FORMAT_TYPE_UINT; + image_description.format.type = ZE_IMAGE_FORMAT_TYPE_UNORM; image_description.format.x = ZE_IMAGE_FORMAT_SWIZZLE_R; image_description.format.y = ZE_IMAGE_FORMAT_SWIZZLE_G; image_description.format.z = ZE_IMAGE_FORMAT_SWIZZLE_B; @@ -242,12 +242,12 @@ TEST_P( args_inhost.push_back(arg); lzt::create_and_execute_function(lzt::zeDevice::get_instance()->get_device(), - module, func_name_inhost, 1, args_inhost, - is_immediate); + module, func_name_inhost, output_width, + output_height, args_inhost, is_immediate); lzt::create_and_execute_function(lzt::zeDevice::get_instance()->get_device(), - module, func_name_inkernel, 1, args_inkernel, - is_immediate); + module, func_name_inkernel, output_width, + output_height, args_inkernel, is_immediate); lzt::copy_image_to_mem(output_xeimage_host, output_inhost); lzt::copy_image_to_mem(output_xeimage_kernel, output_inkernel); diff --git a/utils/test_harness/include/test_harness/test_harness_image.hpp b/utils/test_harness/include/test_harness/test_harness_image.hpp index d157eba6..c333c62c 100644 --- a/utils/test_harness/include/test_harness/test_harness_image.hpp +++ b/utils/test_harness/include/test_harness/test_harness_image.hpp @@ -108,7 +108,7 @@ void destroy_ze_image(ze_image_handle_t image); ze_image_properties_t get_ze_image_properties(ze_image_desc_t image_descriptor); void copy_image_from_mem(lzt::ImagePNG32Bit input, ze_image_handle_t output); -void copy_image_to_mem(ze_image_handle_t input, lzt::ImagePNG32Bit output); +void copy_image_to_mem(ze_image_handle_t input, lzt::ImagePNG32Bit &output); class zeImageCreateCommon { public: diff --git a/utils/test_harness/include/test_harness/test_harness_module.hpp b/utils/test_harness/include/test_harness/test_harness_module.hpp index 3e5b905b..1d6152d4 100644 --- a/utils/test_harness/include/test_harness/test_harness_module.hpp +++ b/utils/test_harness/include/test_harness/test_harness_module.hpp @@ -73,8 +73,13 @@ ze_kernel_properties_t get_kernel_properties(ze_kernel_handle_t kernel); // This function is useful when only a single argument is needed. void create_and_execute_function(ze_device_handle_t device, ze_module_handle_t module, - std::string func_name, int group_size, + std::string func_name, uint32_t group_size_x, void *arg, bool is_immediate); + void create_and_execute_function(ze_device_handle_t device, + ze_module_handle_t module, + std::string func_name, uint32_t group_size_x, + uint32_t group_size_y, void *arg, + bool is_immediate); void kernel_set_indirect_access(ze_kernel_handle_t hKernel, ze_kernel_indirect_access_flags_t flags); void kernel_get_indirect_access(ze_kernel_handle_t hKernel, @@ -91,7 +96,14 @@ struct FunctionArg { // Accepts arbitrary amounts of function arguments void create_and_execute_function(ze_device_handle_t device, ze_module_handle_t module, - std::string func_name, int group_size, + std::string func_name, uint32_t group_size_x, + const std::vector &args, + bool is_immediate); + +void create_and_execute_function(ze_device_handle_t device, + ze_module_handle_t module, + std::string func_name, uint32_t group_size_x, + uint32_t group_size_y, const std::vector &args, bool is_immediate); diff --git a/utils/test_harness/src/test_harness_image.cpp b/utils/test_harness/src/test_harness_image.cpp index 75150ab2..0d3a6e17 100644 --- a/utils/test_harness/src/test_harness_image.cpp +++ b/utils/test_harness/src/test_harness_image.cpp @@ -50,8 +50,7 @@ void copy_image_from_mem(lzt::ImagePNG32Bit input, ze_image_handle_t output) { lzt::destroy_command_list(command_list); } -void copy_image_to_mem(ze_image_handle_t input, lzt::ImagePNG32Bit output) { - +void copy_image_to_mem(ze_image_handle_t input, lzt::ImagePNG32Bit &output) { auto command_list = lzt::create_command_list(); EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListAppendImageCopyToMemory( command_list, output.raw_data(), input, diff --git a/utils/test_harness/src/test_harness_module.cpp b/utils/test_harness/src/test_harness_module.cpp index b10db98b..96aa2ed6 100644 --- a/utils/test_harness/src/test_harness_module.cpp +++ b/utils/test_harness/src/test_harness_module.cpp @@ -260,30 +260,29 @@ ze_kernel_properties_t get_kernel_properties(ze_kernel_handle_t kernel) { // Expand as needed. void create_and_execute_function(ze_device_handle_t device, ze_module_handle_t module, - std::string func_name, int group_size, + std::string func_name, uint32_t group_size_x, void *arg, bool is_immediate) { std::vector args; if (arg != nullptr) { FunctionArg func_arg{sizeof(arg), &arg}; args.push_back(func_arg); } - create_and_execute_function(device, module, func_name, group_size, args, + create_and_execute_function(device, module, func_name, group_size_x, args, is_immediate); } void create_and_execute_function(ze_device_handle_t device, ze_module_handle_t module, - std::string func_name, int group_size, + std::string func_name, uint32_t group_size_x, const std::vector &args, bool is_immediate) { ze_kernel_handle_t function = create_function(module, func_name); zeCommandBundle cmd_bundle = create_command_bundle(device, is_immediate); - uint32_t group_size_x = group_size; uint32_t group_size_y = 1; uint32_t group_size_z = 1; EXPECT_EQ(ZE_RESULT_SUCCESS, - zeKernelSuggestGroupSize(function, group_size, 1, 1, &group_size_x, + zeKernelSuggestGroupSize(function, group_size_x, 1, 1, &group_size_x, &group_size_y, &group_size_z)); EXPECT_EQ( @@ -330,6 +329,73 @@ void create_and_execute_function(ze_device_handle_t device, destroy_command_bundle(cmd_bundle); } +void create_and_execute_function(ze_device_handle_t device, + ze_module_handle_t module, + std::string func_name, uint32_t group_size_x, + uint32_t group_size_y, void *arg, + bool is_immediate) { + std::vector args; + if (arg != nullptr) { + FunctionArg func_arg{sizeof(arg), &arg}; + args.push_back(func_arg); + } + create_and_execute_function(device, module, func_name, group_size_x, group_size_y, args, is_immediate); +} + +void create_and_execute_function(ze_device_handle_t device, + ze_module_handle_t module, + std::string func_name, uint32_t group_size_x, + uint32_t group_size_y, const std::vector &args, + bool is_immediate) { + + ze_kernel_handle_t function = create_function(module, func_name); + zeCommandBundle cmd_bundle = create_command_bundle(device, is_immediate); + + uint32_t group_size_z = 1; + EXPECT_EQ(ZE_RESULT_SUCCESS, + zeKernelSuggestGroupSize(function, group_size_x, group_size_y, group_size_z, &group_size_x, &group_size_y, &group_size_z)); + + EXPECT_EQ( + ZE_RESULT_SUCCESS, + zeKernelSetGroupSize(function, group_size_x, group_size_y, group_size_z)); + + ze_kernel_properties_t function_properties = get_kernel_properties(function); + EXPECT_EQ(function_properties.numKernelArgs, args.size()); + + int i = 0; + for (auto arg : args) { + EXPECT_EQ( + ZE_RESULT_SUCCESS, + zeKernelSetArgumentValue(function, i++, arg.arg_size, arg.arg_value)); + } + + ze_group_count_t thread_group_dimensions; + thread_group_dimensions.groupCountX = group_size_x; + thread_group_dimensions.groupCountY = group_size_y; + thread_group_dimensions.groupCountZ = 1; + + EXPECT_EQ(ZE_RESULT_SUCCESS, + zeCommandListAppendLaunchKernel(cmd_bundle.list, function, + &thread_group_dimensions, nullptr, + 0, nullptr)); + + EXPECT_EQ(ZE_RESULT_SUCCESS, + zeCommandListAppendBarrier(cmd_bundle.list, nullptr, 0, nullptr)); + if (is_immediate) { + EXPECT_EQ(ZE_RESULT_SUCCESS, + zeCommandListHostSynchronize(cmd_bundle.list, UINT64_MAX)); + } else { + EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListClose(cmd_bundle.list)); + + EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandQueueExecuteCommandLists(cmd_bundle.queue, 1, &cmd_bundle.list, nullptr)); + + EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandQueueSynchronize(cmd_bundle.queue, UINT64_MAX)); + } + + destroy_function(function); + destroy_command_bundle(cmd_bundle); +} + void kernel_set_indirect_access(ze_kernel_handle_t hKernel, ze_kernel_indirect_access_flags_t flags) { auto kernel_initial = hKernel;