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,struct format * fmt,MTdata data)389 static int test_image_format_read(
390 cl_context context, cl_command_queue queue,
391 size_t width, size_t height, size_t depth,
392 GLenum target, struct format* fmt, MTdata data)
393 {
394 int error = 0;
395
396 // Determine the maximum number of supported samples
397 GLint samples = 1;
398 if (target == GL_TEXTURE_2D_MULTISAMPLE || target == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
399 samples = get_gl_max_samples(target, fmt->internal);
400
401 // If we're testing a half float format, then we need to determine the
402 // rounding mode of this machine. Punt if we fail to do so.
403
404 if( fmt->type == kHalf )
405 {
406 if( DetectFloatToHalfRoundingMode(queue) )
407 return 1;
408 bool supports_half = false;
409 error = supportsHalf(context, &supports_half);
410 if( error != 0 )
411 return error;
412 if (!supports_half) return 0;
413 }
414 #ifdef GL_VERSION_3_2
415 if (get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE ||
416 get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
417 {
418 bool supports_msaa;
419 error = supportsMsaa(context, &supports_msaa);
420 if( error != 0 ) return error;
421 if (!supports_msaa) return 0;
422 }
423 if (fmt->formattype == GL_DEPTH_COMPONENT ||
424 fmt->formattype == GL_DEPTH_STENCIL)
425 {
426 bool supports_depth;
427 error = supportsDepth(context, &supports_depth);
428 if( error != 0 ) return error;
429 if (!supports_depth) return 0;
430 }
431 #endif
432 size_t w = width, h = height, d = depth;
433
434 // Unpack the format and use it, along with the target, to create an
435 // appropriate GL texture.
436
437 GLenum gl_fmt = fmt->formattype;
438 GLenum gl_internal_fmt = fmt->internal;
439 GLenum gl_type = fmt->datatype;
440 ExplicitType type = fmt->type;
441
442 // Required for most of the texture-backed cases:
443 glTextureWrapper texture;
444
445 // Required for the special case of TextureBuffer textures:
446 glBufferWrapper glbuf;
447
448 // And these are required for the case of Renderbuffer images:
449 glFramebufferWrapper glFramebuffer;
450 glRenderbufferWrapper glRenderbuffer;
451
452 void* buffer = NULL;
453
454 // Use the correct texture creation function depending on the target, and
455 // adjust width, height, depth as appropriate so subsequent size calculations
456 // succeed.
457
458 switch (get_base_gl_target(target)) {
459 case GL_TEXTURE_1D:
460 h = 1; d = 1;
461 buffer = CreateGLTexture1D( width, target, gl_fmt,
462 gl_internal_fmt, gl_type, type, &texture, &error, true, data );
463 break;
464 case GL_TEXTURE_BUFFER:
465 h = 1; d = 1;
466 buffer = CreateGLTextureBuffer(width, target, gl_fmt, gl_internal_fmt,
467 gl_type, type, &texture, &glbuf, &error, true, data);
468 break;
469 case GL_RENDERBUFFER:
470 case GL_COLOR_ATTACHMENT0:
471 d = 1;
472 buffer = CreateGLRenderbuffer(width, height, target, gl_fmt,
473 gl_internal_fmt, gl_type, type, &glFramebuffer, &glRenderbuffer, &error,
474 data, true);
475 break;
476 case GL_TEXTURE_2D:
477 case GL_TEXTURE_RECTANGLE_EXT:
478 case GL_TEXTURE_CUBE_MAP:
479 d = 1;
480 buffer = CreateGLTexture2D(width, height, target, gl_fmt, gl_internal_fmt,
481 gl_type, type, &texture, &error, true, data);
482 break;
483 case GL_TEXTURE_1D_ARRAY:
484 d = 1;
485 buffer = CreateGLTexture1DArray( width, height, target, gl_fmt,
486 gl_internal_fmt, gl_type, type, &texture, &error, true, data );
487 break;
488 case GL_TEXTURE_2D_ARRAY:
489 buffer = CreateGLTexture2DArray( width, height, depth, target, gl_fmt,
490 gl_internal_fmt, gl_type, type, &texture, &error, true, data );
491 break;
492 case GL_TEXTURE_3D:
493 buffer = CreateGLTexture3D( width, height, depth, target, gl_fmt,
494 gl_internal_fmt, gl_type, type, &texture, &error, data, true );
495 break;
496 #ifdef GL_VERSION_3_2
497 case GL_TEXTURE_2D_MULTISAMPLE:
498 d = 1;
499 buffer = CreateGLTexture2DMultisample( width, height, samples, target, gl_fmt,
500 gl_internal_fmt, gl_type, type, &texture, &error, true, data, true );
501 break;
502 case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
503 buffer = CreateGLTexture2DArrayMultisample( width, height, depth, samples, target, gl_fmt,
504 gl_internal_fmt, gl_type, type, &texture, &error, true, data, true );
505 break;
506 #endif
507 default:
508 log_error("Unsupported texture target.");
509 return 1;
510 }
511
512 if ( error == -2 ) {
513 log_info("OpenGL texture couldn't be created, because a texture is too big. Skipping test.\n");
514 return 0;
515 }
516
517 // Check to see if the texture could not be created for some other reason like
518 // GL_FRAMEBUFFER_UNSUPPORTED
519 if (error == GL_FRAMEBUFFER_UNSUPPORTED) {
520 log_info("Skipping...\n");
521 return 0;
522 }
523
524 if ( error != 0 ) {
525 if ((gl_fmt == GL_RGBA_INTEGER_EXT) && (!CheckGLIntegerExtensionSupport())){
526 log_info("OpenGL version does not support GL_RGBA_INTEGER_EXT. "
527 "Skipping test.\n");
528 return 0;
529 } else {
530 return error;
531 }
532 }
533
534 BufferOwningPtr<char> inputBuffer(buffer);
535 if( inputBuffer == NULL )
536 return -1;
537
538 cl_image_format clFormat;
539 ExplicitType actualType;
540 char *outBuffer;
541
542 // Perform the read:
543
544 GLuint globj = texture;
545 if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0) {
546 globj = glRenderbuffer;
547 }
548
549 error = test_image_read( context, queue, target, globj, w, h, d, samples, &clFormat,
550 &actualType, (void **)&outBuffer );
551
552 if( error != 0 )
553 return error;
554
555 BufferOwningPtr<char> actualResults(outBuffer);
556 if( actualResults == NULL )
557 return -1;
558
559 log_info( "- Read [%4d x %4d x %4d x %4d] : GL Texture : %s : %s : %s => CL Image : %s : %s \n",
560 (int)w, (int)h, (int)d, (int)samples, GetGLFormatName( gl_fmt ), GetGLFormatName( gl_internal_fmt ),
561 GetGLTypeName( gl_type ), GetChannelOrderName( clFormat.image_channel_order ),
562 GetChannelTypeName( clFormat.image_channel_data_type ));
563
564 BufferOwningPtr<char> convertedInputs;
565
566 // We have to convert our input buffer to the returned type, so we can validate.
567 // This is necessary because OpenCL might not actually pick an internal format
568 // that actually matches our input format (for example, if it picks a normalized
569 // format, the results will come out as floats instead of going in as ints).
570
571 if ( gl_type == GL_UNSIGNED_INT_2_10_10_10_REV )
572 {
573 cl_uint *p = (cl_uint *)buffer;
574 float *inData = (float *)malloc( w * h * d * samples * sizeof(float) );
575
576 for( size_t i = 0; i < 4 * w * h * d * samples; i += 4 )
577 {
578 inData[ i + 0 ] = (float)( ( p[ 0 ] >> 20 ) & 0x3ff ) / (float)1023;
579 inData[ i + 1 ] = (float)( ( p[ 0 ] >> 10 ) & 0x3ff ) / (float)1023;
580 inData[ i + 2 ] = (float)( p[ 0 ] & 0x3ff ) / (float)1023;
581 p++;
582 }
583
584 convertedInputs.reset( inData );
585 if( convertedInputs == NULL )
586 return -1;
587 }
588 else if ( gl_type == GL_DEPTH24_STENCIL8 )
589 {
590 // GL_DEPTH24_STENCIL8 is treated as CL_UNORM_INT24 + CL_DEPTH_STENCIL where
591 // the stencil is ignored.
592 cl_uint *p = (cl_uint *)buffer;
593 float *inData = (float *)malloc( w * h * d * samples * sizeof(float) );
594
595 for( size_t i = 0; i < w * h * d * samples; i++ )
596 {
597 inData[ i ] = (float)((p[i] >> 8) & 0xffffff) / (float)0xfffffe;
598 }
599
600 convertedInputs.reset( inData );
601 if( convertedInputs == NULL )
602 return -1;
603 }
604 else if ( gl_type == GL_FLOAT_32_UNSIGNED_INT_24_8_REV)
605 {
606 // GL_FLOAT_32_UNSIGNED_INT_24_8_REV is treated as a CL_FLOAT +
607 // unused 24 + CL_DEPTH_STENCIL; we check the float value and ignore the
608 // second word
609
610 float *p = (float *)buffer;
611 float *inData = (float *)malloc( w * h * d * samples * sizeof(float) );
612
613 for( size_t i = 0; i < w * h * d * samples; i++ )
614 {
615 inData[ i ] = p[i*2];
616 }
617
618 convertedInputs.reset( inData );
619 if( convertedInputs == NULL )
620 return -1;
621 }
622 else
623 {
624 convertedInputs.reset(convert_to_expected( inputBuffer,
625 w * h * d * samples, type, actualType, get_channel_order_channel_count(clFormat.image_channel_order) ));
626 if( convertedInputs == NULL )
627 return -1;
628 }
629
630 // Now we validate
631 if( actualType == kFloat )
632 {
633 if ( clFormat.image_channel_data_type == CL_UNORM_INT_101010 )
634 {
635 return validate_float_results_rgb_101010( convertedInputs, actualResults, w, h, d, samples );
636 }
637 else
638 {
639 return validate_float_results( convertedInputs, actualResults, w, h, d, samples, get_channel_order_channel_count(clFormat.image_channel_order) );
640 }
641 }
642 else
643 {
644 return validate_integer_results( convertedInputs, actualResults, w, h, d, samples, get_explicit_type_size( actualType ) );
645 }
646 }
647
test_images_read_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)648 int test_images_read_common( cl_device_id device, cl_context context,
649 cl_command_queue queue, struct format* formats, size_t nformats,
650 GLenum *targets, size_t ntargets, 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