ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

fix the null required work group size test (#2576)

Browse Source 
A 3D required work-group size is always valid, but a 1D or 2D required
work-group size is only valid when the work-group size in those
dimensions is equal to one.

fixes #2575
This commit is contained in:
Ben Ashbaugh
2025-11-26 16:10:47 -08:00
committed by GitHub
parent 02e99f4554
commit afc7e64c57
1 changed files with 29 additions and 26 deletions
Download Patch File Download Diff File Expand all files Collapse all files

55
test_conformance/api/test_kernel_attributes.cpp
View File

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