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 "testBase.h"
17 #include "common.h"
18 #include <limits.h>
19
20 #if defined(__APPLE__)
21 #include <OpenGL/glu.h>
22 #else
23 #include <GL/glu.h>
24 #include <CL/cl_gl.h>
25 #endif
26
27 #pragma mark -
28 #pragma mark Write test kernels
29
30 // clang-format off
31 static const char *kernelpattern_image_write_1D =
32 "__kernel void sample_test( __global %s4 *source, write_only image1d_t dest )\n"
33 "{\n"
34 " uint index = get_global_id(0);\n"
35 " %s4 value = source[index];\n"
36 " write_image%s( dest, index, %s(value));\n"
37 "}\n";
38
39 static const char *kernelpattern_image_write_1D_half =
40 "__kernel void sample_test( __global half4 *source, write_only image1d_t dest )\n"
41 "{\n"
42 " uint index = get_global_id(0);\n"
43 " write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
44 "}\n";
45
46 static const char *kernelpattern_image_write_1D_buffer =
47 "__kernel void sample_test( __global %s4 *source, write_only image1d_buffer_t dest )\n"
48 "{\n"
49 " uint index = get_global_id(0);\n"
50 " %s4 value = source[index];\n"
51 " write_image%s( dest, index, %s(value));\n"
52 "}\n";
53
54 static const char *kernelpattern_image_write_1D_buffer_half =
55 "__kernel void sample_test( __global half4 *source, write_only image1d_buffer_t dest )\n"
56 "{\n"
57 " uint index = get_global_id(0);\n"
58 " write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
59 "}\n";
60
61 static const char *kernelpattern_image_write_2D =
62 "__kernel void sample_test( __global %s4 *source, write_only image2d_t dest )\n"
63 "{\n"
64 " int tidX = get_global_id(0);\n"
65 " int tidY = get_global_id(1);\n"
66 " uint index = tidY * get_image_width( dest ) + tidX;\n"
67 " %s4 value = source[index];\n"
68 " write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
69 "}\n";
70
71 static const char *kernelpattern_image_write_2D_half =
72 "__kernel void sample_test( __global half4 *source, write_only image2d_t dest )\n"
73 "{\n"
74 " int tidX = get_global_id(0);\n"
75 " int tidY = get_global_id(1);\n"
76 " uint index = tidY * get_image_width( dest ) + tidX;\n"
77 " write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
78 "}\n";
79
80 static const char *kernelpattern_image_write_1Darray =
81 "__kernel void sample_test( __global %s4 *source, write_only image1d_array_t dest )\n"
82 "{\n"
83 " int tidX = get_global_id(0);\n"
84 " int tidY = get_global_id(1);\n"
85 " uint index = tidY * get_image_width( dest ) + tidX;\n"
86 " %s4 value = source[index];\n"
87 " write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
88 "}\n";
89
90 static const char *kernelpattern_image_write_1Darray_half =
91 "__kernel void sample_test( __global half4 *source, write_only image1d_array_t dest )\n"
92 "{\n"
93 " int tidX = get_global_id(0);\n"
94 " int tidY = get_global_id(1);\n"
95 " uint index = tidY * get_image_width( dest ) + tidX;\n"
96 " write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
97 "}\n";
98
99 static const char *kernelpattern_image_write_3D =
100 "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
101 "__kernel void sample_test( __global %s4 *source, write_only image3d_t dest )\n"
102 "{\n"
103 " int tidX = get_global_id(0);\n"
104 " int tidY = get_global_id(1);\n"
105 " int tidZ = get_global_id(2);\n"
106 " int width = get_image_width( dest );\n"
107 " int height = get_image_height( dest );\n"
108 " int index = tidZ * width * height + tidY * width + tidX;\n"
109 " %s4 value = source[index];\n"
110 " write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
111 "}\n";
112
113 static const char *kernelpattern_image_write_3D_half =
114 "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
115 "__kernel void sample_test( __global half4 *source, write_only image3d_t dest )\n"
116 "{\n"
117 " int tidX = get_global_id(0);\n"
118 " int tidY = get_global_id(1);\n"
119 " int tidZ = get_global_id(2);\n"
120 " int width = get_image_width( dest );\n"
121 " int height = get_image_height( dest );\n"
122 " int index = tidZ * width * height + tidY * width + tidX;\n"
123 " write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
124 "}\n";
125
126 static const char *kernelpattern_image_write_2Darray =
127 "__kernel void sample_test( __global %s4 *source, write_only image2d_array_t dest )\n"
128 "{\n"
129 " int tidX = get_global_id(0);\n"
130 " int tidY = get_global_id(1);\n"
131 " int tidZ = get_global_id(2);\n"
132 " int width = get_image_width( dest );\n"
133 " int height = get_image_height( dest );\n"
134 " int index = tidZ * width * height + tidY * width + tidX;\n"
135 " %s4 value = source[index];\n"
136 " write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
137 "}\n";
138
139 static const char *kernelpattern_image_write_2Darray_half =
140 "__kernel void sample_test( __global half4 *source, write_only image2d_array_t dest )\n"
141 "{\n"
142 " int tidX = get_global_id(0);\n"
143 " int tidY = get_global_id(1);\n"
144 " int tidZ = get_global_id(2);\n"
145 " int width = get_image_width( dest );\n"
146 " int height = get_image_height( dest );\n"
147 " int index = tidZ * width * height + tidY * width + tidX;\n"
148 " write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
149 "}\n";
150
151 #ifdef GL_VERSION_3_2
152
153 static const char * kernelpattern_image_write_2D_depth =
154 "__kernel void sample_test( __global %s *source, write_only image2d_depth_t dest )\n"
155 "{\n"
156 " int tidX = get_global_id(0);\n"
157 " int tidY = get_global_id(1);\n"
158 " uint index = tidY * get_image_width( dest ) + tidX;\n"
159 " float value = source[index];\n"
160 " write_imagef( dest, (int2)( tidX, tidY ), value);\n"
161 "}\n";
162
163 static const char * kernelpattern_image_write_2D_array_depth =
164 "__kernel void sample_test( __global %s *source, write_only image2d_array_depth_t dest )\n"
165 "{\n"
166 " int tidX = get_global_id(0);\n"
167 " int tidY = get_global_id(1);\n"
168 " int tidZ = get_global_id(2);\n"
169 " int width = get_image_width( dest );\n"
170 " int height = get_image_height( dest );\n"
171 " int index = tidZ * width * height + tidY * width + tidX;\n"
172 " %s value = source[index];\n"
173 " write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
174 "}\n";
175
176
177 #endif
178 // clang-format on
179
180 #pragma mark -
181 #pragma mark Utility functions
182
get_appropriate_write_kernel(GLenum target,ExplicitType type,cl_channel_order channel_order)183 static const char *get_appropriate_write_kernel(GLenum target,
184 ExplicitType type,
185 cl_channel_order channel_order)
186 {
187 switch (get_base_gl_target(target))
188 {
189 case GL_TEXTURE_1D:
190
191 if (type == kHalf)
192 return kernelpattern_image_write_1D_half;
193 else
194 return kernelpattern_image_write_1D;
195 break;
196 case GL_TEXTURE_BUFFER:
197 if (type == kHalf)
198 return kernelpattern_image_write_1D_buffer_half;
199 else
200 return kernelpattern_image_write_1D_buffer;
201 break;
202 case GL_TEXTURE_1D_ARRAY:
203 if (type == kHalf)
204 return kernelpattern_image_write_1Darray_half;
205 else
206 return kernelpattern_image_write_1Darray;
207 break;
208 case GL_COLOR_ATTACHMENT0:
209 case GL_RENDERBUFFER:
210 case GL_TEXTURE_RECTANGLE_EXT:
211 case GL_TEXTURE_2D:
212 case GL_TEXTURE_CUBE_MAP:
213 #ifdef GL_VERSION_3_2
214 if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
215 return kernelpattern_image_write_2D_depth;
216 #endif
217 if (type == kHalf)
218 return kernelpattern_image_write_2D_half;
219 else
220 return kernelpattern_image_write_2D;
221 break;
222
223 case GL_TEXTURE_2D_ARRAY:
224 #ifdef GL_VERSION_3_2
225 if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
226 return kernelpattern_image_write_2D_array_depth;
227 #endif
228 if (type == kHalf)
229 return kernelpattern_image_write_2Darray_half;
230 else
231 return kernelpattern_image_write_2Darray;
232 break;
233
234 case GL_TEXTURE_3D:
235 if (type == kHalf)
236 return kernelpattern_image_write_3D_half;
237 else
238 return kernelpattern_image_write_3D;
239 break;
240
241 default:
242 log_error("Unsupported GL tex target (%s) passed to write test: "
243 "%s (%s):%d",
244 GetGLTargetName(target), __FUNCTION__, __FILE__,
245 __LINE__);
246 return NULL;
247 }
248 }
249
set_dimensions_by_target(GLenum target,size_t * dims,size_t sizes[3],size_t width,size_t height,size_t depth)250 void set_dimensions_by_target(GLenum target, size_t *dims, size_t sizes[3],
251 size_t width, size_t height, size_t depth)
252 {
253 switch (get_base_gl_target(target))
254 {
255 case GL_TEXTURE_1D:
256 sizes[0] = width;
257 *dims = 1;
258 break;
259
260 case GL_TEXTURE_BUFFER:
261 sizes[0] = width;
262 *dims = 1;
263 break;
264
265 case GL_TEXTURE_1D_ARRAY:
266 sizes[0] = width;
267 sizes[1] = height;
268 *dims = 2;
269 break;
270
271 case GL_COLOR_ATTACHMENT0:
272 case GL_RENDERBUFFER:
273 case GL_TEXTURE_RECTANGLE_EXT:
274 case GL_TEXTURE_2D:
275 case GL_TEXTURE_CUBE_MAP:
276
277 sizes[0] = width;
278 sizes[1] = height;
279 *dims = 2;
280 break;
281
282 case GL_TEXTURE_2D_ARRAY:
283 sizes[0] = width;
284 sizes[1] = height;
285 sizes[2] = depth;
286 *dims = 3;
287 break;
288
289 case GL_TEXTURE_3D:
290 sizes[0] = width;
291 sizes[1] = height;
292 sizes[2] = depth;
293 *dims = 3;
294 break;
295
296 default:
297 log_error("Unsupported GL tex target (%s) passed to write test: "
298 "%s (%s):%d",
299 GetGLTargetName(target), __FUNCTION__, __FILE__,
300 __LINE__);
301 }
302 }
303
test_cl_image_write(cl_context context,cl_command_queue queue,GLenum target,cl_mem clImage,size_t width,size_t height,size_t depth,cl_image_format * outFormat,ExplicitType * outType,void ** outSourceBuffer,MTdata d,bool supports_half)304 int test_cl_image_write(cl_context context, cl_command_queue queue,
305 GLenum target, cl_mem clImage, size_t width,
306 size_t height, size_t depth, cl_image_format *outFormat,
307 ExplicitType *outType, void **outSourceBuffer, MTdata d,
308 bool supports_half)
309 {
310 size_t global_dims, global_sizes[3];
311 clProgramWrapper program;
312 clKernelWrapper kernel;
313 clMemWrapper inStream;
314 char *programPtr;
315 int error;
316 char kernelSource[2048];
317
318 // What CL format did we get from the texture?
319
320 error = clGetImageInfo(clImage, CL_IMAGE_FORMAT, sizeof(cl_image_format),
321 outFormat, NULL);
322 test_error(error, "Unable to get the CL image format");
323
324 // Create the kernel source. The target and the data type will influence
325 // which particular kernel we choose.
326
327 *outType = get_write_kernel_type(outFormat);
328 size_t channelSize = get_explicit_type_size(*outType);
329
330 const char *appropriateKernel = get_appropriate_write_kernel(
331 target, *outType, outFormat->image_channel_order);
332 if (*outType == kHalf && !supports_half)
333 {
334 log_info("cl_khr_fp16 isn't supported. Skip this test.\n");
335 return 0;
336 }
337
338 const char *suffix = get_kernel_suffix(outFormat);
339 const char *convert = get_write_conversion(outFormat, *outType);
340
341 sprintf(kernelSource, appropriateKernel, get_explicit_type_name(*outType),
342 get_explicit_type_name(*outType), suffix, convert);
343
344 programPtr = kernelSource;
345 if (create_single_kernel_helper_with_build_options(
346 context, &program, &kernel, 1, (const char **)&programPtr,
347 "sample_test", ""))
348 {
349 return -1;
350 }
351
352 // Create an appropriately-sized output buffer.
353
354 // Check to see if the output buffer will fit on the device
355 size_t bytes = channelSize * 4 * width * height * depth;
356 cl_ulong alloc_size = 0;
357
358 cl_device_id device = NULL;
359 error = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(device),
360 &device, NULL);
361 test_error(error, "Unable to query command queue for device");
362
363 error = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
364 sizeof(alloc_size), &alloc_size, NULL);
365 test_error(error, "Unable to device for max mem alloc size");
366
367 if (bytes > alloc_size)
368 {
369 log_info(" Skipping: Buffer size (%lu) is greater than "
370 "CL_DEVICE_MAX_MEM_ALLOC_SIZE (%lu)\n",
371 bytes, alloc_size);
372 *outSourceBuffer = NULL;
373 return 0;
374 }
375
376 *outSourceBuffer =
377 CreateRandomData(*outType, width * height * depth * 4, d);
378
379 inStream = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
380 channelSize * 4 * width * height * depth,
381 *outSourceBuffer, &error);
382 test_error(error, "Unable to create output buffer");
383
384 clSamplerWrapper sampler = clCreateSampler(
385 context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error);
386 test_error(error, "Unable to create sampler");
387
388 error = clSetKernelArg(kernel, 0, sizeof(inStream), &inStream);
389 test_error(error, "Unable to set kernel arguments");
390
391 error = clSetKernelArg(kernel, 1, sizeof(clImage), &clImage);
392 test_error(error, "Unable to set kernel arguments");
393
394 // Flush and Acquire.
395
396 glFinish();
397
398 error = (*clEnqueueAcquireGLObjects_ptr)(queue, 1, &clImage, 0, NULL, NULL);
399 test_error(error, "Unable to acquire GL obejcts");
400
401 // Execute ( letting OpenCL choose the local size )
402
403 // Setup the global dimensions and sizes based on the target type.
404 set_dimensions_by_target(target, &global_dims, global_sizes, width, height,
405 depth);
406
407 error = clEnqueueNDRangeKernel(queue, kernel, global_dims, NULL,
408 global_sizes, NULL, 0, NULL, NULL);
409 test_error(error, "Unable to execute test kernel");
410
411 clEventWrapper event;
412 error =
413 (*clEnqueueReleaseGLObjects_ptr)(queue, 1, &clImage, 0, NULL, &event);
414 test_error(error, "clEnqueueReleaseGLObjects failed");
415
416 error = clWaitForEvents(1, &event);
417 test_error(error, "clWaitForEvents failed");
418
419 return 0;
420 }
421
test_image_write(cl_context context,cl_command_queue queue,GLenum glTarget,GLuint glTexture,size_t width,size_t height,size_t depth,cl_image_format * outFormat,ExplicitType * outType,void ** outSourceBuffer,MTdata d,bool supports_half)422 static int test_image_write(cl_context context, cl_command_queue queue,
423 GLenum glTarget, GLuint glTexture, size_t width,
424 size_t height, size_t depth,
425 cl_image_format *outFormat, ExplicitType *outType,
426 void **outSourceBuffer, MTdata d,
427 bool supports_half)
428 {
429 int error;
430
431 // Create a CL image from the supplied GL texture
432 clMemWrapper image = (*clCreateFromGLTexture_ptr)(
433 context, CL_MEM_WRITE_ONLY, glTarget, 0, glTexture, &error);
434
435 if (error != CL_SUCCESS)
436 {
437 print_error(error, "Unable to create CL image from GL texture");
438 GLint fmt;
439 glGetTexLevelParameteriv(glTarget, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt);
440 log_error(" Supplied GL texture was base format %s and internal "
441 "format %s\n",
442 GetGLBaseFormatName(fmt), GetGLFormatName(fmt));
443 return error;
444 }
445
446 return test_cl_image_write(context, queue, glTarget, image, width, height,
447 depth, outFormat, outType, outSourceBuffer, d,
448 supports_half);
449 }
450
supportsHalf(cl_context context,bool * supports_half)451 int supportsHalf(cl_context context, bool *supports_half)
452 {
453 int error;
454 cl_uint numDev;
455
456 error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint),
457 &numDev, NULL);
458 test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");
459
460 cl_device_id *devices = new cl_device_id[numDev];
461 error = clGetContextInfo(context, CL_CONTEXT_DEVICES,
462 numDev * sizeof(cl_device_id), devices, NULL);
463 test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");
464
465 *supports_half = is_extension_available(devices[0], "cl_khr_fp16");
466 delete[] devices;
467
468 return error;
469 }
470
supportsMsaa(cl_context context,bool * supports_msaa)471 int supportsMsaa(cl_context context, bool *supports_msaa)
472 {
473 int error;
474 cl_uint numDev;
475
476 error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint),
477 &numDev, NULL);
478 test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");
479
480 cl_device_id *devices = new cl_device_id[numDev];
481 error = clGetContextInfo(context, CL_CONTEXT_DEVICES,
482 numDev * sizeof(cl_device_id), devices, NULL);
483 test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");
484
485 *supports_msaa =
486 is_extension_available(devices[0], "cl_khr_gl_msaa_sharing");
487 delete[] devices;
488
489 return error;
490 }
491
supportsDepth(cl_context context,bool * supports_depth)492 int supportsDepth(cl_context context, bool *supports_depth)
493 {
494 int error;
495 cl_uint numDev;
496
497 error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint),
498 &numDev, NULL);
499 test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");
500
501 cl_device_id *devices = new cl_device_id[numDev];
502 error = clGetContextInfo(context, CL_CONTEXT_DEVICES,
503 numDev * sizeof(cl_device_id), devices, NULL);
504 test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");
505
506 *supports_depth =
507 is_extension_available(devices[0], "cl_khr_gl_depth_images");
508 delete[] devices;
509
510 return error;
511 }
512
test_image_format_write(cl_context context,cl_command_queue queue,size_t width,size_t height,size_t depth,GLenum target,GLenum format,GLenum internalFormat,GLenum glType,ExplicitType type,MTdata d)513 static int test_image_format_write(cl_context context, cl_command_queue queue,
514 size_t width, size_t height, size_t depth,
515 GLenum target, GLenum format,
516 GLenum internalFormat, GLenum glType,
517 ExplicitType type, MTdata d)
518 {
519 int error;
520 // If we're testing a half float format, then we need to determine the
521 // rounding mode of this machine. Punt if we fail to do so.
522
523 if (type == kHalf)
524 if (DetectFloatToHalfRoundingMode(queue)) return 1;
525
526 // Create an appropriate GL texture or renderbuffer, given the target.
527
528 glTextureWrapper glTexture;
529 glBufferWrapper glBuf;
530 glFramebufferWrapper glFramebuffer;
531 glRenderbufferWrapper glRenderbuffer;
532 switch (get_base_gl_target(target))
533 {
534 case GL_TEXTURE_1D:
535 CreateGLTexture1D(width, target, format, internalFormat, glType,
536 type, &glTexture, &error, false, d);
537 break;
538 case GL_TEXTURE_BUFFER:
539 CreateGLTextureBuffer(width, target, format, internalFormat, glType,
540 type, &glTexture, &glBuf, &error, false, d);
541 break;
542 case GL_TEXTURE_1D_ARRAY:
543 CreateGLTexture1DArray(width, height, target, format,
544 internalFormat, glType, type, &glTexture,
545 &error, false, d);
546 break;
547 case GL_TEXTURE_RECTANGLE_EXT:
548 case GL_TEXTURE_2D:
549 case GL_TEXTURE_CUBE_MAP:
550 CreateGLTexture2D(width, height, target, format, internalFormat,
551 glType, type, &glTexture, &error, false, d);
552 break;
553 case GL_COLOR_ATTACHMENT0:
554 case GL_RENDERBUFFER:
555 CreateGLRenderbuffer(width, height, target, format, internalFormat,
556 glType, type, &glFramebuffer, &glRenderbuffer,
557 &error, d, false);
558 case GL_TEXTURE_2D_ARRAY:
559 CreateGLTexture2DArray(width, height, depth, target, format,
560 internalFormat, glType, type, &glTexture,
561 &error, false, d);
562 break;
563 case GL_TEXTURE_3D:
564 CreateGLTexture3D(width, height, depth, target, format,
565 internalFormat, glType, type, &glTexture, &error,
566 d, false);
567 break;
568
569 default:
570 log_error("Unsupported GL tex target (%s) passed to write test: "
571 "%s (%s):%d",
572 GetGLTargetName(target), __FUNCTION__, __FILE__,
573 __LINE__);
574 return -1;
575 }
576
577 // If there was a problem during creation, make sure it isn't a known
578 // cause, and then complain.
579 if (error == -2)
580 {
581 log_info("OpenGL texture couldn't be created, because a texture is too "
582 "big. Skipping test.\n");
583 return 0;
584 }
585
586 if (error != 0)
587 {
588 if ((format == GL_RGBA_INTEGER_EXT)
589 && (!CheckGLIntegerExtensionSupport()))
590 {
591 log_info("OpenGL version does not support GL_RGBA_INTEGER_EXT. "
592 "Skipping test.\n");
593 return 0;
594 }
595 else
596 {
597 return error;
598 }
599 }
600
601 // Run and get the results
602 cl_image_format clFormat;
603 ExplicitType sourceType;
604 ExplicitType validationType;
605 void *outSourceBuffer = NULL;
606
607 GLenum globj = glTexture;
608 if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0)
609 {
610 globj = glRenderbuffer;
611 }
612
613 bool supports_half = false;
614 error = supportsHalf(context, &supports_half);
615 if (error != 0) return error;
616
617 error = test_image_write(context, queue, target, globj, width, height,
618 depth, &clFormat, &sourceType,
619 (void **)&outSourceBuffer, d, supports_half);
620
621 if (error != 0 || ((sourceType == kHalf) && !supports_half))
622 {
623 if (outSourceBuffer) free(outSourceBuffer);
624 return error;
625 }
626
627 if (!outSourceBuffer) return 0;
628
629 // If actual source type was half, convert to float for validation.
630
631 if (sourceType == kHalf)
632 validationType = kFloat;
633 else
634 validationType = sourceType;
635
636 BufferOwningPtr<char> validationSource;
637
638 if (clFormat.image_channel_data_type == CL_UNORM_INT_101010)
639 {
640 validationSource.reset(outSourceBuffer);
641 }
642 else
643 {
644 validationSource.reset(convert_to_expected(
645 outSourceBuffer, width * height * depth, sourceType, validationType,
646 get_channel_order_channel_count(clFormat.image_channel_order)));
647 free(outSourceBuffer);
648 }
649
650 log_info(
651 "- Write for %s [%4ld x %4ld x %4ld] : GL Texture : %s : %s : %s =>"
652 " CL Image : %s : %s \n",
653 GetGLTargetName(target), width, height, depth, GetGLFormatName(format),
654 GetGLFormatName(internalFormat), GetGLTypeName(glType),
655 GetChannelOrderName(clFormat.image_channel_order),
656 GetChannelTypeName(clFormat.image_channel_data_type));
657
658 // Read the results from the GL texture.
659
660 ExplicitType readType = type;
661 BufferOwningPtr<char> glResults(
662 ReadGLTexture(target, glTexture, glBuf, width, format, internalFormat,
663 glType, readType, /* unused */ 1, 1));
664 if (glResults == NULL) return -1;
665
666 // We have to convert our input buffer to the returned type, so we can
667 // validate.
668 BufferOwningPtr<char> convertedGLResults;
669 if (clFormat.image_channel_data_type != CL_UNORM_INT_101010)
670 {
671 convertedGLResults.reset(convert_to_expected(
672 glResults, width * height * depth, readType, validationType,
673 get_channel_order_channel_count(clFormat.image_channel_order),
674 glType));
675 }
676
677 // Validate.
678
679 int valid = 0;
680 if (convertedGLResults)
681 {
682 if (sourceType == kFloat || sourceType == kHalf)
683 {
684 if (clFormat.image_channel_data_type == CL_UNORM_INT_101010)
685 {
686 valid = validate_float_results_rgb_101010(
687 validationSource, glResults, width, height, depth, 1);
688 }
689 else
690 {
691 valid =
692 validate_float_results(validationSource, convertedGLResults,
693 width, height, depth, 1,
694 get_channel_order_channel_count(
695 clFormat.image_channel_order));
696 }
697 }
698 else
699 {
700 valid = validate_integer_results(
701 validationSource, convertedGLResults, width, height, depth, 1,
702 get_explicit_type_size(readType));
703 }
704 }
705
706 return valid;
707 }
708
709 #pragma mark -
710 #pragma mark Write test common entry point
711
712 // This is the main loop for all of the write tests. It iterates over the
713 // given formats & targets, testing a variety of sizes against each
714 // combination.
715
test_images_write_common(cl_device_id device,cl_context context,cl_command_queue queue,const format * formats,size_t nformats,GLenum * targets,size_t ntargets,sizevec_t * sizes,size_t nsizes)716 int test_images_write_common(cl_device_id device, cl_context context,
717 cl_command_queue queue, const format *formats,
718 size_t nformats, GLenum *targets, size_t ntargets,
719 sizevec_t *sizes, size_t nsizes)
720 {
721 int err = 0;
722 int error = 0;
723 RandomSeed seed(gRandomSeed);
724
725 // First, ensure this device supports images.
726
727 if (checkForImageSupport(device))
728 {
729 log_info("Device does not support images. Skipping test.\n");
730 return 0;
731 }
732
733 // Get the value of CL_DEVICE_MAX_MEM_ALLOC_SIZE
734 cl_ulong max_individual_allocation_size = 0;
735 err = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
736 sizeof(max_individual_allocation_size),
737 &max_individual_allocation_size, NULL);
738 if (err)
739 {
740 log_error("ERROR: clGetDeviceInfo failed for "
741 "CL_DEVICE_MAX_MEM_ALLOC_SIZE.\n");
742 error++;
743 return error;
744 }
745
746 size_t total_allocation_size;
747 size_t fidx, tidx, sidx;
748
749 for (fidx = 0; fidx < nformats; fidx++)
750 {
751 for (tidx = 0; tidx < ntargets; tidx++)
752 {
753
754 // Texture buffer only takes an internal format, so the level data
755 // passed by the test and used for verification must match the
756 // internal format
757 if ((targets[tidx] == GL_TEXTURE_BUFFER)
758 && (GetGLFormat(formats[fidx].internal)
759 != formats[fidx].formattype))
760 continue;
761
762 if (formats[fidx].datatype == GL_UNSIGNED_INT_2_10_10_10_REV)
763 {
764 // Check if the RGB 101010 format is supported
765 if (is_rgb_101010_supported(context, targets[tidx]) == 0)
766 continue; // skip
767 }
768
769 if (formats[fidx].datatype == GL_UNSIGNED_INT_24_8)
770 {
771 // check if a implementation supports writing to the depth
772 // stencil formats
773 cl_image_format imageFormat = { CL_DEPTH_STENCIL,
774 CL_UNORM_INT24 };
775 if (!is_image_format_supported(
776 context, CL_MEM_WRITE_ONLY,
777 (targets[tidx] == GL_TEXTURE_2D
778 || targets[tidx] == GL_TEXTURE_RECTANGLE)
779 ? CL_MEM_OBJECT_IMAGE2D
780 : CL_MEM_OBJECT_IMAGE2D_ARRAY,
781 &imageFormat))
782 continue;
783 }
784
785 if (formats[fidx].datatype == GL_FLOAT_32_UNSIGNED_INT_24_8_REV)
786 {
787 // check if a implementation supports writing to the depth
788 // stencil formats
789 cl_image_format imageFormat = { CL_DEPTH_STENCIL, CL_FLOAT };
790 if (!is_image_format_supported(
791 context, CL_MEM_WRITE_ONLY,
792 (targets[tidx] == GL_TEXTURE_2D
793 || targets[tidx] == GL_TEXTURE_RECTANGLE)
794 ? CL_MEM_OBJECT_IMAGE2D
795 : CL_MEM_OBJECT_IMAGE2D_ARRAY,
796 &imageFormat))
797 continue;
798 }
799
800 if (targets[tidx] != GL_TEXTURE_BUFFER)
801 log_info(
802 "Testing image write for GL format %s : %s : %s : %s\n",
803 GetGLTargetName(targets[tidx]),
804 GetGLFormatName(formats[fidx].internal),
805 GetGLBaseFormatName(formats[fidx].formattype),
806 GetGLTypeName(formats[fidx].datatype));
807 else
808 log_info("Testing image write for GL format %s : %s\n",
809 GetGLTargetName(targets[tidx]),
810 GetGLFormatName(formats[fidx].internal));
811
812
813 for (sidx = 0; sidx < nsizes; sidx++)
814 {
815
816 // All tested formats are 4-channel formats
817 total_allocation_size = sizes[sidx].width * sizes[sidx].height
818 * sizes[sidx].depth * 4
819 * get_explicit_type_size(formats[fidx].type);
820
821 if (total_allocation_size > max_individual_allocation_size)
822 {
823 log_info("The requested allocation size (%gMB) is larger "
824 "than the "
825 "maximum individual allocation size (%gMB)\n",
826 total_allocation_size / (1024.0 * 1024.0),
827 max_individual_allocation_size
828 / (1024.0 * 1024.0));
829 log_info("Skipping write test for %s : %s : %s : %s "
830 " and size (%ld, %ld, %ld)\n",
831 GetGLTargetName(targets[tidx]),
832 GetGLFormatName(formats[fidx].internal),
833 GetGLBaseFormatName(formats[fidx].formattype),
834 GetGLTypeName(formats[fidx].datatype),
835 sizes[sidx].width, sizes[sidx].height,
836 sizes[sidx].depth);
837 continue;
838 }
839 #ifdef GL_VERSION_3_2
840 if (get_base_gl_target(targets[tidx])
841 == GL_TEXTURE_2D_MULTISAMPLE
842 || get_base_gl_target(targets[tidx])
843 == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
844 {
845 bool supports_msaa;
846 int errorInGetInfo = supportsMsaa(context, &supports_msaa);
847 if (errorInGetInfo != 0) return errorInGetInfo;
848 if (!supports_msaa) return 0;
849 }
850 if (formats[fidx].formattype == GL_DEPTH_COMPONENT
851 || formats[fidx].formattype == GL_DEPTH_STENCIL)
852 {
853 bool supports_depth;
854 int errorInGetInfo =
855 supportsDepth(context, &supports_depth);
856 if (errorInGetInfo != 0) return errorInGetInfo;
857 if (!supports_depth) return 0;
858 }
859 #endif
860
861 if (test_image_format_write(
862 context, queue, sizes[sidx].width, sizes[sidx].height,
863 sizes[sidx].depth, targets[tidx],
864 formats[fidx].formattype, formats[fidx].internal,
865 formats[fidx].datatype, formats[fidx].type, seed))
866 {
867 log_error(
868 "ERROR: Image write test failed for %s : %s : %s : %s "
869 " and size (%ld, %ld, %ld)\n\n",
870 GetGLTargetName(targets[tidx]),
871 GetGLFormatName(formats[fidx].internal),
872 GetGLBaseFormatName(formats[fidx].formattype),
873 GetGLTypeName(formats[fidx].datatype),
874 sizes[sidx].width, sizes[sidx].height,
875 sizes[sidx].depth);
876
877 error++;
878 break; // Skip other sizes for this combination
879 }
880 }
881
882 // If we passed all sizes (check versus size loop count):
883
884 if (sidx == nsizes)
885 {
886 log_info(
887 "passed: Image write for GL format %s : %s : %s : %s\n\n",
888 GetGLTargetName(targets[tidx]),
889 GetGLFormatName(formats[fidx].internal),
890 GetGLBaseFormatName(formats[fidx].formattype),
891 GetGLTypeName(formats[fidx].datatype));
892 }
893 }
894 }
895
896 return error;
897 }
898