Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix test_sampler #85

Draft
wants to merge 1 commit into
base: master
Choose a base branch
from
Draft

Fix test_sampler #85

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 6 additions & 6 deletions conformance_tests/core/test_sampler/src/test_sampler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion utils/test_harness/include/test_harness/test_harness_image.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down
16 changes: 14 additions & 2 deletions utils/test_harness/include/test_harness/test_harness_module.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand All @@ -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<FunctionArg> &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<FunctionArg> &args,
bool is_immediate);

Expand Down
3 changes: 1 addition & 2 deletions utils/test_harness/src/test_harness_image.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
76 changes: 71 additions & 5 deletions utils/test_harness/src/test_harness_module.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<FunctionArg> 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<FunctionArg> &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(
Expand Down Expand Up @@ -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<FunctionArg> 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<FunctionArg> &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;
Expand Down