| // |
| // 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. |
| // |
| #include "common.h" |
| |
| #include <algorithm> |
| |
| using namespace std; |
| |
| struct image_kernel_data |
| { |
| cl_int width; |
| cl_int height; |
| cl_int depth; |
| cl_int arraySize; |
| cl_int widthDim; |
| cl_int heightDim; |
| cl_int channelType; |
| cl_int channelOrder; |
| cl_int expectedChannelType; |
| cl_int expectedChannelOrder; |
| cl_int numSamples; |
| }; |
| |
| static const char *methodTestKernelPattern = |
| "%s" |
| "typedef struct {\n" |
| " int width;\n" |
| " int height;\n" |
| " int depth;\n" |
| " int arraySize;\n" |
| " int widthDim;\n" |
| " int heightDim;\n" |
| " int channelType;\n" |
| " int channelOrder;\n" |
| " int expectedChannelType;\n" |
| " int expectedChannelOrder;\n" |
| " int numSamples;\n" |
| " } image_kernel_data;\n" |
| "__kernel void sample_kernel( read_only %s input, __global image_kernel_data *outData )\n" |
| "{\n" |
| "%s%s%s%s%s%s%s%s%s%s%s" |
| "}\n"; |
| |
| static const char *arraySizeKernelLine = |
| " outData->arraySize = get_image_array_size( input );\n"; |
| static const char *imageWidthKernelLine = |
| " outData->width = get_image_width( input );\n"; |
| static const char *imageHeightKernelLine = |
| " outData->height = get_image_height( input );\n"; |
| static const char *imageDimKernelLine = |
| " int2 dim = get_image_dim( input );\n"; |
| static const char *imageWidthDimKernelLine = |
| " outData->widthDim = dim.x;\n"; |
| static const char *imageHeightDimKernelLine = |
| " outData->heightDim = dim.y;\n"; |
| static const char *channelTypeKernelLine = |
| " outData->channelType = get_image_channel_data_type( input );\n"; |
| static const char *channelTypeConstLine = |
| " outData->expectedChannelType = CLK_%s;\n"; |
| static const char *channelOrderKernelLine = |
| " outData->channelOrder = get_image_channel_order( input );\n"; |
| static const char *channelOrderConstLine = |
| " outData->expectedChannelOrder = CLK_%s;\n"; |
| static const char *numSamplesKernelLine = |
| " outData->numSamples = get_image_num_samples( input );\n"; |
| static const char *enableMSAAKernelLine = |
| "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"; |
| |
| static int verify(cl_int input, cl_int kernelOutput, const char * description) |
| { |
| if( kernelOutput != input ) |
| { |
| log_error( "ERROR: %s did not validate (expected %d, got %d)\n", description, input, kernelOutput); |
| return -1; |
| } |
| return 0; |
| } |
| |
| extern int supportsMsaa(cl_context context, bool* supports_msaa); |
| extern int supportsDepth(cl_context context, bool* supports_depth); |
| |
| int test_image_format_methods( cl_device_id device, cl_context context, cl_command_queue queue, |
| size_t width, size_t height, size_t arraySize, size_t samples, |
| GLenum target, format format, MTdata d ) |
| { |
| int error, result=0; |
| |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| clMemWrapper image, outDataBuffer; |
| char programSrc[ 10240 ]; |
| |
| image_kernel_data outKernelData; |
| |
| #ifdef GL_VERSION_3_2 |
| if (get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE || |
| get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE_ARRAY) |
| { |
| bool supports_msaa; |
| error = supportsMsaa(context, &supports_msaa); |
| if( error != 0 ) return error; |
| if (!supports_msaa) return 0; |
| } |
| if (format.formattype == GL_DEPTH_COMPONENT || |
| format.formattype == GL_DEPTH_STENCIL) |
| { |
| bool supports_depth; |
| error = supportsDepth(context, &supports_depth); |
| if( error != 0 ) return error; |
| if (!supports_depth) return 0; |
| } |
| #endif |
| DetectFloatToHalfRoundingMode(queue); |
| |
| glTextureWrapper glTexture; |
| switch (get_base_gl_target(target)) { |
| case GL_TEXTURE_2D: |
| CreateGLTexture2D( width, height, target, |
| format.formattype, format.internal, format.datatype, |
| format.type, &glTexture, &error, false, d ); |
| break; |
| case GL_TEXTURE_2D_ARRAY: |
| CreateGLTexture2DArray( width, height, arraySize, target, |
| format.formattype, format.internal, format.datatype, |
| format.type, &glTexture, &error, false, d ); |
| break; |
| case GL_TEXTURE_2D_MULTISAMPLE: |
| CreateGLTexture2DMultisample( width, height, samples, target, |
| format.formattype, format.internal, format.datatype, |
| format.type, &glTexture, &error, false, d, false); |
| break; |
| case GL_TEXTURE_2D_MULTISAMPLE_ARRAY: |
| CreateGLTexture2DArrayMultisample( width, height, arraySize, samples, target, |
| format.formattype, format.internal, format.datatype, |
| format.type, &glTexture, &error, false, d, false); |
| break; |
| |
| default: |
| log_error("Unsupported GL tex target (%s) passed to write test: " |
| "%s (%s):%d", GetGLTargetName(target), __FUNCTION__, |
| __FILE__, __LINE__); |
| } |
| |
| // Check to see if the texture could not be created for some other reason like |
| // GL_FRAMEBUFFER_UNSUPPORTED |
| if (error == GL_FRAMEBUFFER_UNSUPPORTED) { |
| return 0; |
| } |
| |
| // Construct testing source |
| log_info( " - Creating image %d by %d...\n", width, height ); |
| // Create a CL image from the supplied GL texture |
| image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_READ_ONLY, |
| target, 0, glTexture, &error ); |
| |
| if ( error != CL_SUCCESS ) { |
| print_error( error, "Unable to create CL image from GL texture" ); |
| GLint fmt; |
| glGetTexLevelParameteriv( target, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt ); |
| log_error( " Supplied GL texture was base format %s and internal " |
| "format %s\n", GetGLBaseFormatName( fmt ), GetGLFormatName( fmt ) ); |
| return error; |
| } |
| |
| cl_image_format imageFormat; |
| error = clGetImageInfo (image, CL_IMAGE_FORMAT, |
| sizeof(imageFormat), &imageFormat, NULL); |
| test_error(error, "Failed to get image format"); |
| |
| const char * imageType = 0; |
| bool doArraySize = false; |
| bool doImageWidth = false; |
| bool doImageHeight = false; |
| bool doImageChannelDataType = false; |
| bool doImageChannelOrder = false; |
| bool doImageDim = false; |
| bool doNumSamples = false; |
| bool doMSAA = false; |
| switch(target) { |
| case GL_TEXTURE_2D: |
| imageType = "image2d_depth_t"; |
| doImageWidth = true; |
| doImageHeight = true; |
| doImageChannelDataType = true; |
| doImageChannelOrder = true; |
| doImageDim = true; |
| break; |
| case GL_TEXTURE_2D_ARRAY: |
| imageType = "image2d_array_depth_t"; |
| doImageWidth = true; |
| doImageHeight = true; |
| doArraySize = true; |
| doImageChannelDataType = true; |
| doImageChannelOrder = true; |
| doImageDim = true; |
| doArraySize = true; |
| break; |
| case GL_TEXTURE_2D_MULTISAMPLE: |
| doNumSamples = true; |
| doMSAA = true; |
| if(format.formattype == GL_DEPTH_COMPONENT) { |
| doImageWidth = true; |
| imageType = "image2d_msaa_depth_t"; |
| } else { |
| imageType = "image2d_msaa_t"; |
| } |
| break; |
| case GL_TEXTURE_2D_MULTISAMPLE_ARRAY: |
| doMSAA = true; |
| if(format.formattype == GL_DEPTH_COMPONENT) { |
| doImageWidth = true; |
| imageType = "image2d_msaa_array_depth_t"; |
| } else { |
| imageType = "image2d_array_msaa_t"; |
| } |
| break; |
| } |
| |
| |
| |
| char channelTypeConstKernelLine[512] = {0}; |
| char channelOrderConstKernelLine[512] = {0}; |
| const char* channelTypeName=0; |
| const char* channelOrderName=0; |
| if(doImageChannelDataType) { |
| channelTypeName = GetChannelTypeName( imageFormat.image_channel_data_type ); |
| if(channelTypeName && strlen(channelTypeName)) { |
| // replace CL_* with CLK_* |
| sprintf(channelTypeConstKernelLine, channelTypeConstLine, &channelTypeName[3]); |
| } |
| } |
| if(doImageChannelOrder) { |
| channelOrderName = GetChannelOrderName( imageFormat.image_channel_order ); |
| if(channelOrderName && strlen(channelOrderName)) { |
| // replace CL_* with CLK_* |
| sprintf(channelOrderConstKernelLine, channelOrderConstLine, &channelOrderName[3]); |
| } |
| } |
| |
| // Create a program to run against |
| sprintf(programSrc, |
| methodTestKernelPattern, |
| ( doMSAA ) ? enableMSAAKernelLine : "", |
| imageType, |
| ( doArraySize ) ? arraySizeKernelLine : "", |
| ( doImageWidth ) ? imageWidthKernelLine : "", |
| ( doImageHeight ) ? imageHeightKernelLine : "", |
| ( doImageChannelDataType ) ? channelTypeKernelLine : "", |
| ( doImageChannelDataType ) ? channelTypeConstKernelLine : "", |
| ( doImageChannelOrder ) ? channelOrderKernelLine : "", |
| ( doImageChannelOrder ) ? channelOrderConstKernelLine : "", |
| ( doImageDim ) ? imageDimKernelLine : "", |
| ( doImageDim && doImageWidth ) ? imageWidthDimKernelLine : "", |
| ( doImageDim && doImageHeight ) ? imageHeightDimKernelLine : "", |
| ( doNumSamples ) ? numSamplesKernelLine : ""); |
| |
| |
| //log_info("-----------------------------------\n%s\n", programSrc); |
| error = clFinish(queue); |
| if (error) |
| print_error(error, "clFinish failed.\n"); |
| const char *ptr = programSrc; |
| error = create_single_kernel_helper( context, &program, &kernel, 1, &ptr, "sample_kernel" ); |
| test_error( error, "Unable to create kernel to test against" ); |
| |
| // Create an output buffer |
| outDataBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| sizeof(outKernelData), NULL, &error); |
| test_error( error, "Unable to create output buffer" ); |
| |
| // Set up arguments and run |
| error = clSetKernelArg( kernel, 0, sizeof( image ), &image ); |
| test_error( error, "Unable to set kernel argument" ); |
| error = clSetKernelArg( kernel, 1, sizeof( outDataBuffer ), &outDataBuffer ); |
| test_error( error, "Unable to set kernel argument" ); |
| |
| // Finish and Acquire. |
| glFinish(); |
| error = (*clEnqueueAcquireGLObjects_ptr)(queue, 1, &image, 0, NULL, NULL); |
| test_error(error, "Unable to acquire GL obejcts"); |
| |
| size_t threads[1] = { 1 }, localThreads[1] = { 1 }; |
| |
| error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL ); |
| test_error( error, "Unable to run kernel" ); |
| |
| error = clEnqueueReadBuffer( queue, outDataBuffer, CL_TRUE, 0, sizeof( outKernelData ), &outKernelData, 0, NULL, NULL ); |
| test_error( error, "Unable to read data buffer" ); |
| |
| // Verify the results now |
| if( doImageWidth ) |
| result |= verify(width, outKernelData.width, "width"); |
| if( doImageHeight) |
| result |= verify(height, outKernelData.height, "height"); |
| if( doImageDim && doImageWidth ) |
| result |= verify(width, outKernelData.widthDim, "width from get_image_dim"); |
| if( doImageDim && doImageHeight ) |
| result |= verify(height, outKernelData.heightDim, "height from get_image_dim"); |
| if( doImageChannelDataType ) |
| result |= verify(outKernelData.channelType, outKernelData.expectedChannelType, channelTypeName); |
| if( doImageChannelOrder ) |
| result |= verify(outKernelData.channelOrder, outKernelData.expectedChannelOrder, channelOrderName); |
| if( doArraySize ) |
| result |= verify(arraySize, outKernelData.arraySize, "array size"); |
| if( doNumSamples ) |
| result |= verify(samples, outKernelData.numSamples, "samples"); |
| if(result) { |
| log_error("Test image methods failed"); |
| } |
| |
| clEventWrapper event; |
| error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &image, 0, NULL, &event ); |
| test_error(error, "clEnqueueReleaseGLObjects failed"); |
| |
| error = clWaitForEvents( 1, &event ); |
| test_error(error, "clWaitForEvents failed"); |
| |
| return result; |
| } |
| |
| int test_image_methods_depth( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ){ |
| if (!is_extension_available(device, "cl_khr_gl_depth_images")) { |
| log_info("Test not run because 'cl_khr_gl_depth_images' extension is not supported by the tested device\n"); |
| return 0; |
| } |
| |
| size_t pixelSize; |
| int result = 0; |
| GLenum depth_targets[] = {GL_TEXTURE_2D, GL_TEXTURE_2D_ARRAY}; |
| size_t ntargets = sizeof(depth_targets) / sizeof(depth_targets[0]); |
| size_t nformats = sizeof(depth_formats) / sizeof(depth_formats[0]); |
| |
| const size_t nsizes = 5; |
| sizevec_t sizes[nsizes]; |
| // Need to limit texture size according to GL device properties |
| GLint maxTextureSize = 4096, maxTextureRectangleSize = 4096, maxTextureLayers = 16, size; |
| glGetIntegerv(GL_MAX_TEXTURE_SIZE, &maxTextureSize); |
| glGetIntegerv(GL_MAX_RECTANGLE_TEXTURE_SIZE_EXT, &maxTextureRectangleSize); |
| glGetIntegerv(GL_MAX_ARRAY_TEXTURE_LAYERS, &maxTextureLayers); |
| |
| size = min(maxTextureSize, maxTextureRectangleSize); |
| |
| RandomSeed seed( gRandomSeed ); |
| |
| // Generate some random sizes (within reasonable ranges) |
| for (size_t i = 0; i < nsizes; i++) { |
| sizes[i].width = random_in_range( 2, min(size, 1<<(i+4)), seed ); |
| sizes[i].height = random_in_range( 2, min(size, 1<<(i+4)), seed ); |
| sizes[i].depth = random_in_range( 2, min(maxTextureLayers, 1<<(i+4)), seed ); |
| } |
| |
| for (size_t i = 0; i < nsizes; i++) { |
| for(size_t itarget = 0; itarget < ntargets; ++itarget) { |
| for(size_t iformat = 0; iformat < nformats; ++iformat) |
| result |= test_image_format_methods(device, context, queue, sizes[i].width, sizes[i].height, (depth_targets[itarget] == GL_TEXTURE_2D_ARRAY) ? sizes[i].depth: 1, 0, |
| depth_targets[itarget], depth_formats[iformat], seed ); |
| } |
| } |
| return result; |
| } |
| |
| int test_image_methods_multisample( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ){ |
| if (!is_extension_available(device, "cl_khr_gl_msaa_sharing")) { |
| log_info("Test not run because 'cl_khr_gl_msaa_sharing' extension is not supported by the tested device\n"); |
| return 0; |
| } |
| |
| size_t pixelSize; |
| int result = 0; |
| GLenum targets[] = {GL_TEXTURE_2D_MULTISAMPLE, GL_TEXTURE_2D_MULTISAMPLE_ARRAY}; |
| size_t ntargets = sizeof(targets) / sizeof(targets[0]); |
| size_t nformats = sizeof(common_formats) / sizeof(common_formats[0]); |
| |
| const size_t nsizes = 5; |
| sizevec_t sizes[nsizes]; |
| GLint maxTextureLayers = 16, maxTextureSize = 4096; |
| glGetIntegerv(GL_MAX_ARRAY_TEXTURE_LAYERS, &maxTextureLayers); |
| glGetIntegerv(GL_MAX_TEXTURE_SIZE, &maxTextureSize); |
| |
| RandomSeed seed( gRandomSeed ); |
| |
| // Generate some random sizes (within reasonable ranges) |
| for (size_t i = 0; i < nsizes; i++) { |
| sizes[i].width = random_in_range( 2, min(maxTextureSize, 1<<(i+4)), seed ); |
| sizes[i].height = random_in_range( 2, min(maxTextureSize, 1<<(i+4)), seed ); |
| sizes[i].depth = random_in_range( 2, min(maxTextureLayers, 1<<(i+4)), seed ); |
| } |
| |
| glEnable(GL_MULTISAMPLE); |
| |
| for (size_t i = 0; i < nsizes; i++) { |
| for(size_t itarget = 0; itarget < ntargets; ++itarget) { |
| for(size_t iformat = 0; iformat < nformats; ++iformat) { |
| GLint samples = get_gl_max_samples(targets[itarget], common_formats[iformat].internal); |
| result |= test_image_format_methods(device, context, queue, sizes[i].width, sizes[i].height, (targets[ntargets] == GL_TEXTURE_2D_MULTISAMPLE_ARRAY) ? sizes[i].depth: 1, |
| samples, targets[itarget], common_formats[iformat], seed ); |
| } |
| } |
| } |
| return result; |
| } |