// // Copyright (c) 2017 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // // http://www.apache.org/licenses/LICENSE-2.0 // // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // See the License for the specific language governing permissions and // limitations under the License. // #ifndef TEST_CONFORMANCE_CLCPP_IMAGES_TEST_SAMPLE_HPP #define TEST_CONFORMANCE_CLCPP_IMAGES_TEST_SAMPLE_HPP #include #include #include #include #include "common.hpp" namespace test_images_sample { enum class sampler_source { param, program_scope }; const sampler_source sampler_sources[] = { sampler_source::param, sampler_source::program_scope }; template struct image_test : image_test_base { cl_channel_order channel_order; sampler_source source; image_test(cl_channel_order channel_order, sampler_source source) : channel_order(channel_order), source(source) { } // ----------------------------------------------------------------------------------- // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------ // ----------------------------------------------------------------------------------- #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS) std::string generate_source() { std::stringstream s; s << R"( typedef )" << type_name() << R"( element_type; )"; std::string sampler; if (source == sampler_source::program_scope) { s << R"( constant sampler_t sampler_program_scope = CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE; )"; sampler = "sampler_program_scope"; } else if (source == sampler_source::param) { sampler = "sampler_param"; } s << R"( kernel void test( read_only )" << image_test::image_type_name() << R"(_t img, const global int4 *coords, global element_type *output, sampler_t sampler_param ) { const ulong gid = get_global_linear_id(); output[gid] = read_image)" << image_test::function_suffix() << "(img, " << sampler << ", coords[gid]." << image_test::coord_accessor() << R"(); } )"; return s.str(); } #else std::string generate_source() { std::stringstream s; s << R"( #include #include #include #include using namespace cl; )"; s << R"( typedef )" << type_name() << R"( element_type; )"; std::string sampler; if (source == sampler_source::program_scope) { s << R"( sampler sampler_program_scope = make_sampler(); )"; sampler = "sampler_program_scope"; } else if (source == sampler_source::param) { sampler = "sampler_param"; } s << R"( kernel void test( const )" << image_test::image_type_name() << R"( img, const global_ptr coords, global_ptr output, sampler sampler_param ) { const ulong gid = get_global_linear_id(); output[gid] = img.sample()" << sampler << ", coords[gid]." << image_test::coord_accessor() << R"(); } )"; return s.str(); } #endif int run(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { int error = CL_SUCCESS; cl_program program; cl_kernel kernel; std::string kernel_name = "test"; std::string source = generate_source(); // ----------------------------------------------------------------------------------- // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------ // ----------------------------------------------------------------------------------- // Only OpenCL C++ to SPIR-V compilation #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION) error = create_opencl_kernel( context, &program, &kernel, source, kernel_name ); RETURN_ON_ERROR(error) return error; // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code) #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS) error = create_opencl_kernel( context, &program, &kernel, source, kernel_name, "-cl-std=CL2.0", false ); RETURN_ON_ERROR(error) // Normal run #else error = create_opencl_kernel( context, &program, &kernel, source, kernel_name ); RETURN_ON_ERROR(error) #endif using element_type = typename image_test::element_type; using coord_type = cl_int4; using scalar_element_type = typename scalar_type::type; using channel_type = typename image_test::channel_type; cl_image_format image_format; image_format.image_channel_order = channel_order; image_format.image_channel_data_type = ChannelType; const size_t pixel_size = get_pixel_size(&image_format); const size_t channel_count = get_channel_order_channel_count(image_format.image_channel_order); cl_image_desc image_desc; image_desc.image_type = ImageType; if (ImageType == CL_MEM_OBJECT_IMAGE1D) { image_desc.image_width = 2048; image_desc.image_height = 1; image_desc.image_depth = 1; } else if (ImageType == CL_MEM_OBJECT_IMAGE2D) { image_desc.image_width = 256; image_desc.image_height = 256; image_desc.image_depth = 1; } else if (ImageType == CL_MEM_OBJECT_IMAGE3D) { image_desc.image_width = 64; image_desc.image_height = 64; image_desc.image_depth = 64; } image_desc.image_array_size = 0; image_desc.image_row_pitch = image_desc.image_width * pixel_size; image_desc.image_slice_pitch = image_desc.image_row_pitch * image_desc.image_height; image_desc.num_mip_levels = 0; image_desc.num_samples = 0; image_desc.mem_object = NULL; image_descriptor image_info = create_image_descriptor(image_desc, &image_format); std::vector image_values = generate_input( image_desc.image_width * image_desc.image_height * image_desc.image_depth * channel_count, image_test::channel_min(), image_test::channel_max(), std::vector() ); const size_t count = num_elements; std::vector coords = generate_input( count, detail::make_value(0), coord_type { static_cast(image_desc.image_width - 1), static_cast(image_desc.image_height - 1), static_cast(image_desc.image_depth - 1), 0 }, std::vector() ); cl_mem img = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &image_format, &image_desc, static_cast(image_values.data()), &error); RETURN_ON_CL_ERROR(error, "clCreateImage") cl_mem coords_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(coord_type) * count, static_cast(coords.data()), &error); RETURN_ON_CL_ERROR(error, "clCreateBuffer") cl_mem output_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(element_type) * count, NULL, &error); RETURN_ON_CL_ERROR(error, "clCreateBuffer") const cl_sampler_properties sampler_properties[] = { CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_NONE, CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, 0 }; cl_sampler sampler = clCreateSamplerWithProperties(context, sampler_properties, &error); RETURN_ON_CL_ERROR(error, "clCreateSamplerWithProperties") error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &img); RETURN_ON_CL_ERROR(error, "clSetKernelArg") error = clSetKernelArg(kernel, 1, sizeof(coords_buffer), &coords_buffer); RETURN_ON_CL_ERROR(error, "clSetKernelArg") error = clSetKernelArg(kernel, 2, sizeof(output_buffer), &output_buffer); RETURN_ON_CL_ERROR(error, "clSetKernelArg") error = clSetKernelArg(kernel, 3, sizeof(sampler), &sampler); RETURN_ON_CL_ERROR(error, "clSetKernelArg") const size_t global_size = count; error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL); RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel") std::vector output(count); error = clEnqueueReadBuffer( queue, output_buffer, CL_TRUE, 0, sizeof(element_type) * count, static_cast(output.data()), 0, NULL, NULL ); RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer") for (size_t i = 0; i < count; i++) { const coord_type c = coords[i]; const element_type result = output[i]; element_type expected; read_image_pixel(static_cast(image_values.data()), &image_info, c.s[0], c.s[1], c.s[2], expected.s); if (!are_equal(result, expected)) { RETURN_ON_ERROR_MSG(-1, "Sampling from coordinates %s failed. Expected: %s, got: %s", format_value(c).c_str(), format_value(expected).c_str(), format_value(result).c_str() ); } } clReleaseMemObject(img); clReleaseMemObject(coords_buffer); clReleaseMemObject(output_buffer); clReleaseSampler(sampler); clReleaseKernel(kernel); clReleaseProgram(program); return error; } }; template int run_test_cases(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { if (!is_test_supported(device)) return CL_SUCCESS; int error = CL_SUCCESS; for (auto channel_order : get_channel_orders(device)) for (auto source : sampler_sources) { error = image_test(channel_order, source) .run(device, context, queue, num_elements); RETURN_ON_ERROR(error) error = image_test(channel_order, source) .run(device, context, queue, num_elements); RETURN_ON_ERROR(error) error = image_test(channel_order, source) .run(device, context, queue, num_elements); RETURN_ON_ERROR(error) error = image_test(channel_order, source) .run(device, context, queue, num_elements); RETURN_ON_ERROR(error) error = image_test(channel_order, source) .run(device, context, queue, num_elements); RETURN_ON_ERROR(error) error = image_test(channel_order, source) .run(device, context, queue, num_elements); RETURN_ON_ERROR(error) error = image_test(channel_order, source) .run(device, context, queue, num_elements); RETURN_ON_ERROR(error) } return error; } AUTO_TEST_CASE(test_images_sample_1d) (cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { return run_test_cases(device, context, queue, num_elements); } AUTO_TEST_CASE(test_images_sample_2d) (cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { return run_test_cases(device, context, queue, num_elements); } AUTO_TEST_CASE(test_images_sample_3d) (cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { return run_test_cases(device, context, queue, num_elements); } } // namespace #endif // TEST_CONFORMANCE_CLCPP_IMAGES_TEST_SAMPLE_HPP