diff --git a/test_conformance/api/test_mem_object_properties_queries.cpp b/test_conformance/api/test_mem_object_properties_queries.cpp index a6ed7ad1..2360d83e 100644 --- a/test_conformance/api/test_mem_object_properties_queries.cpp +++ b/test_conformance/api/test_mem_object_properties_queries.cpp @@ -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 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 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_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_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) {