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