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