// // 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; }; // clang-format off 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"; // clang-format on 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; } 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; } 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; }