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 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_READ_WRITE,
281 sizeof(outKernelData), NULL, &error);
282 test_error( error, "Unable to create output buffer" );
283
284 // Set up arguments and run
285 error = clSetKernelArg( kernel, 0, sizeof( image ), &image );
286 test_error( error, "Unable to set kernel argument" );
287 error = clSetKernelArg( kernel, 1, sizeof( outDataBuffer ), &outDataBuffer );
288 test_error( error, "Unable to set kernel argument" );
289
290 // Finish and Acquire.
291 glFinish();
292 error = (*clEnqueueAcquireGLObjects_ptr)(queue, 1, &image, 0, NULL, NULL);
293 test_error(error, "Unable to acquire GL obejcts");
294
295 size_t threads[1] = { 1 }, localThreads[1] = { 1 };
296
297 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
298 test_error( error, "Unable to run kernel" );
299
300 error = clEnqueueReadBuffer( queue, outDataBuffer, CL_TRUE, 0, sizeof( outKernelData ), &outKernelData, 0, NULL, NULL );
301 test_error( error, "Unable to read data buffer" );
302
303 // Verify the results now
304 if( doImageWidth )
305 result |= verify(width, outKernelData.width, "width");
306 if( doImageHeight)
307 result |= verify(height, outKernelData.height, "height");
308 if( doImageDim && doImageWidth )
309 result |= verify(width, outKernelData.widthDim, "width from get_image_dim");
310 if( doImageDim && doImageHeight )
311 result |= verify(height, outKernelData.heightDim, "height from get_image_dim");
312 if( doImageChannelDataType )
313 result |= verify(outKernelData.channelType, outKernelData.expectedChannelType, channelTypeName);
314 if( doImageChannelOrder )
315 result |= verify(outKernelData.channelOrder, outKernelData.expectedChannelOrder, channelOrderName);
316 if( doArraySize )
317 result |= verify(arraySize, outKernelData.arraySize, "array size");
318 if( doNumSamples )
319 result |= verify(samples, outKernelData.numSamples, "samples");
320 if(result) {
321 log_error("Test image methods failed");
322 }
323
324 clEventWrapper event;
325 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &image, 0, NULL, &event );
326 test_error(error, "clEnqueueReleaseGLObjects failed");
327
328 error = clWaitForEvents( 1, &event );
329 test_error(error, "clWaitForEvents failed");
330
331 return result;
332 }
333
test_image_methods_depth(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)334 int test_image_methods_depth( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ){
335 if (!is_extension_available(device, "cl_khr_gl_depth_images")) {
336 log_info("Test not run because 'cl_khr_gl_depth_images' extension is not supported by the tested device\n");
337 return 0;
338 }
339
340 size_t pixelSize;
341 int result = 0;
342 GLenum depth_targets[] = {GL_TEXTURE_2D, GL_TEXTURE_2D_ARRAY};
343 size_t ntargets = sizeof(depth_targets) / sizeof(depth_targets[0]);
344 size_t nformats = sizeof(depth_formats) / sizeof(depth_formats[0]);
345
346 const size_t nsizes = 5;
347 sizevec_t sizes[nsizes];
348 // Need to limit texture size according to GL device properties
349 GLint maxTextureSize = 4096, maxTextureRectangleSize = 4096, maxTextureLayers = 16, size;
350 glGetIntegerv(GL_MAX_TEXTURE_SIZE, &maxTextureSize);
351 glGetIntegerv(GL_MAX_RECTANGLE_TEXTURE_SIZE_EXT, &maxTextureRectangleSize);
352 glGetIntegerv(GL_MAX_ARRAY_TEXTURE_LAYERS, &maxTextureLayers);
353
354 size = min(maxTextureSize, maxTextureRectangleSize);
355
356 RandomSeed seed( gRandomSeed );
357
358 // Generate some random sizes (within reasonable ranges)
359 for (size_t i = 0; i < nsizes; i++) {
360 sizes[i].width = random_in_range( 2, min(size, 1<<(i+4)), seed );
361 sizes[i].height = random_in_range( 2, min(size, 1<<(i+4)), seed );
362 sizes[i].depth = random_in_range( 2, min(maxTextureLayers, 1<<(i+4)), seed );
363 }
364
365 for (size_t i = 0; i < nsizes; i++) {
366 for(size_t itarget = 0; itarget < ntargets; ++itarget) {
367 for(size_t iformat = 0; iformat < nformats; ++iformat)
368 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,
369 depth_targets[itarget], depth_formats[iformat], seed );
370 }
371 }
372 return result;
373 }
374
test_image_methods_multisample(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)375 int test_image_methods_multisample( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ){
376 if (!is_extension_available(device, "cl_khr_gl_msaa_sharing")) {
377 log_info("Test not run because 'cl_khr_gl_msaa_sharing' extension is not supported by the tested device\n");
378 return 0;
379 }
380
381 size_t pixelSize;
382 int result = 0;
383 GLenum targets[] = {GL_TEXTURE_2D_MULTISAMPLE, GL_TEXTURE_2D_MULTISAMPLE_ARRAY};
384 size_t ntargets = sizeof(targets) / sizeof(targets[0]);
385 size_t nformats = sizeof(common_formats) / sizeof(common_formats[0]);
386
387 const size_t nsizes = 5;
388 sizevec_t sizes[nsizes];
389 GLint maxTextureLayers = 16, maxTextureSize = 4096;
390 glGetIntegerv(GL_MAX_ARRAY_TEXTURE_LAYERS, &maxTextureLayers);
391 glGetIntegerv(GL_MAX_TEXTURE_SIZE, &maxTextureSize);
392
393 RandomSeed seed( gRandomSeed );
394
395 // Generate some random sizes (within reasonable ranges)
396 for (size_t i = 0; i < nsizes; i++) {
397 sizes[i].width = random_in_range( 2, min(maxTextureSize, 1<<(i+4)), seed );
398 sizes[i].height = random_in_range( 2, min(maxTextureSize, 1<<(i+4)), seed );
399 sizes[i].depth = random_in_range( 2, min(maxTextureLayers, 1<<(i+4)), seed );
400 }
401
402 glEnable(GL_MULTISAMPLE);
403
404 for (size_t i = 0; i < nsizes; i++) {
405 for(size_t itarget = 0; itarget < ntargets; ++itarget) {
406 for(size_t iformat = 0; iformat < nformats; ++iformat) {
407 GLint samples = get_gl_max_samples(targets[itarget], common_formats[iformat].internal);
408 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,
409 samples, targets[itarget], common_formats[iformat], seed );
410 }
411 }
412 }
413 return result;
414 }
415