// // 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 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; }