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
18 #include "gl_headers.h"
19
20 static const char *imageReadKernelPattern =
21 "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" /* added support for half floats */
22 "__kernel void sample_test( read_only image2d_t source, sampler_t sampler, __global %s4 *results )\n"
23 "{\n"
24 " int tidX = get_global_id(0);\n"
25 " int tidY = get_global_id(1);\n"
26 " results[ tidY * get_image_width( source ) + tidX ] = read_image%s( source, sampler, (int2)( tidX, tidY ) );\n"
27 "}\n";
28
29 static const char *imageWriteKernelPattern =
30 "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" /* added support for half floats */
31 "__kernel void sample_test( __global %s4 *source, write_only image2d_t dest )\n"
32 "{\n"
33 " int tidX = get_global_id(0);\n"
34 " int tidY = get_global_id(1);\n"
35 " uint index = tidY * get_image_width( dest ) + tidX;\n"
36 " %s4 value = source[index];\n"
37 " write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
38 "}\n";
39
test_cl_image_read(cl_context context,cl_command_queue queue,cl_mem clImage,size_t imageWidth,size_t imageHeight,cl_image_format * outFormat,ExplicitType * outType,void ** outResultBuffer)40 int test_cl_image_read( cl_context context, cl_command_queue queue, cl_mem clImage,
41 size_t imageWidth, size_t imageHeight, cl_image_format *outFormat, ExplicitType *outType, void **outResultBuffer )
42 {
43 clProgramWrapper program;
44 clKernelWrapper kernel;
45 clMemWrapper outStream;
46
47 int error;
48 size_t threads[ 2 ], localThreads[ 2 ];
49 char kernelSource[10240];
50 char *programPtr;
51
52
53 // Determine data type and format that CL came up with
54 error = clGetImageInfo( clImage, CL_IMAGE_FORMAT, sizeof( cl_image_format ), outFormat, NULL );
55 test_error( error, "Unable to get CL image format" );
56
57 /* Create the source */
58 *outType = get_read_kernel_type( outFormat );
59 size_t channelSize = get_explicit_type_size( *outType );
60
61 sprintf( kernelSource, imageReadKernelPattern, get_explicit_type_name( *outType ), get_kernel_suffix( outFormat ) );
62
63 #ifdef GLES_DEBUG
64 log_info("-- start cl image read kernel --\n");
65 log_info("%s", kernelSource);
66 log_info("-- end cl image read kernel --\n");
67 #endif
68
69 /* Create kernel */
70 programPtr = kernelSource;
71 if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
72 {
73 return -1;
74 }
75
76
77 // Create a vanilla output buffer
78 outStream = clCreateBuffer( context, CL_MEM_READ_WRITE, channelSize * 4 * imageWidth * imageHeight, NULL, &error );
79 test_error( error, "Unable to create output buffer" );
80
81
82 /* Assign streams and execute */
83 clSamplerWrapper sampler = clCreateSampler( context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
84 test_error( error, "Unable to create sampler" );
85
86 error = clSetKernelArg( kernel, 0, sizeof( clImage ), &clImage );
87 test_error( error, "Unable to set kernel arguments" );
88 error = clSetKernelArg( kernel, 1, sizeof( sampler ), &sampler );
89 test_error( error, "Unable to set kernel arguments" );
90 error = clSetKernelArg( kernel, 2, sizeof( outStream ), &outStream );
91 test_error( error, "Unable to set kernel arguments" );
92
93 glFlush();
94
95 error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &clImage, 0, NULL, NULL);
96 test_error( error, "Unable to acquire GL obejcts");
97
98 /* Run the kernel */
99 threads[ 0 ] = imageWidth;
100 threads[ 1 ] = imageHeight;
101
102 error = get_max_common_2D_work_group_size( context, kernel, threads, localThreads );
103 test_error( error, "Unable to get work group size to use" );
104
105 error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, localThreads, 0, NULL, NULL );
106 test_error( error, "Unable to execute test kernel" );
107
108
109 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &clImage, 0, NULL, NULL );
110 test_error(error, "clEnqueueReleaseGLObjects failed");
111
112 // Read results from the CL buffer
113 *outResultBuffer = malloc(channelSize * 4 * imageWidth * imageHeight);
114 error = clEnqueueReadBuffer( queue, outStream, CL_TRUE, 0, channelSize * 4 * imageWidth * imageHeight,
115 *outResultBuffer, 0, NULL, NULL );
116 test_error( error, "Unable to read output CL buffer!" );
117
118 return 0;
119 }
120
test_image_read(cl_context context,cl_command_queue queue,GLenum glTarget,GLuint glTexture,size_t imageWidth,size_t imageHeight,cl_image_format * outFormat,ExplicitType * outType,void ** outResultBuffer)121 static int test_image_read( cl_context context, cl_command_queue queue, GLenum glTarget, GLuint glTexture,
122 size_t imageWidth, size_t imageHeight, cl_image_format *outFormat, ExplicitType *outType, void **outResultBuffer )
123 {
124 // Create a CL image from the supplied GL texture
125 int error;
126 clMemWrapper image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_READ_ONLY, glTarget, 0, glTexture, &error );
127 if( error != CL_SUCCESS )
128 {
129 print_error( error, "Unable to create CL image from GL texture" );
130 #ifndef GL_ES_VERSION_2_0
131 GLint fmt;
132 glGetTexLevelParameteriv( glTarget, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt );
133 log_error( " Supplied GL texture was baseformat %s and internalformat %s\n", GetGLBaseFormatName( fmt ), GetGLFormatName( fmt ) );
134 #endif
135 return error;
136 }
137
138 return test_cl_image_read( context, queue, image, imageWidth, imageHeight, outFormat, outType, outResultBuffer );
139 }
140
test_image_format_read(cl_context context,cl_command_queue queue,size_t width,size_t height,GLenum target,GLenum format,GLenum internalFormat,GLenum glType,ExplicitType type,MTdata d)141 int test_image_format_read( cl_context context, cl_command_queue queue,
142 size_t width, size_t height, GLenum target,
143 GLenum format, GLenum internalFormat,
144 GLenum glType, ExplicitType type, MTdata d )
145 {
146 int error;
147
148
149 // Create the GL texture
150 glTextureWrapper glTexture;
151 void *tmp = CreateGLTexture2D( width, height, target, format, internalFormat, glType, type, &glTexture, &error, true, d );
152 BufferOwningPtr<char> inputBuffer(tmp);
153 if( error != 0 )
154 {
155 return error;
156 }
157
158 /* skip formats not supported by OpenGL */
159 if(!tmp)
160 {
161 return 0;
162 }
163
164 // Run and get the results
165 cl_image_format clFormat;
166 ExplicitType actualType;
167 char *outBuffer;
168 error = test_image_read( context, queue, target, glTexture, width, height, &clFormat, &actualType, (void **)&outBuffer );
169 if( error != 0 )
170 return error;
171 BufferOwningPtr<char> actualResults(outBuffer);
172
173 log_info( "- Read [%4d x %4d] : GL Texture : %s : %s : %s => CL Image : %s : %s \n", (int)width, (int)height,
174 GetGLFormatName( format ), GetGLFormatName( internalFormat ), GetGLTypeName( glType),
175 GetChannelOrderName( clFormat.image_channel_order ), GetChannelTypeName( clFormat.image_channel_data_type ));
176
177 // We have to convert our input buffer to the returned type, so we can validate.
178 BufferOwningPtr<char> convertedInputs(convert_to_expected( inputBuffer, width * height, type, actualType ));
179
180 // Now we validate
181 int valid = 0;
182 if(convertedInputs) {
183 if( actualType == kFloat )
184 valid = validate_float_results( convertedInputs, actualResults, width, height );
185 else
186 valid = validate_integer_results( convertedInputs, actualResults, width, height, get_explicit_type_size( actualType ) );
187 }
188
189 return valid;
190 }
191
test_images_read(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)192 int test_images_read( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
193 {
194 GLenum targets[] =
195 #ifdef GL_ES_VERSION_2_0
196 { GL_TEXTURE_2D };
197 #else // GL_ES_VERSION_2_0
198 { GL_TEXTURE_2D, GL_TEXTURE_RECTANGLE_EXT };
199 #endif // GL_ES_VERSION_2_0
200
201 struct {
202 GLenum internal;
203 GLenum format;
204 GLenum datatype;
205 ExplicitType type;
206
207 } formats[] = {
208 { GL_RGBA, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
209 { GL_RGBA, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
210 { GL_RGBA, GL_RGBA, GL_FLOAT, kFloat },
211 };
212
213 size_t fmtIdx, tgtIdx;
214 int error = 0;
215 size_t iter = 6;
216 RandomSeed seed(gRandomSeed );
217
218 // Check if images are supported
219 if (checkForImageSupport(device)) {
220 log_info("Device does not support images. Skipping test.\n");
221 return 0;
222 }
223
224 // Loop through a set of GL formats, testing a set of sizes against each one
225 for( fmtIdx = 0; fmtIdx < sizeof( formats ) / sizeof( formats[ 0 ] ); fmtIdx++ )
226 {
227 for( tgtIdx = 0; tgtIdx < sizeof( targets ) / sizeof( targets[ 0 ] ); tgtIdx++ )
228 {
229 size_t i;
230
231 log_info( "Testing image read for GL format %s : %s : %s : %s\n",
232 GetGLTargetName( targets[ tgtIdx ] ),
233 GetGLFormatName( formats[ fmtIdx ].internal ),
234 GetGLBaseFormatName( formats[ fmtIdx ].format ),
235 GetGLTypeName( formats[ fmtIdx ].datatype ) );
236
237 for( i = 0; i < iter; i++ )
238 {
239 size_t width = random_in_range( 16, 512, seed );
240 size_t height = random_in_range( 16, 512, seed );
241
242 if( test_image_format_read( context, queue, width, height,
243 targets[ tgtIdx ],
244 formats[ fmtIdx ].format,
245 formats[ fmtIdx ].internal,
246 formats[ fmtIdx ].datatype,
247 formats[ fmtIdx ].type, seed ) )
248 {
249 log_error( "ERROR: Image read test failed for %s : %s : %s : %s\n\n",
250 GetGLTargetName( targets[ tgtIdx ] ),
251 GetGLFormatName( formats[ fmtIdx ].internal ),
252 GetGLBaseFormatName( formats[ fmtIdx ].format ),
253 GetGLTypeName( formats[ fmtIdx ].datatype ) );
254
255 error++;
256 break; // Skip other sizes for this combination
257 }
258 }
259 if( i == iter )
260 {
261 log_info( "passed: Image read for GL format %s : %s : %s : %s\n\n",
262 GetGLTargetName( targets[ tgtIdx ] ),
263 GetGLFormatName( formats[ fmtIdx ].internal ),
264 GetGLBaseFormatName( formats[ fmtIdx ].format ),
265 GetGLTypeName( formats[ fmtIdx ].datatype ) );
266 }
267 }
268 }
269
270 return error;
271 }
272
test_images_read_cube(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)273 int test_images_read_cube( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
274 {
275 GLenum targets[] = {
276 GL_TEXTURE_CUBE_MAP_POSITIVE_X,
277 GL_TEXTURE_CUBE_MAP_POSITIVE_Y,
278 GL_TEXTURE_CUBE_MAP_POSITIVE_Z,
279 GL_TEXTURE_CUBE_MAP_NEGATIVE_X,
280 GL_TEXTURE_CUBE_MAP_NEGATIVE_Y,
281 GL_TEXTURE_CUBE_MAP_NEGATIVE_Z };
282
283 struct {
284 GLenum internal;
285 GLenum format;
286 GLenum datatype;
287 ExplicitType type;
288
289 } formats[] = {
290 #ifdef GL_ES_VERSION_2_0
291 { GL_RGBA, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
292 { GL_RGBA, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
293 // XXX add others
294 #else // GL_ES_VERSION_2_0
295 { GL_RGBA, GL_BGRA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
296 { GL_RGBA, GL_RGBA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
297 { GL_RGBA8, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
298 { GL_RGBA16, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
299 { GL_RGBA8I_EXT, GL_RGBA_INTEGER_EXT, GL_BYTE, kChar },
300 { GL_RGBA16I_EXT, GL_RGBA_INTEGER_EXT, GL_SHORT, kShort },
301 { GL_RGBA32I_EXT, GL_RGBA_INTEGER_EXT, GL_INT, kInt },
302 { GL_RGBA8UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, kUChar },
303 { GL_RGBA16UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_SHORT, kUShort },
304 { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT, kUInt },
305 { GL_RGBA32F_ARB, GL_RGBA, GL_FLOAT, kFloat }
306 #endif
307 };
308
309 size_t sizes[] = { 16, 32, 64, 128, 256, 512, 1024, 2048, 4096 };
310
311 size_t fmtIdx, tgtIdx;
312 int error = 0;
313 size_t iter = 6;
314 RandomSeed seed(gRandomSeed);
315
316 // Check if images are supported
317 if (checkForImageSupport(device)) {
318 log_info("Device does not support images. Skipping test.\n");
319 return 0;
320 }
321
322 // Loop through a set of GL formats, testing a set of sizes against each one
323 for( fmtIdx = 0; fmtIdx < sizeof( formats ) / sizeof( formats[ 0 ] ); fmtIdx++ )
324 {
325 for( tgtIdx = 0; tgtIdx < sizeof( targets ) / sizeof( targets[ 0 ] ); tgtIdx++ )
326 {
327 size_t i;
328
329 log_info( "Testing image read cubemap for GL format %s : %s : %s : %s\n\n",
330 GetGLTargetName( targets[ tgtIdx ] ),
331 GetGLFormatName( formats[ fmtIdx ].internal ),
332 GetGLBaseFormatName( formats[ fmtIdx ].format ),
333 GetGLTypeName( formats[ fmtIdx ].datatype ) );
334
335 for( i = 0; i < iter; i++ )
336 {
337 if( test_image_format_read( context, queue, sizes[i], sizes[i],
338 targets[ tgtIdx ],
339 formats[ fmtIdx ].format,
340 formats[ fmtIdx ].internal,
341 formats[ fmtIdx ].datatype,
342 formats[ fmtIdx ].type, seed ) )
343 {
344 log_error( "ERROR: Image read cubemap test failed for %s : %s : %s : %s\n\n",
345 GetGLTargetName( targets[ tgtIdx ] ),
346 GetGLFormatName( formats[ fmtIdx ].internal ),
347 GetGLBaseFormatName( formats[ fmtIdx ].format ),
348 GetGLTypeName( formats[ fmtIdx ].datatype ) );
349
350 error++;
351 break; // Skip other sizes for this combination
352 }
353 }
354 if( i == iter )
355 {
356 log_info( "passed: Image read cubemap for GL format %s : %s : %s : %s\n\n",
357 GetGLTargetName( targets[ tgtIdx ] ),
358 GetGLFormatName( formats[ fmtIdx ].internal ),
359 GetGLBaseFormatName( formats[ fmtIdx ].format ),
360 GetGLTypeName( formats[ fmtIdx ].datatype ) );
361
362 }
363 else
364 break; // Skip other cube map targets; they're unlikely to pass either
365 }
366 }
367
368 return error;
369 }
370
371
372 #pragma mark -------------------- Write tests -------------------------
373
374
test_cl_image_write(cl_context context,cl_command_queue queue,cl_mem clImage,size_t imageWidth,size_t imageHeight,cl_image_format * outFormat,ExplicitType * outType,void ** outSourceBuffer,MTdata d)375 int test_cl_image_write( cl_context context, cl_command_queue queue, cl_mem clImage,
376 size_t imageWidth, size_t imageHeight, cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer, MTdata d )
377 {
378 clProgramWrapper program;
379 clKernelWrapper kernel;
380 clMemWrapper inStream;
381
382 int error;
383 size_t threads[ 2 ], localThreads[ 2 ];
384 char kernelSource[10240];
385 char *programPtr;
386
387 // Determine data type and format that CL came up with
388 error = clGetImageInfo( clImage, CL_IMAGE_FORMAT, sizeof( cl_image_format ), outFormat, NULL );
389 test_error( error, "Unable to get CL image format" );
390
391 /* Create the source */
392 *outType = get_write_kernel_type( outFormat );
393 size_t channelSize = get_explicit_type_size( *outType );
394
395 const char* suffix = get_kernel_suffix( outFormat );
396 const char* convert = get_write_conversion( outFormat, *outType );
397
398 sprintf( kernelSource, imageWriteKernelPattern, get_explicit_type_name( *outType ), get_explicit_type_name( *outType ), suffix, convert);
399
400 #ifdef GLES_DEBUG
401 log_info("-- start cl image write kernel --\n");
402 log_info("%s", kernelSource);
403 log_info("-- end cl image write kernel --\n");
404 #endif
405
406 /* Create kernel */
407 programPtr = kernelSource;
408 if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
409 {
410 return -1;
411 }
412
413 // Generate some source data based on the input type we need
414 *outSourceBuffer = CreateRandomData(*outType, imageWidth * imageHeight * 4, d);
415
416 // Create a vanilla input buffer
417 inStream = clCreateBuffer( context, CL_MEM_COPY_HOST_PTR, channelSize * 4 * imageWidth * imageHeight, *outSourceBuffer, &error );
418 test_error( error, "Unable to create output buffer" );
419
420 /* Assign streams and execute */
421 clSamplerWrapper sampler = clCreateSampler( context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
422 test_error( error, "Unable to create sampler" );
423
424 error = clSetKernelArg( kernel, 0, sizeof( inStream ), &inStream );
425 test_error( error, "Unable to set kernel arguments" );
426 error = clSetKernelArg( kernel, 1, sizeof( clImage ), &clImage );
427 test_error( error, "Unable to set kernel arguments" );
428
429 glFlush();
430
431 error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &clImage, 0, NULL, NULL);
432 test_error( error, "Unable to acquire GL obejcts");
433
434 /* Run the kernel */
435 threads[ 0 ] = imageWidth;
436 threads[ 1 ] = imageHeight;
437
438 error = get_max_common_2D_work_group_size( context, kernel, threads, localThreads );
439 test_error( error, "Unable to get work group size to use" );
440
441 error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, localThreads, 0, NULL, NULL );
442 test_error( error, "Unable to execute test kernel" );
443
444 clEventWrapper event;
445 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &clImage, 0, NULL, &event );
446 test_error(error, "clEnqueueReleaseGLObjects failed");
447
448 error = clWaitForEvents( 1, &event );
449 test_error(error, "clWaitForEvents failed");
450
451 #ifdef GLES_DEBUG
452 int i;
453 size_t origin[] = {0, 0, 0,};
454 size_t region[] = {imageWidth, imageHeight, 1 };
455 void* cldata = malloc( channelSize * 4 * imageWidth * imageHeight );
456 clEnqueueReadImage( queue, clImage, 1, origin, region, 0, 0, cldata, 0, 0, 0);
457 log_info("- start CL Image Data -- \n");
458 DumpGLBuffer(GetGLTypeForExplicitType(*outType), imageWidth, imageHeight, cldata);
459 log_info("- end CL Image Data -- \n");
460 free(cldata);
461 #endif
462
463 // All done!
464 return 0;
465 }
466
test_image_write(cl_context context,cl_command_queue queue,GLenum glTarget,GLuint glTexture,size_t imageWidth,size_t imageHeight,cl_image_format * outFormat,ExplicitType * outType,void ** outSourceBuffer,MTdata d)467 int test_image_write( cl_context context, cl_command_queue queue, GLenum glTarget, GLuint glTexture,
468 size_t imageWidth, size_t imageHeight, cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer, MTdata d )
469 {
470 int error;
471
472 // Create a CL image from the supplied GL texture
473 clMemWrapper image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_WRITE_ONLY, glTarget, 0, glTexture, &error );
474 if( error != CL_SUCCESS )
475 {
476 print_error( error, "Unable to create CL image from GL texture" );
477 #ifndef GL_ES_VERSION_2_0
478 GLint fmt;
479 glGetTexLevelParameteriv( glTarget, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt );
480 log_error( " Supplied GL texture was baseformat %s and internalformat %s\n", GetGLBaseFormatName( fmt ), GetGLFormatName( fmt ) );
481 #endif
482 return error;
483 }
484
485 return test_cl_image_write( context, queue, image, imageWidth, imageHeight, outFormat, outType, outSourceBuffer, d );
486 }
487
488
test_image_format_write(cl_context context,cl_command_queue queue,size_t width,size_t height,GLenum target,GLenum format,GLenum internalFormat,GLenum glType,ExplicitType type,MTdata d)489 int test_image_format_write( cl_context context, cl_command_queue queue,
490 size_t width, size_t height, GLenum target,
491 GLenum format, GLenum internalFormat,
492 GLenum glType, ExplicitType type, MTdata d )
493 {
494 int error;
495
496 // Create the GL texture
497 glTextureWrapper glTexture;
498 void *tmp = CreateGLTexture2D( width, height, target, format, internalFormat, glType, type, &glTexture, &error, true, d );
499 BufferOwningPtr<char> inputBuffer(tmp);
500 if( error != 0 )
501 {
502 return error;
503 }
504
505 /* skip formats not supported by OpenGL */
506 if(!tmp)
507 {
508 return 0;
509 }
510
511 // Run and get the results
512 cl_image_format clFormat;
513 ExplicitType sourceType;
514 void *outSourceBuffer;
515 error = test_image_write( context, queue, target, glTexture, width, height, &clFormat, &sourceType, (void **)&outSourceBuffer, d );
516 if( error != 0 )
517 return error;
518
519 BufferOwningPtr<char> actualSource(outSourceBuffer);
520
521 log_info( "- Write [%4d x %4d] : GL Texture : %s : %s : %s => CL Image : %s : %s \n", (int)width, (int)height,
522 GetGLFormatName( format ), GetGLFormatName( internalFormat ), GetGLTypeName( glType),
523 GetChannelOrderName( clFormat.image_channel_order ), GetChannelTypeName( clFormat.image_channel_data_type ));
524
525 // Now read the results from the GL texture
526 ExplicitType readType = type;
527 BufferOwningPtr<char> glResults( ReadGLTexture( target, glTexture, format, internalFormat, glType, readType, width, height ) );
528
529 // We have to convert our input buffer to the returned type, so we can validate.
530 BufferOwningPtr<char> convertedGLResults( convert_to_expected( glResults, width * height, readType, sourceType ) );
531
532 #ifdef GLES_DEBUG
533 log_info("- start read GL data -- \n");
534 DumpGLBuffer(glType, width, height, glResults);
535 log_info("- end read GL data -- \n");
536
537 log_info("- start converted data -- \n");
538 DumpGLBuffer(glType, width, height, convertedGLResults);
539 log_info("- end converted data -- \n");
540 #endif
541
542 // Now we validate
543 int valid = 0;
544 if(convertedGLResults) {
545 if( sourceType == kFloat )
546 valid = validate_float_results( actualSource, convertedGLResults, width, height );
547 else
548 valid = validate_integer_results( actualSource, convertedGLResults, width, height, get_explicit_type_size( readType ) );
549 }
550
551 return valid;
552 }
553
test_images_write(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)554 int test_images_write( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
555 {
556 GLenum targets[] =
557 #ifdef GL_ES_VERSION_2_0
558 { GL_TEXTURE_2D };
559 #else // GL_ES_VERSION_2_0
560 { GL_TEXTURE_2D, GL_TEXTURE_RECTANGLE_EXT };
561 #endif
562
563 struct {
564 GLenum internal;
565 GLenum format;
566 GLenum datatype;
567 ExplicitType type;
568
569 } formats[] = {
570 #ifdef GL_ES_VERSION_2_0
571 { GL_RGBA, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
572 { GL_RGBA, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
573 // XXX add others
574 #else // GL_ES_VERSION_2_0
575 { GL_RGBA, GL_BGRA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
576 { GL_RGBA, GL_RGBA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
577 { GL_RGBA8, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
578 { GL_RGBA16, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
579 { GL_RGBA8I_EXT, GL_RGBA_INTEGER_EXT, GL_BYTE, kChar },
580 { GL_RGBA16I_EXT, GL_RGBA_INTEGER_EXT, GL_SHORT, kShort },
581 { GL_RGBA32I_EXT, GL_RGBA_INTEGER_EXT, GL_INT, kInt },
582 { GL_RGBA8UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, kUChar },
583 { GL_RGBA16UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_SHORT, kUShort },
584 { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT, kUInt },
585 { GL_RGBA32F_ARB, GL_RGBA, GL_FLOAT, kFloat }
586 #endif
587 };
588
589 size_t fmtIdx, tgtIdx;
590 int error = 0;
591 size_t iter = 6;
592 RandomSeed seed(gRandomSeed);
593
594 // Check if images are supported
595 if (checkForImageSupport(device)) {
596 log_info("Device does not support images. Skipping test.\n");
597 return 0;
598 }
599
600 // Loop through a set of GL formats, testing a set of sizes against each one
601 for( fmtIdx = 0; fmtIdx < sizeof( formats ) / sizeof( formats[ 0 ] ); fmtIdx++ )
602 {
603 for( tgtIdx = 0; tgtIdx < sizeof( targets ) / sizeof( targets[ 0 ] ); tgtIdx++ )
604 {
605 log_info( "Testing image write test for %s : %s : %s : %s\n",
606 GetGLTargetName( targets[ tgtIdx ] ),
607 GetGLFormatName( formats[ fmtIdx ].internal ),
608 GetGLBaseFormatName( formats[ fmtIdx ].format ),
609 GetGLTypeName( formats[ fmtIdx ].datatype ) );
610
611 size_t i;
612 for( i = 0; i < iter; i++ )
613 {
614 size_t width = random_in_range( 16, 512, seed );
615 size_t height = random_in_range( 16, 512, seed );
616
617 if( targets[ tgtIdx ] == GL_TEXTURE_2D )
618 width = height;
619
620 if( test_image_format_write( context, queue, width, height,
621 targets[ tgtIdx ],
622 formats[ fmtIdx ].format,
623 formats[ fmtIdx ].internal,
624 formats[ fmtIdx ].datatype,
625 formats[ fmtIdx ].type, seed ) )
626 {
627 log_error( "ERROR: Image write test failed for %s : %s : %s : %s\n\n",
628 GetGLTargetName( targets[ tgtIdx ] ),
629 GetGLFormatName( formats[ fmtIdx ].internal ),
630 GetGLBaseFormatName( formats[ fmtIdx ].format ),
631 GetGLTypeName( formats[ fmtIdx ].datatype ) );
632
633 error++;
634 break; // Skip other sizes for this combination
635 }
636 }
637 if( i == 6 )
638 {
639 log_info( "passed: Image write for GL format %s : %s : %s : %s\n\n",
640 GetGLTargetName( targets[ tgtIdx ] ),
641 GetGLFormatName( formats[ fmtIdx ].internal ),
642 GetGLBaseFormatName( formats[ fmtIdx ].format ),
643 GetGLTypeName( formats[ fmtIdx ].datatype ) );
644
645 }
646 }
647 }
648
649 return error;
650 }
651
test_images_write_cube(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)652 int test_images_write_cube( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
653 {
654 GLenum targets[] = {
655 GL_TEXTURE_CUBE_MAP_POSITIVE_X,
656 GL_TEXTURE_CUBE_MAP_POSITIVE_Y,
657 GL_TEXTURE_CUBE_MAP_POSITIVE_Z,
658 GL_TEXTURE_CUBE_MAP_NEGATIVE_X,
659 GL_TEXTURE_CUBE_MAP_NEGATIVE_Y,
660 GL_TEXTURE_CUBE_MAP_NEGATIVE_Z };
661
662 struct {
663 GLenum internal;
664 GLenum format;
665 GLenum datatype;
666 ExplicitType type;
667
668 } formats[] = {
669 #ifdef GL_ES_VERSION_2_0
670 { GL_RGBA, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
671 { GL_RGBA, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
672 // XXX add others
673 #else // GL_ES_VERSION_2_0
674 { GL_RGBA, GL_BGRA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
675 { GL_RGBA, GL_RGBA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
676 { GL_RGBA8, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
677 { GL_RGBA16, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
678 { GL_RGBA8I_EXT, GL_RGBA_INTEGER_EXT, GL_BYTE, kChar },
679 { GL_RGBA16I_EXT, GL_RGBA_INTEGER_EXT, GL_SHORT, kShort },
680 { GL_RGBA32I_EXT, GL_RGBA_INTEGER_EXT, GL_INT, kInt },
681 { GL_RGBA8UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, kUChar },
682 { GL_RGBA16UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_SHORT, kUShort },
683 { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT, kUInt },
684 { GL_RGBA32F_ARB, GL_RGBA, GL_FLOAT, kFloat }
685 #endif
686 };
687
688 size_t sizes[] = { 16, 32, 64, 128, 256, 512, 1024, 2048, 4096 };
689
690 size_t fmtIdx, tgtIdx;
691 int error = 0;
692 size_t iter = 6;
693 RandomSeed seed( gRandomSeed );
694
695 // Check if images are supported
696 if (checkForImageSupport(device)) {
697 log_info("Device does not support images. Skipping test.\n");
698 return 0;
699 }
700
701 // Loop through a set of GL formats, testing a set of sizes against each one
702 for( fmtIdx = 0; fmtIdx < sizeof( formats ) / sizeof( formats[ 0 ] ); fmtIdx++ )
703 {
704 for( tgtIdx = 0; tgtIdx < sizeof( targets ) / sizeof( targets[ 0 ] ); tgtIdx++ )
705 {
706 size_t i;
707 log_info( "Testing image write cubemap test for %s : %s : %s : %s\n",
708 GetGLTargetName( targets[ tgtIdx ] ),
709 GetGLFormatName( formats[ fmtIdx ].internal ),
710 GetGLBaseFormatName( formats[ fmtIdx ].format ),
711 GetGLTypeName( formats[ fmtIdx ].datatype ) );
712
713 for( i = 0; i < iter; i++ )
714 {
715 if( test_image_format_write( context, queue, sizes[i], sizes[i],
716 targets[ tgtIdx ],
717 formats[ fmtIdx ].format,
718 formats[ fmtIdx ].internal,
719 formats[ fmtIdx ].datatype,
720 formats[ fmtIdx ].type, seed ) )
721 {
722 log_error( "ERROR: Image write cubemap test failed for %s : %s : %s : %s\n\n",
723 GetGLTargetName( targets[ tgtIdx ] ),
724 GetGLFormatName( formats[ fmtIdx ].internal ),
725 GetGLBaseFormatName( formats[ fmtIdx ].format ),
726 GetGLTypeName( formats[ fmtIdx ].datatype ) );
727
728
729 error++;
730 break; // Skip other sizes for this combination
731 }
732 }
733 if( i == iter )
734 {
735 log_info( "passed: Image write cubemap for GL format %s : %s : %s : %s\n\n",
736 GetGLTargetName( targets[ tgtIdx ] ),
737 GetGLFormatName( formats[ fmtIdx ].internal ),
738 GetGLBaseFormatName( formats[ fmtIdx ].format ),
739 GetGLTypeName( formats[ fmtIdx ].datatype ) );
740 }
741 else
742 break; // Skip other cube map targets; they're unlikely to pass either
743 }
744 }
745
746 return error;
747 }
748