• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 //    http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "common.h"
17 
18 #include <algorithm>
19 
20 using namespace std;
21 
22 typedef struct image_kernel_data
23 {
24     cl_int width;
25     cl_int height;
26     cl_int depth;
27   cl_int arraySize;
28     cl_int widthDim;
29     cl_int heightDim;
30     cl_int channelType;
31     cl_int channelOrder;
32     cl_int expectedChannelType;
33     cl_int expectedChannelOrder;
34   cl_int numSamples;
35 };
36 
37 static const char *methodTestKernelPattern =
38 "%s"
39 "typedef struct {\n"
40 "    int width;\n"
41 "    int height;\n"
42 "    int depth;\n"
43 "    int arraySize;\n"
44 "    int widthDim;\n"
45 "    int heightDim;\n"
46 "    int channelType;\n"
47 "    int channelOrder;\n"
48 "    int expectedChannelType;\n"
49 "    int expectedChannelOrder;\n"
50 "    int numSamples;\n"
51 " } image_kernel_data;\n"
52 "__kernel void sample_kernel( read_only %s input, __global image_kernel_data *outData )\n"
53 "{\n"
54 "%s%s%s%s%s%s%s%s%s%s%s"
55 "}\n";
56 
57 static const char *arraySizeKernelLine =
58 "   outData->arraySize = get_image_array_size( input );\n";
59 static const char *imageWidthKernelLine =
60 "   outData->width = get_image_width( input );\n";
61 static const char *imageHeightKernelLine =
62 "   outData->height = get_image_height( input );\n";
63 static const char *imageDimKernelLine =
64 "   int2 dim = get_image_dim( input );\n";
65 static const char *imageWidthDimKernelLine =
66 "   outData->widthDim = dim.x;\n";
67 static const char *imageHeightDimKernelLine =
68 "   outData->heightDim = dim.y;\n";
69 static const char *channelTypeKernelLine =
70 "   outData->channelType = get_image_channel_data_type( input );\n";
71 static const char *channelTypeConstLine =
72 "   outData->expectedChannelType = CLK_%s;\n";
73 static const char *channelOrderKernelLine =
74 "   outData->channelOrder = get_image_channel_order( input );\n";
75 static const char *channelOrderConstLine =
76 "   outData->expectedChannelOrder = CLK_%s;\n";
77 static const char *numSamplesKernelLine =
78 "   outData->numSamples = get_image_num_samples( input );\n";
79 static const char *enableMSAAKernelLine =
80 "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n";
81 
verify(cl_int input,cl_int kernelOutput,const char * description)82 static int verify(cl_int input, cl_int kernelOutput, const char * description)
83 {
84   if( kernelOutput != input )
85   {
86     log_error( "ERROR: %s did not validate (expected %d, got %d)\n", description, input, kernelOutput);
87       return -1;
88   }
89   return 0;
90 }
91 
92 extern int supportsMsaa(cl_context context, bool* supports_msaa);
93 extern int supportsDepth(cl_context context, bool* supports_depth);
94 
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)95 int test_image_format_methods( cl_device_id device, cl_context context, cl_command_queue queue,
96                        size_t width, size_t height, size_t arraySize, size_t samples,
97                         GLenum target, format format, MTdata d )
98 {
99     int error, result=0;
100 
101     clProgramWrapper program;
102     clKernelWrapper kernel;
103     clMemWrapper image, outDataBuffer;
104     char programSrc[ 10240 ];
105 
106     image_kernel_data    outKernelData;
107 
108 #ifdef GL_VERSION_3_2
109     if (get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE ||
110         get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
111     {
112         bool supports_msaa;
113         error = supportsMsaa(context, &supports_msaa);
114         if( error != 0 ) return error;
115         if (!supports_msaa) return 0;
116     }
117     if (format.formattype == GL_DEPTH_COMPONENT ||
118         format.formattype == GL_DEPTH_STENCIL)
119     {
120         bool supports_depth;
121         error = supportsDepth(context, &supports_depth);
122         if( error != 0 ) return error;
123         if (!supports_depth) return 0;
124     }
125 #endif
126   DetectFloatToHalfRoundingMode(queue);
127 
128   glTextureWrapper glTexture;
129   switch (get_base_gl_target(target)) {
130     case GL_TEXTURE_2D:
131       CreateGLTexture2D( width, height, target,
132                         format.formattype, format.internal, format.datatype,
133                         format.type, &glTexture, &error, false, d );
134       break;
135     case GL_TEXTURE_2D_ARRAY:
136       CreateGLTexture2DArray( width, height, arraySize, target,
137                              format.formattype, format.internal, format.datatype,
138                              format.type, &glTexture, &error, false, d );
139       break;
140     case GL_TEXTURE_2D_MULTISAMPLE:
141       CreateGLTexture2DMultisample( width, height, samples, target,
142                                    format.formattype, format.internal, format.datatype,
143                                    format.type, &glTexture, &error, false, d, false);
144       break;
145     case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
146       CreateGLTexture2DArrayMultisample( width, height, arraySize, samples, target,
147                                         format.formattype, format.internal, format.datatype,
148                                         format.type, &glTexture, &error, false, d, false);
149       break;
150 
151     default:
152       log_error("Unsupported GL tex target (%s) passed to write test: "
153                 "%s (%s):%d", GetGLTargetName(target), __FUNCTION__,
154                 __FILE__, __LINE__);
155   }
156 
157   // Check to see if the texture could not be created for some other reason like
158   // GL_FRAMEBUFFER_UNSUPPORTED
159   if (error == GL_FRAMEBUFFER_UNSUPPORTED) {
160     return 0;
161   }
162 
163     // Construct testing source
164   log_info( " - Creating image %d by %d...\n", width, height );
165   // Create a CL image from the supplied GL texture
166   image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_READ_ONLY,
167                                         target, 0, glTexture, &error );
168 
169   if ( error != CL_SUCCESS ) {
170     print_error( error, "Unable to create CL image from GL texture" );
171     GLint fmt;
172     glGetTexLevelParameteriv( target, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt );
173     log_error( "    Supplied GL texture was base format %s and internal "
174               "format %s\n", GetGLBaseFormatName( fmt ), GetGLFormatName( fmt ) );
175     return error;
176   }
177 
178   cl_image_format imageFormat;
179   error = clGetImageInfo (image, CL_IMAGE_FORMAT,
180                           sizeof(imageFormat), &imageFormat, NULL);
181   test_error(error, "Failed to get image format");
182 
183   const char * imageType = 0;
184   bool doArraySize = false;
185   bool doImageWidth = false;
186   bool doImageHeight = false;
187   bool doImageChannelDataType = false;
188   bool doImageChannelOrder = false;
189   bool doImageDim = false;
190   bool doNumSamples = false;
191   bool doMSAA = false;
192   switch(target) {
193     case GL_TEXTURE_2D:
194       imageType = "image2d_depth_t";
195       doImageWidth = true;
196       doImageHeight = true;
197       doImageChannelDataType = true;
198       doImageChannelOrder = true;
199       doImageDim = true;
200       break;
201     case GL_TEXTURE_2D_ARRAY:
202       imageType = "image2d_array_depth_t";
203       doImageWidth = true;
204       doImageHeight = true;
205       doArraySize = true;
206       doImageChannelDataType = true;
207       doImageChannelOrder = true;
208       doImageDim = true;
209       doArraySize = true;
210       break;
211     case GL_TEXTURE_2D_MULTISAMPLE:
212       doNumSamples = true;
213       doMSAA = true;
214       if(format.formattype == GL_DEPTH_COMPONENT) {
215         doImageWidth = true;
216         imageType = "image2d_msaa_depth_t";
217       } else {
218         imageType = "image2d_msaa_t";
219       }
220       break;
221     case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
222       doMSAA = true;
223       if(format.formattype == GL_DEPTH_COMPONENT) {
224         doImageWidth = true;
225         imageType = "image2d_msaa_array_depth_t";
226       } else {
227         imageType = "image2d_array_msaa_t";
228       }
229       break;
230   }
231 
232 
233 
234   char channelTypeConstKernelLine[512] = {0};
235   char channelOrderConstKernelLine[512] = {0};
236   const char* channelTypeName=0;
237   const char* channelOrderName=0;
238   if(doImageChannelDataType) {
239     channelTypeName = GetChannelTypeName( imageFormat.image_channel_data_type );
240     if(channelTypeName && strlen(channelTypeName)) {
241       // replace CL_* with CLK_*
242       sprintf(channelTypeConstKernelLine, channelTypeConstLine, &channelTypeName[3]);
243     }
244   }
245   if(doImageChannelOrder) {
246     channelOrderName = GetChannelOrderName( imageFormat.image_channel_order );
247     if(channelOrderName && strlen(channelOrderName)) {
248       // replace CL_* with CLK_*
249       sprintf(channelOrderConstKernelLine, channelOrderConstLine, &channelOrderName[3]);
250     }
251   }
252 
253 	// Create a program to run against
254 	sprintf(programSrc,
255           methodTestKernelPattern,
256           ( doMSAA ) ? enableMSAAKernelLine : "",
257 	        imageType,
258           ( doArraySize ) ? arraySizeKernelLine : "",
259           ( doImageWidth ) ? imageWidthKernelLine : "",
260           ( doImageHeight ) ? imageHeightKernelLine : "",
261           ( doImageChannelDataType ) ? channelTypeKernelLine : "",
262           ( doImageChannelDataType ) ? channelTypeConstKernelLine : "",
263           ( doImageChannelOrder ) ? channelOrderKernelLine : "",
264           ( doImageChannelOrder ) ? channelOrderConstKernelLine : "",
265           ( doImageDim ) ? imageDimKernelLine : "",
266           ( doImageDim && doImageWidth ) ? imageWidthDimKernelLine : "",
267           ( doImageDim && doImageHeight ) ? imageHeightDimKernelLine : "",
268           ( doNumSamples ) ? numSamplesKernelLine : "");
269 
270 
271   //log_info("-----------------------------------\n%s\n", programSrc);
272   error = clFinish(queue);
273   if (error)
274     print_error(error, "clFinish failed.\n");
275     const char *ptr = programSrc;
276     error = create_single_kernel_helper( context, &program, &kernel, 1, &ptr, "sample_kernel" );
277     test_error( error, "Unable to create kernel to test against" );
278 
279     // Create an output buffer
280     outDataBuffer = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof( outKernelData ), NULL, &error );
281     test_error( error, "Unable to create output buffer" );
282 
283     // Set up arguments and run
284     error = clSetKernelArg( kernel, 0, sizeof( image ), &image );
285     test_error( error, "Unable to set kernel argument" );
286     error = clSetKernelArg( kernel, 1, sizeof( outDataBuffer ), &outDataBuffer );
287     test_error( error, "Unable to set kernel argument" );
288 
289   // Flush and Acquire.
290   glFlush();
291   error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &image, 0, NULL, NULL);
292   test_error( error, "Unable to acquire GL obejcts");
293 
294     size_t threads[1] = { 1 }, localThreads[1] = { 1 };
295 
296     error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
297     test_error( error, "Unable to run kernel" );
298 
299     error = clEnqueueReadBuffer( queue, outDataBuffer, CL_TRUE, 0, sizeof( outKernelData ), &outKernelData, 0, NULL, NULL );
300     test_error( error, "Unable to read data buffer" );
301 
302     // Verify the results now
303   if( doImageWidth )
304     result |= verify(width, outKernelData.width, "width");
305   if( doImageHeight)
306     result |= verify(height, outKernelData.height, "height");
307   if( doImageDim && doImageWidth )
308     result |= verify(width, outKernelData.widthDim, "width from get_image_dim");
309   if( doImageDim && doImageHeight )
310     result |= verify(height, outKernelData.heightDim, "height from get_image_dim");
311   if( doImageChannelDataType )
312     result |= verify(outKernelData.channelType, outKernelData.expectedChannelType, channelTypeName);
313   if( doImageChannelOrder )
314     result |= verify(outKernelData.channelOrder, outKernelData.expectedChannelOrder, channelOrderName);
315   if( doArraySize )
316     result |= verify(arraySize, outKernelData.arraySize, "array size");
317   if( doNumSamples )
318     result |= verify(samples, outKernelData.numSamples, "samples");
319   if(result) {
320     log_error("Test image methods failed");
321   }
322 
323   clEventWrapper event;
324   error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &image, 0, NULL, &event );
325   test_error(error, "clEnqueueReleaseGLObjects failed");
326 
327   error = clWaitForEvents( 1, &event );
328   test_error(error, "clWaitForEvents failed");
329 
330     return result;
331 }
332 
test_image_methods_depth(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)333 int test_image_methods_depth( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ){
334   if (!is_extension_available(device, "cl_khr_gl_depth_images")) {
335     log_info("Test not run because 'cl_khr_gl_depth_images' extension is not supported by the tested device\n");
336     return 0;
337   }
338 
339     size_t pixelSize;
340     int result = 0;
341   GLenum depth_targets[] = {GL_TEXTURE_2D, GL_TEXTURE_2D_ARRAY};
342   size_t ntargets = sizeof(depth_targets) / sizeof(depth_targets[0]);
343   size_t nformats = sizeof(depth_formats) / sizeof(depth_formats[0]);
344 
345   const size_t nsizes = 5;
346   sizevec_t sizes[nsizes];
347   // Need to limit texture size according to GL device properties
348   GLint maxTextureSize = 4096, maxTextureRectangleSize = 4096, maxTextureLayers = 16, size;
349   glGetIntegerv(GL_MAX_TEXTURE_SIZE, &maxTextureSize);
350   glGetIntegerv(GL_MAX_RECTANGLE_TEXTURE_SIZE_EXT, &maxTextureRectangleSize);
351   glGetIntegerv(GL_MAX_ARRAY_TEXTURE_LAYERS, &maxTextureLayers);
352 
353   size = min(maxTextureSize, maxTextureRectangleSize);
354 
355   RandomSeed seed( gRandomSeed );
356 
357   // Generate some random sizes (within reasonable ranges)
358   for (size_t i = 0; i < nsizes; i++) {
359     sizes[i].width  = random_in_range( 2, min(size, 1<<(i+4)), seed );
360     sizes[i].height = random_in_range( 2, min(size, 1<<(i+4)), seed );
361     sizes[i].depth  = random_in_range( 2, min(maxTextureLayers, 1<<(i+4)), seed );
362   }
363 
364   for (size_t i = 0; i < nsizes; i++) {
365     for(size_t itarget = 0; itarget < ntargets; ++itarget) {
366       for(size_t iformat = 0; iformat < nformats; ++iformat)
367         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,
368                                   depth_targets[itarget], depth_formats[iformat], seed );
369     }
370   }
371     return result;
372 }
373 
test_image_methods_multisample(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)374 int test_image_methods_multisample( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ){
375   if (!is_extension_available(device, "cl_khr_gl_msaa_sharing")) {
376     log_info("Test not run because 'cl_khr_gl_msaa_sharing' extension is not supported by the tested device\n");
377     return 0;
378   }
379 
380     size_t pixelSize;
381   int result = 0;
382   GLenum targets[] = {GL_TEXTURE_2D_MULTISAMPLE, GL_TEXTURE_2D_MULTISAMPLE_ARRAY};
383   size_t ntargets = sizeof(targets) / sizeof(targets[0]);
384   size_t nformats = sizeof(common_formats) / sizeof(common_formats[0]);
385 
386   const size_t nsizes = 5;
387   sizevec_t sizes[nsizes];
388   GLint maxTextureLayers = 16, maxTextureSize = 4096;
389   glGetIntegerv(GL_MAX_ARRAY_TEXTURE_LAYERS, &maxTextureLayers);
390   glGetIntegerv(GL_MAX_TEXTURE_SIZE, &maxTextureSize);
391 
392   RandomSeed seed( gRandomSeed );
393 
394   // Generate some random sizes (within reasonable ranges)
395   for (size_t i = 0; i < nsizes; i++) {
396     sizes[i].width  = random_in_range( 2, min(maxTextureSize, 1<<(i+4)), seed );
397     sizes[i].height = random_in_range( 2, min(maxTextureSize, 1<<(i+4)), seed );
398     sizes[i].depth  = random_in_range( 2, min(maxTextureLayers, 1<<(i+4)), seed );
399         }
400 
401   glEnable(GL_MULTISAMPLE);
402 
403   for (size_t i = 0; i < nsizes; i++) {
404     for(size_t itarget = 0; itarget < ntargets; ++itarget) {
405       for(size_t iformat = 0; iformat < nformats; ++iformat) {
406         GLint samples = get_gl_max_samples(targets[itarget], common_formats[iformat].internal);
407         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,
408                                   samples, targets[itarget], common_formats[iformat], seed );
409       }
410     }
411   }
412     return result;
413 }
414