mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
add a test case for subbuffer property queries (#961)
This commit is contained in:
Ben Ashbaugh
@@ -22,16 +22,18 @@
|
||||
typedef enum
|
||||
{
|
||||
image,
|
||||
buffer
|
||||
} mem_obj_type;
|
||||
image_with_properties,
|
||||
buffer,
|
||||
buffer_with_properties,
|
||||
subbuffer,
|
||||
} test_type;
|
||||
|
||||
struct test_data
|
||||
{
|
||||
mem_obj_type obj_t;
|
||||
test_type type;
|
||||
std::vector<cl_mem_properties> properties;
|
||||
std::string description;
|
||||
std::string src;
|
||||
std::string kernel_name;
|
||||
cl_kernel kernel;
|
||||
};
|
||||
|
||||
static int create_object_and_check_properties(cl_context context,
|
||||
@@ -43,47 +45,94 @@ static int create_object_and_check_properties(cl_context context,
|
||||
{
|
||||
cl_int error = CL_SUCCESS;
|
||||
|
||||
if (test_case.obj_t == image)
|
||||
switch (test_case.type)
|
||||
{
|
||||
cl_image_format format;
|
||||
format.image_channel_order = CL_RGBA;
|
||||
format.image_channel_data_type = CL_UNSIGNED_INT32;
|
||||
cl_image_desc desc;
|
||||
memset(&desc, 0x0, sizeof(cl_image_desc));
|
||||
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
|
||||
desc.image_width = size_x;
|
||||
desc.image_height = size_y;
|
||||
case image: {
|
||||
cl_image_format format = { 0 };
|
||||
format.image_channel_order = CL_RGBA;
|
||||
format.image_channel_data_type = CL_UNSIGNED_INT32;
|
||||
test_object = clCreateImage2D(context, flags, &format, size_x,
|
||||
size_y, 0, local_data.data(), &error);
|
||||
test_error(error, "clCreateImage2D failed");
|
||||
}
|
||||
break;
|
||||
case image_with_properties: {
|
||||
cl_image_format format = { 0 };
|
||||
format.image_channel_order = CL_RGBA;
|
||||
format.image_channel_data_type = CL_UNSIGNED_INT32;
|
||||
cl_image_desc desc = { 0 };
|
||||
desc.image_type = CL_MEM_OBJECT_IMAGE2D;
|
||||
desc.image_width = size_x;
|
||||
desc.image_height = size_y;
|
||||
|
||||
if (test_case.properties.size() == 0)
|
||||
{
|
||||
test_object =
|
||||
clCreateImageWithProperties(context, NULL, flags, &format,
|
||||
&desc, local_data.data(), &error);
|
||||
if (test_case.properties.size() == 0)
|
||||
{
|
||||
test_object = clCreateImageWithProperties(
|
||||
context, NULL, flags, &format, &desc, local_data.data(),
|
||||
&error);
|
||||
}
|
||||
else
|
||||
{
|
||||
test_object = clCreateImageWithProperties(
|
||||
context, test_case.properties.data(), flags, &format, &desc,
|
||||
local_data.data(), &error);
|
||||
}
|
||||
test_error(error, "clCreateImageWithProperties failed");
|
||||
}
|
||||
else
|
||||
{
|
||||
test_object = clCreateImageWithProperties(
|
||||
context, test_case.properties.data(), flags, &format, &desc,
|
||||
local_data.data(), &error);
|
||||
break;
|
||||
case buffer: {
|
||||
test_object = clCreateBuffer(context, flags,
|
||||
local_data.size() * sizeof(cl_uint),
|
||||
local_data.data(), &error);
|
||||
test_error(error, "clCreateBuffer failed");
|
||||
}
|
||||
test_error(error, "clCreateImageWithProperties failed");
|
||||
}
|
||||
if (test_case.obj_t == buffer)
|
||||
{
|
||||
if (test_case.properties.size() == 0)
|
||||
{
|
||||
test_object = clCreateBufferWithProperties(
|
||||
context, NULL, flags, local_data.size() * sizeof(cl_uint),
|
||||
local_data.data(), &error);
|
||||
}
|
||||
else
|
||||
{
|
||||
test_object = clCreateBufferWithProperties(
|
||||
context, test_case.properties.data(), flags,
|
||||
local_data.size() * sizeof(cl_uint), local_data.data(), &error);
|
||||
case buffer_with_properties: {
|
||||
if (test_case.properties.size() == 0)
|
||||
{
|
||||
test_object = clCreateBufferWithProperties(
|
||||
context, NULL, flags, local_data.size() * sizeof(cl_uint),
|
||||
local_data.data(), &error);
|
||||
}
|
||||
else
|
||||
{
|
||||
test_object = clCreateBufferWithProperties(
|
||||
context, test_case.properties.data(), flags,
|
||||
local_data.size() * sizeof(cl_uint), local_data.data(),
|
||||
&error);
|
||||
}
|
||||
test_error(error, "clCreateBufferWithProperties failed.");
|
||||
}
|
||||
break;
|
||||
case subbuffer: {
|
||||
clMemWrapper parent_object;
|
||||
if (test_case.properties.size() == 0)
|
||||
{
|
||||
parent_object = clCreateBufferWithProperties(
|
||||
context, NULL, flags, local_data.size() * sizeof(cl_uint),
|
||||
local_data.data(), &error);
|
||||
}
|
||||
else
|
||||
{
|
||||
parent_object = clCreateBufferWithProperties(
|
||||
context, test_case.properties.data(), flags,
|
||||
local_data.size() * sizeof(cl_uint), local_data.data(),
|
||||
&error);
|
||||
}
|
||||
test_error(error, "clCreateBufferWithProperties failed.");
|
||||
|
||||
test_error(error, "clCreateBufferWithProperties failed.");
|
||||
cl_mem_flags subbuffer_flags = flags
|
||||
& (CL_MEM_READ_WRITE | CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY);
|
||||
|
||||
cl_buffer_region region = { 0 };
|
||||
region.origin = 0;
|
||||
region.size = local_data.size() * sizeof(cl_uint);
|
||||
test_object = clCreateSubBuffer(parent_object, subbuffer_flags,
|
||||
CL_BUFFER_CREATE_TYPE_REGION,
|
||||
®ion, &error);
|
||||
test_error(error, "clCreateSubBuffer failed.");
|
||||
}
|
||||
break;
|
||||
default: log_error("Unknown test type!"); return TEST_FAIL;
|
||||
}
|
||||
|
||||
std::vector<cl_mem_properties> check_properties;
|
||||
@@ -94,6 +143,22 @@ static int create_object_and_check_properties(cl_context context,
|
||||
test_error(error,
|
||||
"clGetMemObjectInfo failed asking for CL_MEM_PROPERTIES size.");
|
||||
|
||||
// Buffers, subbuffers, and images must return no properties.
|
||||
if (test_case.type == buffer || test_case.type == subbuffer
|
||||
|| test_case.type == image)
|
||||
{
|
||||
if (set_size == 0)
|
||||
{
|
||||
return TEST_PASS;
|
||||
}
|
||||
else
|
||||
{
|
||||
log_error("Queried properties must have size equal to zero for "
|
||||
"buffers, subbuffers, and images.");
|
||||
return TEST_FAIL;
|
||||
}
|
||||
}
|
||||
|
||||
if (set_size == 0 && test_case.properties.size() == 0)
|
||||
{
|
||||
return TEST_PASS;
|
||||
@@ -123,8 +188,6 @@ static int run_test_query_properties(cl_context context, cl_command_queue queue,
|
||||
int error = CL_SUCCESS;
|
||||
log_info("\nTC description: %s\n", test_case.description.c_str());
|
||||
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
clMemWrapper obj_src;
|
||||
clMemWrapper obj_dst;
|
||||
clEventWrapper event;
|
||||
@@ -144,12 +207,6 @@ static int run_test_query_properties(cl_context context, cl_command_queue queue,
|
||||
generate_random_data(kUInt, size, init_generator, dst_data.data());
|
||||
free_mtdata(init_generator);
|
||||
init_generator = NULL;
|
||||
const char* kernel_src = test_case.src.c_str();
|
||||
error =
|
||||
create_single_kernel_helper(context, &program, &kernel, 1, &kernel_src,
|
||||
test_case.kernel_name.c_str());
|
||||
|
||||
test_error(error, "create_single_kernel_helper failed");
|
||||
|
||||
flags = (cl_mem_flags)(CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR);
|
||||
error = create_object_and_check_properties(context, obj_src, test_case,
|
||||
@@ -161,37 +218,44 @@ static int run_test_query_properties(cl_context context, cl_command_queue queue,
|
||||
flags, dst_data, size_x, size_y);
|
||||
test_error(error, "create_object_and_check_properties obj_dst failed.");
|
||||
|
||||
error = clSetKernelArg(kernel, 0, sizeof(obj_src), &obj_src);
|
||||
error = clSetKernelArg(test_case.kernel, 0, sizeof(obj_src), &obj_src);
|
||||
test_error(error, "clSetKernelArg 0 failed.");
|
||||
|
||||
error = clSetKernelArg(kernel, 1, sizeof(obj_dst), &obj_dst);
|
||||
error = clSetKernelArg(test_case.kernel, 1, sizeof(obj_dst), &obj_dst);
|
||||
test_error(error, "clSetKernelArg 1 failed.");
|
||||
if (test_case.obj_t == image)
|
||||
switch (test_case.type)
|
||||
{
|
||||
error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global_dim, NULL,
|
||||
0, NULL, &event);
|
||||
test_error(error, "clEnqueueNDRangeKernel failed.");
|
||||
case image:
|
||||
case image_with_properties: {
|
||||
error = clEnqueueNDRangeKernel(queue, test_case.kernel, 2, NULL,
|
||||
global_dim, NULL, 0, NULL, &event);
|
||||
test_error(error, "clEnqueueNDRangeKernel failed.");
|
||||
|
||||
error = clWaitForEvents(1, &event);
|
||||
test_error(error, "clWaitForEvents failed.");
|
||||
error = clWaitForEvents(1, &event);
|
||||
test_error(error, "clWaitForEvents failed.");
|
||||
|
||||
error = clEnqueueReadImage(queue, obj_dst, CL_TRUE, origin, region, 0,
|
||||
0, dst_data.data(), 0, NULL, NULL);
|
||||
test_error(error, "clEnqueueReadImage failed.");
|
||||
}
|
||||
if (test_case.obj_t == buffer)
|
||||
{
|
||||
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL, 0,
|
||||
NULL, &event);
|
||||
test_error(error, "clEnqueueNDRangeKernel failed.");
|
||||
error = clEnqueueReadImage(queue, obj_dst, CL_TRUE, origin, region,
|
||||
0, 0, dst_data.data(), 0, NULL, NULL);
|
||||
test_error(error, "clEnqueueReadImage failed.");
|
||||
}
|
||||
break;
|
||||
case buffer:
|
||||
case buffer_with_properties:
|
||||
case subbuffer: {
|
||||
error = clEnqueueNDRangeKernel(queue, test_case.kernel, 1, NULL,
|
||||
&size, NULL, 0, NULL, &event);
|
||||
test_error(error, "clEnqueueNDRangeKernel failed.");
|
||||
|
||||
error = clWaitForEvents(1, &event);
|
||||
test_error(error, "clWaitForEvents failed.");
|
||||
error = clWaitForEvents(1, &event);
|
||||
test_error(error, "clWaitForEvents failed.");
|
||||
|
||||
error = clEnqueueReadBuffer(queue, obj_dst, CL_TRUE, 0,
|
||||
dst_data.size() * sizeof(cl_uint),
|
||||
dst_data.data(), 0, NULL, NULL);
|
||||
test_error(error, "clEnqueueReadBuffer failed.");
|
||||
error = clEnqueueReadBuffer(queue, obj_dst, CL_TRUE, 0,
|
||||
dst_data.size() * sizeof(cl_uint),
|
||||
dst_data.data(), 0, NULL, NULL);
|
||||
test_error(error, "clEnqueueReadBuffer failed.");
|
||||
}
|
||||
break;
|
||||
default: log_error("Unknown test type!"); return TEST_FAIL;
|
||||
}
|
||||
|
||||
for (size_t i = 0; i < size; ++i)
|
||||
@@ -223,21 +287,31 @@ int test_image_properties_queries(cl_device_id deviceID, cl_context context,
|
||||
return TEST_SKIPPED_ITSELF;
|
||||
}
|
||||
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
|
||||
const char* kernel_src = R"CLC(
|
||||
__kernel void data_copy(read_only image2d_t src, write_only image2d_t dst)
|
||||
{
|
||||
int tid_x = get_global_id(0);
|
||||
int tid_y = get_global_id(1);
|
||||
int2 coords = (int2)(tid_x, tid_y);
|
||||
uint4 val = read_imageui(src, coords);
|
||||
write_imageui(dst, coords, val);
|
||||
|
||||
}
|
||||
)CLC";
|
||||
|
||||
error = create_single_kernel_helper(context, &program, &kernel, 1,
|
||||
&kernel_src, "data_copy");
|
||||
test_error(error, "create_single_kernel_helper failed");
|
||||
|
||||
std::vector<test_data> test_cases;
|
||||
std::string test_kernel = { "__kernel void data_copy(read_only image2d_t "
|
||||
"src, write_only image2d_t dst)\n"
|
||||
"{\n"
|
||||
" int tid_x = get_global_id(0);\n"
|
||||
" int tid_y = get_global_id(1);\n"
|
||||
" int2 coords = (int2)(tid_x, tid_y);\n"
|
||||
" uint4 val = read_imageui(src, coords);\n"
|
||||
" write_imageui(dst, coords, val);\n"
|
||||
"\n"
|
||||
"}\n" };
|
||||
test_cases.push_back({ image, {}, "regular image", kernel });
|
||||
test_cases.push_back(
|
||||
{ image, { 0 }, "image, 0 properties", test_kernel, "data_copy" });
|
||||
{ image_with_properties, { 0 }, "image, 0 properties", kernel });
|
||||
test_cases.push_back(
|
||||
{ image, {}, "image, NULL properties", test_kernel, "data_copy" });
|
||||
{ image_with_properties, {}, "image, NULL properties", kernel });
|
||||
|
||||
for (auto test_case : test_cases)
|
||||
{
|
||||
@@ -251,20 +325,33 @@ int test_buffer_properties_queries(cl_device_id deviceID, cl_context context,
|
||||
cl_command_queue queue, int num_elements)
|
||||
{
|
||||
int error = CL_SUCCESS;
|
||||
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel;
|
||||
|
||||
const char* kernel_src = R"CLC(
|
||||
__kernel void data_copy(__global int *src, __global int *dst)
|
||||
{
|
||||
int tid = get_global_id(0);
|
||||
|
||||
dst[tid] = src[tid];
|
||||
|
||||
}
|
||||
)CLC";
|
||||
error = create_single_kernel_helper(context, &program, &kernel, 1,
|
||||
&kernel_src, "data_copy");
|
||||
test_error(error, "create_single_kernel_helper failed");
|
||||
|
||||
std::vector<test_data> test_cases;
|
||||
std::string test_kernel = {
|
||||
"__kernel void data_copy(__global int *src, __global int *dst)\n"
|
||||
"{\n"
|
||||
" int tid = get_global_id(0);\n"
|
||||
"\n"
|
||||
" dst[tid] = src[tid];\n"
|
||||
"\n"
|
||||
"}\n"
|
||||
};
|
||||
test_cases.push_back({ buffer, {}, "regular buffer", kernel });
|
||||
test_cases.push_back(
|
||||
{ buffer, { 0 }, "buffer, 0 properties", test_kernel, "data_copy" });
|
||||
{ buffer_with_properties, { 0 }, "buffer with 0 properties", kernel });
|
||||
test_cases.push_back(
|
||||
{ buffer, {}, "buffer, NULL properties", test_kernel, "data_copy" });
|
||||
{ buffer_with_properties, {}, "buffer with NULL properties", kernel });
|
||||
test_cases.push_back(
|
||||
{ subbuffer, { 0 }, "subbuffer with 0 properties", kernel });
|
||||
test_cases.push_back(
|
||||
{ subbuffer, {}, "subbuffer with NULL properties", kernel });
|
||||
|
||||
for (auto test_case : test_cases)
|
||||
{
|
||||
|
||||
Reference in New Issue
Block a user