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