diff --git a/test_common/harness/imageHelpers.cpp b/test_common/harness/imageHelpers.cpp index 64bd92af..f4118e30 100644 --- a/test_common/harness/imageHelpers.cpp +++ b/test_common/harness/imageHelpers.cpp @@ -21,7 +21,6 @@ #if !defined (_WIN32) && !defined(__APPLE__) #include #endif -#include #include #include #if !defined (_WIN32) @@ -3664,16 +3663,15 @@ bool find_format( cl_image_format *formatList, unsigned int numFormats, cl_image return false; } -bool check_minimum_supported(cl_image_format *formatList, - unsigned int numFormats, - cl_mem_flags flags, - cl_mem_object_type image_type, - cl_device_id device) +void build_required_image_formats(cl_mem_flags flags, + cl_mem_object_type image_type, + cl_device_id device, + std::vector& formatsToSupport) { - bool passed = true; - std::vector formatsToSupport; Version version = get_device_cl_version(device); + formatsToSupport.clear(); + // Required embedded formats. static std::vector embeddedProfReadOrWriteFormats { @@ -3810,18 +3808,26 @@ bool check_minimum_supported(cl_image_format *formatList, } } } +} - for (auto &format: formatsToSupport) - { - if( !find_format( formatList, numFormats, &format ) ) - { - log_error( "ERROR: Format required by OpenCL %s is not supported: ", version.to_string().c_str() ); - print_header( &format, true ); - passed = false; - } - } +bool is_image_format_required(cl_image_format format, + cl_mem_flags flags, + cl_mem_object_type image_type, + cl_device_id device) +{ + std::vector formatsToSupport; + build_required_image_formats(flags, image_type, device, formatsToSupport); - return passed; + for (auto &formatItr: formatsToSupport) + { + if (formatItr.image_channel_order == format.image_channel_order && + formatItr.image_channel_data_type == format.image_channel_data_type) + { + return true; + } + } + + return false; } cl_uint compute_max_mip_levels( size_t width, size_t height, size_t depth) diff --git a/test_common/harness/imageHelpers.h b/test_common/harness/imageHelpers.h index d922a689..1ac1faca 100644 --- a/test_common/harness/imageHelpers.h +++ b/test_common/harness/imageHelpers.h @@ -23,6 +23,7 @@ #include #include #include +#include #if !defined(_WIN32) #include @@ -71,11 +72,14 @@ extern void print_read_header( cl_image_format *format, image_sampler_data *samp extern void print_write_header( cl_image_format *format, bool err); extern void print_header( cl_image_format *format, bool err ); extern bool find_format( cl_image_format *formatList, unsigned int numFormats, cl_image_format *formatToFind ); -extern bool check_minimum_supported(cl_image_format *formatList, - unsigned int numFormats, - cl_mem_flags flags, - cl_mem_object_type image_type, - cl_device_id device); +extern bool is_image_format_required(cl_image_format format, + cl_mem_flags flags, + cl_mem_object_type image_type, + cl_device_id device); +extern void build_required_image_formats(cl_mem_flags flags, + cl_mem_object_type image_type, + cl_device_id device, + std::vector& formatsToSupport); extern size_t get_format_type_size( const cl_image_format *format ); extern size_t get_channel_data_type_size( cl_channel_type channelType ); diff --git a/test_conformance/basic/test_readimage3d.cpp b/test_conformance/basic/test_readimage3d.cpp index 95a5b5a1..1337c9fb 100644 --- a/test_conformance/basic/test_readimage3d.cpp +++ b/test_conformance/basic/test_readimage3d.cpp @@ -1,6 +1,6 @@ // // 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 @@ -14,6 +14,7 @@ // limitations under the License. // #include "harness/compat.h" +#include "harness/imageHelpers.h" #include #include @@ -22,7 +23,6 @@ #include #include - #include "procs.h" static const char *bgra8888_kernel_code = @@ -43,7 +43,6 @@ static const char *bgra8888_kernel_code = "\n" "}\n"; - static const char *rgba8888_kernel_code = "\n" "__kernel void test_rgba8888(read_only image3d_t srcimg, __global float4 *dst, sampler_t sampler)\n" @@ -63,7 +62,6 @@ static const char *rgba8888_kernel_code = "\n" "}\n"; - static unsigned char * generate_3d_image8(int w, int h, int d, MTdata data) { @@ -110,123 +108,106 @@ prepare_reference(unsigned char * input_ptr, int w, int h, int d) return ptr; } - int test_readimage3d(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) { - cl_mem streams[3]; - cl_program program[2]; - cl_kernel kernel[2]; - cl_image_format img_format; - unsigned char *input_ptr[2]; - float *output_ptr; - double *ref_ptr[2]; - size_t threads[3]; - int img_width = 64; - int img_height = 64; - int img_depth = 64; - int i, err; - size_t origin[3] = {0, 0, 0}; - size_t region[3] = {img_width, img_height, img_depth}; - size_t length = img_width * img_height * img_depth * 4 * sizeof(float); + cl_mem streams[2]; + cl_program program; + cl_kernel kernel; + cl_sampler sampler; + struct testFormat + { + const char* kernelName; + const char* kernelSourceString; + const cl_image_format img_format; + }; + static testFormat formatsToTest[] = + { + { + "test_bgra8888", + bgra8888_kernel_code, + {CL_BGRA, CL_UNORM_INT8}, + }, + { + "test_rgba8888", + rgba8888_kernel_code, + {CL_RGBA, CL_UNORM_INT8}, + }, + }; - PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device ) + unsigned char *input_ptr; + float *output_ptr; + double *ref_ptr; + size_t threads[3]; + int img_width = 64; + int img_height = 64; + int img_depth = 64; + int i, err; + size_t origin[3] = {0, 0, 0}; + size_t region[3] = {img_width, img_height, img_depth}; + size_t length = img_width * img_height * img_depth * 4 * sizeof(float); - MTdata d = init_genrand( gRandomSeed ); - input_ptr[0] = generate_3d_image8(img_width, img_height, img_depth, d); - input_ptr[1] = generate_3d_image8(img_width, img_height, img_depth, d); - ref_ptr[0] = prepare_reference(input_ptr[0], img_width, img_height, img_depth); - ref_ptr[1] = prepare_reference(input_ptr[1], img_width, img_height, img_depth); - free_mtdata(d); d = NULL; - output_ptr = (float*)malloc(length); + PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device ) - img_format.image_channel_order = CL_BGRA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[0] = create_image_3d(context, CL_MEM_READ_ONLY, &img_format, img_width, img_height, img_depth, 0, 0, NULL, &err); - test_error(err, "create_image_3d failed"); + for (uint32_t i = 0; i < ARRAY_SIZE(formatsToTest); i++) + { + if (!is_image_format_required(formatsToTest[i].img_format, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, device)) + continue; - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[1] = create_image_3d(context, CL_MEM_READ_ONLY, &img_format, img_width, img_height, img_depth, 0, 0, NULL, &err); - test_error(err, "create_image_3d failed"); + MTdata d = init_genrand( gRandomSeed ); + input_ptr = generate_3d_image8(img_width, img_height, img_depth, d); + ref_ptr = prepare_reference(input_ptr, img_width, img_height, img_depth); + output_ptr = (float*)malloc(length); - streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err); - test_error(err, "clCreateBuffer failed"); + streams[0] = create_image_3d(context, CL_MEM_READ_ONLY, &formatsToTest[i].img_format, img_width, img_height, img_depth, 0, 0, NULL, &err); + test_error(err, "create_image_3d failed"); - err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr[0], 0, NULL, NULL); - test_error(err, "clEnqueueWriteImage failed"); + streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err); + test_error(err, "clCreateBuffer failed"); - err = clEnqueueWriteImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, input_ptr[1], 0, NULL, NULL); - test_error(err, "clEnqueueWriteImage failed"); + sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); - err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &bgra8888_kernel_code, "test_bgra8888" ); - if (err) - return -1; + err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr, 0, NULL, NULL); + test_error(err, "clEnqueueWriteImage failed"); - err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &rgba8888_kernel_code, "test_rgba8888" ); - if (err) - return -1; + err = create_single_kernel_helper(context, &program, &kernel, 1, &formatsToTest[i].kernelSourceString, formatsToTest[i].kernelName); + test_error(err, "create_single_kernel_helper failed"); - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); + err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); + err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); + err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); + test_error(err, "clSetKernelArg failed"); - err = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]); - err |= clSetKernelArg(kernel[0], 1, sizeof streams[2], &streams[2]); - err |= clSetKernelArg(kernel[0], 2, sizeof sampler, &sampler); - test_error(err, "clSetKernelArg failed"); + threads[0] = (unsigned int)img_width; + threads[1] = (unsigned int)img_height; + threads[2] = (unsigned int)img_depth; - err = clSetKernelArg(kernel[1], 0, sizeof streams[1], &streams[1]); - err |= clSetKernelArg(kernel[1], 1, sizeof streams[2], &streams[2]); - err |= clSetKernelArg(kernel[1], 2, sizeof sampler, &sampler); - test_error(err, "clSetKernelArg failed"); + err = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, threads, NULL, 0, NULL, NULL); + test_error(err, "clEnqueueNDRangeKernel failed"); - threads[0] = (unsigned int)img_width; - threads[1] = (unsigned int)img_height; - threads[2] = (unsigned int)img_depth; + err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL); + test_error(err, "clEnqueueReadBuffer failed"); - for (i=0; i<2; i++) - { - err = clEnqueueNDRangeKernel(queue, kernel[i], 3, NULL, threads, NULL, 0, NULL, NULL); - test_error(err, "clEnqueueNDRangeKernel failed"); + err = verify_3d_image8(ref_ptr, output_ptr, img_width, img_height, img_depth); + if ( err == 0 ) + { + log_info("READ_IMAGE3D_%s_%s test passed\n", + GetChannelTypeName(formatsToTest[i].img_format.image_channel_data_type), + GetChannelOrderName(formatsToTest[i].img_format.image_channel_order)); + } - err = clEnqueueReadBuffer(queue, streams[2], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL); - test_error(err, "clEnqueueReadBuffer failed"); + clReleaseSampler(sampler); + clReleaseMemObject(streams[0]); + clReleaseMemObject(streams[1]); + clReleaseKernel(kernel); + clReleaseProgram(program); + free_mtdata(d); + d = NULL; + free(input_ptr); + free(ref_ptr); + free(output_ptr); + } - switch (i) - { - case 0: - err = verify_3d_image8(ref_ptr[i], output_ptr, img_width, img_height, img_depth); - if ( err != 0 ) - log_info("READ_IMAGE3D_BGRA_UNORM_INT8 test passed\n"); - break; - case 1: - err = verify_3d_image8(ref_ptr[i], output_ptr, img_width, img_height, img_depth); - if ( err != 0 ) - log_info("READ_IMAGE3D_RGBA_UNORM_INT8 test passed\n"); - break; - } - - if (err) - break; - } - - // cleanup - clReleaseSampler(sampler); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - clReleaseMemObject(streams[2]); - for (i=0; i<2; i++) - { - clReleaseKernel(kernel[i]); - clReleaseProgram(program[i]); - } - free(input_ptr[0]); - free(input_ptr[1]); - free(output_ptr); - free(ref_ptr[0]); - free(ref_ptr[1]); - - return err; + return err; } - - diff --git a/test_conformance/images/clGetInfo/test_loops.cpp b/test_conformance/images/clGetInfo/test_loops.cpp index 28f732c7..a9d20723 100644 --- a/test_conformance/images/clGetInfo/test_loops.cpp +++ b/test_conformance/images/clGetInfo/test_loops.cpp @@ -14,6 +14,9 @@ // limitations under the License. // #include "../testBase.h" +#include "harness/imageHelpers.h" +#include +#include extern cl_filter_mode gFilterModeToUse; extern cl_addressing_mode gAddressModeToUse; @@ -36,6 +39,30 @@ static const char *str_3d_image = "3D"; static const char *str_1d_image_array = "1D array"; static const char *str_2d_image_array = "2D array"; +static bool check_minimum_supported(cl_image_format *formatList, + unsigned int numFormats, + cl_mem_flags flags, + cl_mem_object_type image_type, + cl_device_id device) +{ + bool passed = true; + Version version = get_device_cl_version(device); + std::vector formatsToSupport; + build_required_image_formats(flags, image_type, device, formatsToSupport); + + for (auto &format: formatsToSupport) + { + if( !find_format( formatList, numFormats, &format ) ) + { + log_error( "ERROR: Format required by OpenCL %s is not supported: ", version.to_string().c_str() ); + print_header( &format, true ); + passed = false; + } + } + + return passed; +} + static const char *convert_image_type_to_string(cl_mem_object_type image_type) { const char *p;