diff --git a/test_conformance/api/test_kernel_attributes.cpp b/test_conformance/api/test_kernel_attributes.cpp index 86b3595c..dd50e0f9 100644 --- a/test_conformance/api/test_kernel_attributes.cpp +++ b/test_conformance/api/test_kernel_attributes.cpp @@ -385,14 +385,14 @@ REGISTER_TEST(null_required_work_group_size) struct KernelAttribInfo { - std::string str; - cl_uint max_dim; + cl_int wgs[3]; + cl_uint min_dim; }; std::vector attribs; - attribs.push_back({ "__attribute__((reqd_work_group_size(2,1,1)))", 1 }); - attribs.push_back({ "__attribute__((reqd_work_group_size(2,3,1)))", 2 }); - attribs.push_back({ "__attribute__((reqd_work_group_size(2,3,4)))", 3 }); + attribs.push_back({ { 2, 1, 1 }, 1 }); + attribs.push_back({ { 2, 3, 1 }, 2 }); + attribs.push_back({ { 2, 3, 4 }, 3 }); const std::string body_str = R"( __kernel void wg_size(__global int* dst) @@ -409,7 +409,11 @@ REGISTER_TEST(null_required_work_group_size) for (auto& attrib : attribs) { - const std::string source_str = attrib.str + body_str; + const std::string attrib_str = "__attribute__((reqd_work_group_size(" + + std::to_string(attrib.wgs[0]) + "," + + std::to_string(attrib.wgs[1]) + "," + + std::to_string(attrib.wgs[2]) + ")))"; + const std::string source_str = attrib_str + body_str; const char* source = source_str.c_str(); clProgramWrapper program; @@ -421,21 +425,19 @@ REGISTER_TEST(null_required_work_group_size) error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &dst); test_error(error, "clSetKernelArg failed"); - for (cl_uint work_dim = 1; work_dim <= attrib.max_dim; work_dim++) + for (cl_uint work_dim = attrib.min_dim; work_dim <= 3; work_dim++) { - const cl_int expected[3] = { 2, work_dim >= 2 ? 3 : 1, - work_dim >= 3 ? 4 : 1 }; const size_t test_work_group_size = - expected[0] * expected[1] * expected[2]; - if ((size_t)expected[0] > device_max_work_item_sizes[0] - || (size_t)expected[1] > device_max_work_item_sizes[1] - || (size_t)expected[2] > device_max_work_item_sizes[2] + attrib.wgs[0] * attrib.wgs[1] * attrib.wgs[2]; + if ((size_t)attrib.wgs[0] > device_max_work_item_sizes[0] + || (size_t)attrib.wgs[1] > device_max_work_item_sizes[1] + || (size_t)attrib.wgs[2] > device_max_work_item_sizes[2] || test_work_group_size > device_max_work_group_size) { log_info("Skipping test for work_dim = %u: required work group " "size (%i, %i, %i) (total %zu) exceeds device max " "work group size (%zu, %zu, %zu) (total %zu)\n", - work_dim, expected[0], expected[1], expected[2], + work_dim, attrib.wgs[0], attrib.wgs[1], attrib.wgs[2], test_work_group_size, device_max_work_item_sizes[0], device_max_work_item_sizes[1], device_max_work_item_sizes[2], @@ -444,8 +446,9 @@ REGISTER_TEST(null_required_work_group_size) } const cl_int zero = 0; - error = clEnqueueFillBuffer(queue, dst, &zero, sizeof(zero), 0, - sizeof(expected), 0, nullptr, nullptr); + error = + clEnqueueFillBuffer(queue, dst, &zero, sizeof(zero), 0, + sizeof(attrib.wgs), 0, nullptr, nullptr); const size_t global_work_size[3] = { 2 * 32, 3 * 32, 4 * 32 }; error = clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, @@ -458,12 +461,12 @@ REGISTER_TEST(null_required_work_group_size) results, 0, nullptr, nullptr); test_error(error, "clEnqueueReadBuffer failed"); - if (results[0] != expected[0] || results[1] != expected[1] - || results[2] != expected[2]) + if (results[0] != attrib.wgs[0] || results[1] != attrib.wgs[1] + || results[2] != attrib.wgs[2]) { log_error("Executed local size mismatch with work_dim = %u: " "Expected (%d,%d,%d) got (%d,%d,%d)\n", - work_dim, expected[0], expected[1], expected[2], + work_dim, attrib.wgs[0], attrib.wgs[1], attrib.wgs[2], results[0], results[1], results[2]); return TEST_FAIL; } @@ -477,15 +480,15 @@ REGISTER_TEST(null_required_work_group_size) test_error(error, "clGetKernelSuggestedLocalWorkSizeKHR failed"); - if ((cl_int)suggested[0] != expected[0] - || (cl_int)suggested[1] != expected[1] - || (cl_int)suggested[2] != expected[2]) + if (suggested[0] != (size_t)attrib.wgs[0] + || suggested[1] != (size_t)attrib.wgs[1] + || suggested[2] != (size_t)attrib.wgs[2]) { log_error("Suggested local size mismatch with work_dim = " - "%u: Expected (%d,%d,%d) got (%d,%d,%d)\n", - work_dim, expected[0], expected[1], expected[2], - (cl_int)suggested[0], (cl_int)suggested[1], - (cl_int)suggested[2]); + "%u: Expected (%d,%d,%d) got (%zu,%zu,%zu)\n", + work_dim, attrib.wgs[0], attrib.wgs[1], + attrib.wgs[2], suggested[0], suggested[1], + suggested[2]); return TEST_FAIL; } }