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 #include "testBase.h"
18
19 #if defined( __APPLE__ )
20 #include <OpenGL/glu.h>
21 #else
22 #include <GL/glu.h>
23 #include <CL/cl_gl.h>
24 #endif
25
26 extern int supportsHalf(cl_context context, bool* supports_half);
27 extern int supportsMsaa(cl_context context, bool* supports_msaa);
28 extern int supportsDepth(cl_context context, bool* supports_depth);
29
30 static const char *kernelpattern_image_read_1d =
31 "__kernel void sample_test( read_only image1d_t source, sampler_t sampler, __global %s4 *results )\n"
32 "{\n"
33 " int offset = get_global_id(0);\n"
34 " results[ offset ] = read_image%s( source, sampler, offset );\n"
35 "}\n";
36
37 static const char *kernelpattern_image_read_1d_buffer =
38 "__kernel void sample_test( read_only image1d_buffer_t source, sampler_t sampler, __global %s4 *results )\n"
39 "{\n"
40 " int offset = get_global_id(0);\n"
41 " results[ offset ] = read_image%s( source, offset );\n"
42 "}\n";
43
44 static const char *kernelpattern_image_read_1darray =
45 "__kernel void sample_test( read_only image1d_array_t source, sampler_t sampler, __global %s4 *results )\n"
46 "{\n"
47 " int tidX = get_global_id(0);\n"
48 " int tidY = get_global_id(1);\n"
49 " results[ tidY * get_image_width( source ) + tidX ] = read_image%s( source, sampler, (int2)( tidX, tidY ) );\n"
50 "}\n";
51
52 static const char *kernelpattern_image_read_2d =
53 "__kernel void sample_test( read_only image2d_t source, sampler_t sampler, __global %s4 *results )\n"
54 "{\n"
55 " int tidX = get_global_id(0);\n"
56 " int tidY = get_global_id(1);\n"
57 " results[ tidY * get_image_width( source ) + tidX ] = read_image%s( source, sampler, (int2)( tidX, tidY ) );\n"
58 "}\n";
59
60 static const char *kernelpattern_image_read_2darray =
61 "__kernel void sample_test( read_only image2d_array_t source, sampler_t sampler, __global %s4 *results )\n"
62 "{\n"
63 " int tidX = get_global_id(0);\n"
64 " int tidY = get_global_id(1);\n"
65 " int tidZ = get_global_id(2);\n"
66 " int width = get_image_width( source );\n"
67 " int height = get_image_height( source );\n"
68 " int offset = tidZ * width * height + tidY * width + tidX;\n"
69 "\n"
70 " results[ offset ] = read_image%s( source, sampler, (int4)( tidX, tidY, tidZ, 0 ) );\n"
71 "}\n";
72
73 static const char *kernelpattern_image_read_3d =
74 "__kernel void sample_test( read_only image3d_t source, sampler_t sampler, __global %s4 *results )\n"
75 "{\n"
76 " int tidX = get_global_id(0);\n"
77 " int tidY = get_global_id(1);\n"
78 " int tidZ = get_global_id(2);\n"
79 " int width = get_image_width( source );\n"
80 " int height = get_image_height( source );\n"
81 " int offset = tidZ * width * height + tidY * width + tidX;\n"
82 "\n"
83 " results[ offset ] = read_image%s( source, sampler, (int4)( tidX, tidY, tidZ, 0 ) );\n"
84 "}\n";
85
86 static const char *kernelpattern_image_read_2d_depth =
87 "__kernel void sample_test( read_only image2d_depth_t source, sampler_t sampler, __global %s *results )\n"
88 "{\n"
89 " int tidX = get_global_id(0);\n"
90 " int tidY = get_global_id(1);\n"
91 " results[ tidY * get_image_width( source ) + tidX ] = read_image%s( source, sampler, (int2)( tidX, tidY ) );\n"
92 "}\n";
93
94 static const char *kernelpattern_image_read_2darray_depth =
95 "__kernel void sample_test( read_only image2d_array_depth_t source, sampler_t sampler, __global %s *results )\n"
96 "{\n"
97 " int tidX = get_global_id(0);\n"
98 " int tidY = get_global_id(1);\n"
99 " int tidZ = get_global_id(2);\n"
100 " int width = get_image_width( source );\n"
101 " int height = get_image_height( source );\n"
102 " int offset = tidZ * width * height + tidY * width + tidX;\n"
103 "\n"
104 " results[ offset ] = read_image%s( source, sampler, (int4)( tidX, tidY, tidZ, 0 ) );\n"
105 "}\n";
106
107 static const char *kernelpattern_image_multisample_read_2d =
108 "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
109 "__kernel void sample_test( read_only image2d_msaa_t source, sampler_t sampler, __global %s4 *results )\n"
110 "{\n"
111 " int tidX = get_global_id(0);\n"
112 " int tidY = get_global_id(1);\n"
113 " int width = get_image_width( source );\n"
114 " int height = get_image_height( source );\n"
115 " int num_samples = get_image_num_samples( source );\n"
116 " for(size_t sample = 0; sample < num_samples; sample++ ) {\n"
117 " int offset = sample * width * height + tidY * width + tidX;\n"
118 " results[ offset ] = read_image%s( source, (int2)( tidX, tidY ), sample );\n"
119 " }\n"
120 "}\n";
121
122 static const char *kernelpattern_image_multisample_read_2d_depth =
123 "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
124 "__kernel void sample_test( read_only image2d_msaa_depth_t source, sampler_t sampler, __global %s *results )\n"
125 "{\n"
126 " int tidX = get_global_id(0);\n"
127 " int tidY = get_global_id(1);\n"
128 " int width = get_image_width( source );\n"
129 " int height = get_image_height( source );\n"
130 " int num_samples = get_image_num_samples( source );\n"
131 " for(size_t sample = 0; sample < num_samples; sample++ ) {\n"
132 " int offset = sample * width * height + tidY * width + tidX;\n"
133 " results[ offset ] = read_image%s( source, (int2)( tidX, tidY ), sample );\n"
134 " }\n"
135 "}\n";
136
137 static const char *kernelpattern_image_multisample_read_2darray =
138 "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
139 "__kernel void sample_test( read_only image2d_array_msaa_t source, sampler_t sampler, __global %s4 *results )\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 num_samples = get_image_num_samples( source );\n"
145 " int width = get_image_width( source );\n"
146 " int height = get_image_height( source );\n"
147 " int array_size = get_image_array_size( source );\n"
148 " for(size_t sample = 0; sample< num_samples; ++sample) {\n"
149 " int offset = (array_size * width * height) * sample + (width * height) * tidZ + tidY * width + tidX;\n"
150 " results[ offset ] = read_image%s( source, (int4)( tidX, tidY, tidZ, 1 ), sample );\n"
151 " }\n"
152 "}\n";
153
154 static const char *kernelpattern_image_multisample_read_2darray_depth =
155 "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
156 "__kernel void sample_test( read_only image2d_array_msaa_depth_t source, sampler_t sampler, __global %s *results )\n"
157 "{\n"
158 " int tidX = get_global_id(0);\n"
159 " int tidY = get_global_id(1);\n"
160 " int tidZ = get_global_id(2);\n"
161 " int num_samples = get_image_num_samples( source );\n"
162 " int width = get_image_width( source );\n"
163 " int height = get_image_height( source );\n"
164 " int array_size = get_image_array_size( source );\n"
165 " for(size_t sample = 0; sample < num_samples; ++sample) {\n"
166 " int offset = (array_size * width * height) * sample + (width * height) * tidZ + tidY * width + tidX;\n"
167 " results[ offset ] = read_image%s( source, (int4)( tidX, tidY, tidZ, 1 ), sample );\n"
168 " }\n"
169 "}\n";
170
get_appropriate_kernel_for_target(GLenum target,cl_channel_order channel_order)171 static const char* get_appropriate_kernel_for_target(GLenum target, cl_channel_order channel_order) {
172
173 switch (get_base_gl_target(target)) {
174 case GL_TEXTURE_1D:
175 return kernelpattern_image_read_1d;
176 case GL_TEXTURE_BUFFER:
177 return kernelpattern_image_read_1d_buffer;
178 case GL_TEXTURE_1D_ARRAY:
179 return kernelpattern_image_read_1darray;
180 case GL_TEXTURE_RECTANGLE_EXT:
181 case GL_TEXTURE_2D:
182 case GL_COLOR_ATTACHMENT0:
183 case GL_RENDERBUFFER:
184 case GL_TEXTURE_CUBE_MAP:
185 #ifdef GL_VERSION_3_2
186 if(channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
187 return kernelpattern_image_read_2d_depth;
188 #endif
189 return kernelpattern_image_read_2d;
190 case GL_TEXTURE_2D_ARRAY:
191 #ifdef GL_VERSION_3_2
192 if(channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
193 return kernelpattern_image_read_2darray_depth;
194 #endif
195 return kernelpattern_image_read_2darray;
196 case GL_TEXTURE_3D:
197 return kernelpattern_image_read_3d;
198 case GL_TEXTURE_2D_MULTISAMPLE:
199 #ifdef GL_VERSION_3_2
200 if(channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
201 return kernelpattern_image_multisample_read_2d_depth;
202 #endif
203 return kernelpattern_image_multisample_read_2d;
204 break;
205 case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
206 #ifdef GL_VERSION_3_2
207 if(channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
208 return kernelpattern_image_multisample_read_2darray_depth;
209 #endif
210 return kernelpattern_image_multisample_read_2darray;
211 break;
212 default:
213 log_error("Unsupported texture target (%s); cannot determine "
214 "appropriate kernel.", GetGLTargetName(target));
215 return NULL;
216 }
217 }
218
test_cl_image_read(cl_context context,cl_command_queue queue,GLenum gl_target,cl_mem image,size_t width,size_t height,size_t depth,size_t sampleNum,cl_image_format * outFormat,ExplicitType * outType,void ** outResultBuffer)219 int test_cl_image_read( cl_context context, cl_command_queue queue,
220 GLenum gl_target, cl_mem image, size_t width, size_t height, size_t depth, size_t sampleNum,
221 cl_image_format *outFormat, ExplicitType *outType, void **outResultBuffer )
222 {
223 clProgramWrapper program;
224 clKernelWrapper kernel;
225 clMemWrapper streams[ 2 ];
226
227 int error;
228 char kernelSource[2048];
229 char *programPtr;
230
231 // Use the image created from the GL texture.
232 streams[ 0 ] = image;
233
234 // Determine data type and format that CL came up with
235 error = clGetImageInfo( streams[ 0 ], CL_IMAGE_FORMAT, sizeof( cl_image_format ), outFormat, NULL );
236 test_error( error, "Unable to get CL image format" );
237
238 // Determine the number of samples
239 cl_uint samples = 0;
240 error = clGetImageInfo( streams[ 0 ], CL_IMAGE_NUM_SAMPLES, sizeof( samples ), &samples, NULL );
241 test_error( error, "Unable to get CL_IMAGE_NUM_SAMPLES" );
242
243 // Create the source
244 *outType = get_read_kernel_type( outFormat );
245 size_t channelSize = get_explicit_type_size( *outType );
246
247 const char* source = get_appropriate_kernel_for_target(gl_target, outFormat->image_channel_order);
248
249 sprintf( kernelSource, source, get_explicit_type_name( *outType ),
250 get_kernel_suffix( outFormat ) );
251
252 programPtr = kernelSource;
253 if( create_single_kernel_helper( context, &program, &kernel, 1,
254 (const char **)&programPtr, "sample_test", "" ) )
255 {
256 return -1;
257 }
258
259 // Create a vanilla output buffer
260 cl_device_id device;
261 error = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(device), &device, NULL);
262 test_error( error, "Unable to get queue device" );
263
264 cl_ulong maxAllocSize = 0;
265 error = clGetDeviceInfo( device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof( maxAllocSize ), &maxAllocSize, NULL );
266 test_error( error, "Unable to get CL_DEVICE_MAX_MEM_ALLOC_SIZE" );
267
268 size_t buffer_bytes = channelSize * get_channel_order_channel_count(outFormat->image_channel_order) * width * height * depth * sampleNum;
269 if (buffer_bytes > maxAllocSize) {
270 log_info("Output buffer size %d is too large for device (max alloc size %d) Skipping...\n",
271 (int)buffer_bytes, (int)maxAllocSize);
272 return 1;
273 }
274
275 streams[ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, buffer_bytes, NULL, &error );
276 test_error( error, "Unable to create output buffer" );
277
278 /* Assign streams and execute */
279 clSamplerWrapper sampler = clCreateSampler( context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
280 test_error( error, "Unable to create sampler" );
281
282 error = clSetKernelArg( kernel, 0, sizeof( streams[ 0 ] ), &streams[ 0 ] );
283 test_error( error, "Unable to set kernel arguments" );
284 error = clSetKernelArg( kernel, 1, sizeof( sampler ), &sampler );
285 test_error( error, "Unable to set kernel arguments" );
286 error = clSetKernelArg( kernel, 2, sizeof( streams[ 1 ] ), &streams[ 1 ] );
287 test_error( error, "Unable to set kernel arguments" );
288
289 glFinish();
290
291 error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &streams[ 0 ], 0, NULL, NULL);
292 test_error( error, "Unable to acquire GL obejcts");
293
294 // The ND range we use is a function of the dimensionality of the image.
295 size_t global_range[3] = { width, height, depth };
296 size_t *local_range = NULL;
297 int ndim = 1;
298
299 switch (get_base_gl_target(gl_target)) {
300 case GL_TEXTURE_1D:
301 case GL_TEXTURE_BUFFER:
302 ndim = 1;
303 break;
304 case GL_TEXTURE_RECTANGLE_EXT:
305 case GL_TEXTURE_2D:
306 case GL_TEXTURE_1D_ARRAY:
307 case GL_COLOR_ATTACHMENT0:
308 case GL_RENDERBUFFER:
309 case GL_TEXTURE_CUBE_MAP:
310 ndim = 2;
311 break;
312 case GL_TEXTURE_3D:
313 case GL_TEXTURE_2D_ARRAY:
314 #ifdef GL_VERSION_3_2
315 case GL_TEXTURE_2D_MULTISAMPLE:
316 case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
317 ndim = 3;
318 break;
319 #endif
320 default:
321 log_error("Test error: Unsupported texture target.\n");
322 return 1;
323 }
324
325 // 2D and 3D images have a special way to set the local size (legacy).
326 // Otherwise, we let CL select by leaving local_range as NULL.
327
328 if (gl_target == GL_TEXTURE_2D) {
329 local_range = (size_t*)malloc(sizeof(size_t) * ndim);
330 get_max_common_2D_work_group_size( context, kernel, global_range, local_range );
331
332 } else if (gl_target == GL_TEXTURE_3D) {
333 local_range = (size_t*)malloc(sizeof(size_t) * ndim);
334 get_max_common_3D_work_group_size( context, kernel, global_range, local_range );
335 }
336
337 error = clEnqueueNDRangeKernel( queue, kernel, ndim, NULL, global_range,
338 local_range, 0, NULL, NULL );
339 test_error( error, "Unable to execute test kernel" );
340
341 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &streams[ 0 ],
342 0, NULL, NULL );
343 test_error(error, "clEnqueueReleaseGLObjects failed");
344
345 // Read results from the CL buffer
346 *outResultBuffer = (void *)( new char[ channelSize * get_channel_order_channel_count(outFormat->image_channel_order) * width * height * depth * sampleNum] );
347 error = clEnqueueReadBuffer( queue, streams[ 1 ], CL_TRUE, 0,
348 channelSize * get_channel_order_channel_count(outFormat->image_channel_order) * width * height * depth * sampleNum, *outResultBuffer, 0, NULL, NULL );
349 test_error( error, "Unable to read output CL buffer!" );
350
351 // free the ranges
352 if (local_range) free(local_range);
353
354 return 0;
355 }
356
test_image_read(cl_context context,cl_command_queue queue,GLenum target,GLuint globj,size_t width,size_t height,size_t depth,size_t sampleNum,cl_image_format * outFormat,ExplicitType * outType,void ** outResultBuffer)357 static int test_image_read( cl_context context, cl_command_queue queue,
358 GLenum target, GLuint globj, size_t width, size_t height, size_t depth, size_t sampleNum,
359 cl_image_format *outFormat, ExplicitType *outType, void **outResultBuffer )
360 {
361 int error;
362
363 // Create a CL image from the supplied GL texture or renderbuffer.
364 cl_mem image;
365 if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0) {
366 image = (*clCreateFromGLRenderbuffer_ptr)( context, CL_MEM_READ_ONLY, globj, &error );
367 } else {
368 image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_READ_ONLY,
369 target, 0, globj, &error );
370 }
371
372 if( error != CL_SUCCESS ) {
373 if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0) {
374 print_error( error, "Unable to create CL image from GL renderbuffer" );
375 } else {
376 print_error( error, "Unable to create CL image from GL texture" );
377 GLint fmt;
378 glGetTexLevelParameteriv( target, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt );
379 log_error( " Supplied GL texture was base format %s and internal "
380 "format %s\n", GetGLBaseFormatName( fmt ), GetGLFormatName( fmt ) );
381 }
382 return error;
383 }
384
385 return test_cl_image_read( context, queue, target, image,
386 width, height, depth, sampleNum, outFormat, outType, outResultBuffer );
387 }
388
test_image_format_read(cl_context context,cl_command_queue queue,size_t width,size_t height,size_t depth,GLenum target,const format * fmt,MTdata data)389 static int test_image_format_read(cl_context context, cl_command_queue queue,
390 size_t width, size_t height, size_t depth,
391 GLenum target, const format *fmt, MTdata data)
392 {
393 int error = 0;
394
395 // Determine the maximum number of supported samples
396 GLint samples = 1;
397 if (target == GL_TEXTURE_2D_MULTISAMPLE || target == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
398 samples = get_gl_max_samples(target, fmt->internal);
399
400 // If we're testing a half float format, then we need to determine the
401 // rounding mode of this machine. Punt if we fail to do so.
402
403 if( fmt->type == kHalf )
404 {
405 if( DetectFloatToHalfRoundingMode(queue) )
406 return 1;
407 bool supports_half = false;
408 error = supportsHalf(context, &supports_half);
409 if( error != 0 )
410 return error;
411 if (!supports_half) return 0;
412 }
413 #ifdef GL_VERSION_3_2
414 if (get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE ||
415 get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
416 {
417 bool supports_msaa;
418 error = supportsMsaa(context, &supports_msaa);
419 if( error != 0 ) return error;
420 if (!supports_msaa) return 0;
421 }
422 if (fmt->formattype == GL_DEPTH_COMPONENT ||
423 fmt->formattype == GL_DEPTH_STENCIL)
424 {
425 bool supports_depth;
426 error = supportsDepth(context, &supports_depth);
427 if( error != 0 ) return error;
428 if (!supports_depth) return 0;
429 }
430 #endif
431 size_t w = width, h = height, d = depth;
432
433 // Unpack the format and use it, along with the target, to create an
434 // appropriate GL texture.
435
436 GLenum gl_fmt = fmt->formattype;
437 GLenum gl_internal_fmt = fmt->internal;
438 GLenum gl_type = fmt->datatype;
439 ExplicitType type = fmt->type;
440
441 // Required for most of the texture-backed cases:
442 glTextureWrapper texture;
443
444 // Required for the special case of TextureBuffer textures:
445 glBufferWrapper glbuf;
446
447 // And these are required for the case of Renderbuffer images:
448 glFramebufferWrapper glFramebuffer;
449 glRenderbufferWrapper glRenderbuffer;
450
451 void* buffer = NULL;
452
453 // Use the correct texture creation function depending on the target, and
454 // adjust width, height, depth as appropriate so subsequent size calculations
455 // succeed.
456
457 switch (get_base_gl_target(target)) {
458 case GL_TEXTURE_1D:
459 h = 1; d = 1;
460 buffer = CreateGLTexture1D( width, target, gl_fmt,
461 gl_internal_fmt, gl_type, type, &texture, &error, true, data );
462 break;
463 case GL_TEXTURE_BUFFER:
464 h = 1; d = 1;
465 buffer = CreateGLTextureBuffer(width, target, gl_fmt, gl_internal_fmt,
466 gl_type, type, &texture, &glbuf, &error, true, data);
467 break;
468 case GL_RENDERBUFFER:
469 case GL_COLOR_ATTACHMENT0:
470 d = 1;
471 buffer = CreateGLRenderbuffer(width, height, target, gl_fmt,
472 gl_internal_fmt, gl_type, type, &glFramebuffer, &glRenderbuffer, &error,
473 data, true);
474 break;
475 case GL_TEXTURE_2D:
476 case GL_TEXTURE_RECTANGLE_EXT:
477 case GL_TEXTURE_CUBE_MAP:
478 d = 1;
479 buffer = CreateGLTexture2D(width, height, target, gl_fmt, gl_internal_fmt,
480 gl_type, type, &texture, &error, true, data);
481 break;
482 case GL_TEXTURE_1D_ARRAY:
483 d = 1;
484 buffer = CreateGLTexture1DArray( width, height, target, gl_fmt,
485 gl_internal_fmt, gl_type, type, &texture, &error, true, data );
486 break;
487 case GL_TEXTURE_2D_ARRAY:
488 buffer = CreateGLTexture2DArray( width, height, depth, target, gl_fmt,
489 gl_internal_fmt, gl_type, type, &texture, &error, true, data );
490 break;
491 case GL_TEXTURE_3D:
492 buffer = CreateGLTexture3D( width, height, depth, target, gl_fmt,
493 gl_internal_fmt, gl_type, type, &texture, &error, data, true );
494 break;
495 #ifdef GL_VERSION_3_2
496 case GL_TEXTURE_2D_MULTISAMPLE:
497 d = 1;
498 buffer = CreateGLTexture2DMultisample( width, height, samples, target, gl_fmt,
499 gl_internal_fmt, gl_type, type, &texture, &error, true, data, true );
500 break;
501 case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
502 buffer = CreateGLTexture2DArrayMultisample( width, height, depth, samples, target, gl_fmt,
503 gl_internal_fmt, gl_type, type, &texture, &error, true, data, true );
504 break;
505 #endif
506 default:
507 log_error("Unsupported texture target.");
508 return 1;
509 }
510
511 if ( error == -2 ) {
512 log_info("OpenGL texture couldn't be created, because a texture is too big. Skipping test.\n");
513 return 0;
514 }
515
516 // Check to see if the texture could not be created for some other reason like
517 // GL_FRAMEBUFFER_UNSUPPORTED
518 if (error == GL_FRAMEBUFFER_UNSUPPORTED) {
519 log_info("Skipping...\n");
520 return 0;
521 }
522
523 if ( error != 0 ) {
524 if ((gl_fmt == GL_RGBA_INTEGER_EXT) && (!CheckGLIntegerExtensionSupport())){
525 log_info("OpenGL version does not support GL_RGBA_INTEGER_EXT. "
526 "Skipping test.\n");
527 return 0;
528 } else {
529 return error;
530 }
531 }
532
533 BufferOwningPtr<char> inputBuffer(buffer);
534 if( inputBuffer == NULL )
535 return -1;
536
537 cl_image_format clFormat;
538 ExplicitType actualType;
539 char *outBuffer;
540
541 // Perform the read:
542
543 GLuint globj = texture;
544 if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0) {
545 globj = glRenderbuffer;
546 }
547
548 error = test_image_read( context, queue, target, globj, w, h, d, samples, &clFormat,
549 &actualType, (void **)&outBuffer );
550
551 if( error != 0 )
552 return error;
553
554 BufferOwningPtr<char> actualResults(outBuffer);
555 if( actualResults == NULL )
556 return -1;
557
558 log_info( "- Read [%4d x %4d x %4d x %4d] : GL Texture : %s : %s : %s => CL Image : %s : %s \n",
559 (int)w, (int)h, (int)d, (int)samples, GetGLFormatName( gl_fmt ), GetGLFormatName( gl_internal_fmt ),
560 GetGLTypeName( gl_type ), GetChannelOrderName( clFormat.image_channel_order ),
561 GetChannelTypeName( clFormat.image_channel_data_type ));
562
563 BufferOwningPtr<char> convertedInputs;
564
565 // We have to convert our input buffer to the returned type, so we can validate.
566 // This is necessary because OpenCL might not actually pick an internal format
567 // that actually matches our input format (for example, if it picks a normalized
568 // format, the results will come out as floats instead of going in as ints).
569
570 if ( gl_type == GL_UNSIGNED_INT_2_10_10_10_REV )
571 {
572 cl_uint *p = (cl_uint *)buffer;
573 float *inData = (float *)malloc( w * h * d * samples * sizeof(float) );
574
575 for( size_t i = 0; i < 4 * w * h * d * samples; i += 4 )
576 {
577 inData[ i + 0 ] = (float)( ( p[ 0 ] >> 20 ) & 0x3ff ) / (float)1023;
578 inData[ i + 1 ] = (float)( ( p[ 0 ] >> 10 ) & 0x3ff ) / (float)1023;
579 inData[ i + 2 ] = (float)( p[ 0 ] & 0x3ff ) / (float)1023;
580 p++;
581 }
582
583 convertedInputs.reset( inData );
584 if( convertedInputs == NULL )
585 return -1;
586 }
587 else if ( gl_type == GL_DEPTH24_STENCIL8 )
588 {
589 // GL_DEPTH24_STENCIL8 is treated as CL_UNORM_INT24 + CL_DEPTH_STENCIL where
590 // the stencil is ignored.
591 cl_uint *p = (cl_uint *)buffer;
592 float *inData = (float *)malloc( w * h * d * samples * sizeof(float) );
593
594 for( size_t i = 0; i < w * h * d * samples; i++ )
595 {
596 inData[ i ] = (float)((p[i] >> 8) & 0xffffff) / (float)0xfffffe;
597 }
598
599 convertedInputs.reset( inData );
600 if( convertedInputs == NULL )
601 return -1;
602 }
603 else if ( gl_type == GL_FLOAT_32_UNSIGNED_INT_24_8_REV)
604 {
605 // GL_FLOAT_32_UNSIGNED_INT_24_8_REV is treated as a CL_FLOAT +
606 // unused 24 + CL_DEPTH_STENCIL; we check the float value and ignore the
607 // second word
608
609 float *p = (float *)buffer;
610 float *inData = (float *)malloc( w * h * d * samples * sizeof(float) );
611
612 for( size_t i = 0; i < w * h * d * samples; i++ )
613 {
614 inData[ i ] = p[i*2];
615 }
616
617 convertedInputs.reset( inData );
618 if( convertedInputs == NULL )
619 return -1;
620 }
621 else
622 {
623 convertedInputs.reset(convert_to_expected( inputBuffer,
624 w * h * d * samples, type, actualType, get_channel_order_channel_count(clFormat.image_channel_order) ));
625 if( convertedInputs == NULL )
626 return -1;
627 }
628
629 // Now we validate
630 if( actualType == kFloat )
631 {
632 if ( clFormat.image_channel_data_type == CL_UNORM_INT_101010 )
633 {
634 return validate_float_results_rgb_101010( convertedInputs, actualResults, w, h, d, samples );
635 }
636 else
637 {
638 return validate_float_results( convertedInputs, actualResults, w, h, d, samples, get_channel_order_channel_count(clFormat.image_channel_order) );
639 }
640 }
641 else
642 {
643 return validate_integer_results( convertedInputs, actualResults, w, h, d, samples, get_explicit_type_size( actualType ) );
644 }
645 }
646
test_images_read_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)647 int test_images_read_common(cl_device_id device, cl_context context,
648 cl_command_queue queue, const format *formats,
649 size_t nformats, GLenum *targets, size_t ntargets,
650 sizevec_t *sizes, size_t nsizes)
651 {
652 int error = 0;
653 RandomSeed seed(gRandomSeed);
654
655 // First, ensure this device supports images.
656
657 if (checkForImageSupport(device)) {
658 log_info("Device does not support images. Skipping test.\n");
659 return 0;
660 }
661
662 size_t fidx, tidx, sidx;
663
664 // Test each format on every target, every size.
665
666 for ( fidx = 0; fidx < nformats; fidx++ ) {
667 for ( tidx = 0; tidx < ntargets; tidx++ ) {
668
669 // Texture buffer only takes an internal format, so the level data passed
670 // by the test and used for verification must match the internal format
671 if ((targets[tidx] == GL_TEXTURE_BUFFER) && (GetGLFormat(formats[ fidx ].internal) != formats[fidx].formattype))
672 continue;
673
674 if ( formats[ fidx ].datatype == GL_UNSIGNED_INT_2_10_10_10_REV )
675 {
676 // Check if the RGB 101010 format is supported
677 if ( is_rgb_101010_supported( context, targets[ tidx ] ) == 0 )
678 break; // skip
679 }
680
681 if (targets[tidx] != GL_TEXTURE_BUFFER)
682 log_info( "Testing image read for GL format %s : %s : %s : %s\n",
683 GetGLTargetName( targets[ tidx ] ),
684 GetGLFormatName( formats[ fidx ].internal ),
685 GetGLBaseFormatName( formats[ fidx ].formattype ),
686 GetGLTypeName( formats[ fidx ].datatype ) );
687 else
688 log_info( "Testing image read for GL format %s : %s\n",
689 GetGLTargetName( targets[ tidx ] ),
690 GetGLFormatName( formats[ fidx ].internal ));
691
692 for ( sidx = 0; sidx < nsizes; sidx++ ) {
693
694 // Test this format + size:
695 int err;
696 if ((err = test_image_format_read(context, queue,
697 sizes[sidx].width, sizes[sidx].height, sizes[sidx].depth,
698 targets[tidx], &formats[fidx], seed) ))
699 {
700 // Negative return values are errors, positive mean the test was skipped
701 if (err < 0) {
702
703 // We land here in the event of test failure.
704
705 log_error( "ERROR: Image read test failed for %s : %s : %s : %s\n\n",
706 GetGLTargetName( targets[ tidx ] ),
707 GetGLFormatName( formats[ fidx ].internal ),
708 GetGLBaseFormatName( formats[ fidx ].formattype ),
709 GetGLTypeName( formats[ fidx ].datatype ) );
710 error++;
711 }
712
713 // Skip the other sizes for this format.
714 printf("Skipping remaining sizes for this format\n");
715
716 break;
717 }
718 }
719
720 // Note a successful format test, if we passed every size.
721
722 if( sidx == sizeof (sizes) / sizeof( sizes[0] ) ) {
723 log_info( "passed: Image read test for GL format %s : %s : %s : %s\n\n",
724 GetGLTargetName( targets[ tidx ] ),
725 GetGLFormatName( formats[ fidx ].internal ),
726 GetGLBaseFormatName( formats[ fidx ].formattype ),
727 GetGLTypeName( formats[ fidx ].datatype ) );
728 }
729 }
730 }
731
732 return error;
733 }
734