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 #ifdef __APPLE__
373 #pragma mark -------------------- Write tests -------------------------
374 #endif
375
376
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)377 int test_cl_image_write( cl_context context, cl_command_queue queue, cl_mem clImage,
378 size_t imageWidth, size_t imageHeight, cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer, MTdata d )
379 {
380 clProgramWrapper program;
381 clKernelWrapper kernel;
382 clMemWrapper inStream;
383
384 int error;
385 size_t threads[ 2 ], localThreads[ 2 ];
386 char kernelSource[10240];
387 char *programPtr;
388
389 // Determine data type and format that CL came up with
390 error = clGetImageInfo( clImage, CL_IMAGE_FORMAT, sizeof( cl_image_format ), outFormat, NULL );
391 test_error( error, "Unable to get CL image format" );
392
393 /* Create the source */
394 *outType = get_write_kernel_type( outFormat );
395 size_t channelSize = get_explicit_type_size( *outType );
396
397 const char* suffix = get_kernel_suffix( outFormat );
398 const char* convert = get_write_conversion( outFormat, *outType );
399
400 sprintf( kernelSource, imageWriteKernelPattern, get_explicit_type_name( *outType ), get_explicit_type_name( *outType ), suffix, convert);
401
402 #ifdef GLES_DEBUG
403 log_info("-- start cl image write kernel --\n");
404 log_info("%s", kernelSource);
405 log_info("-- end cl image write kernel --\n");
406 #endif
407
408 /* Create kernel */
409 programPtr = kernelSource;
410 if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
411 {
412 return -1;
413 }
414
415 // Generate some source data based on the input type we need
416 *outSourceBuffer = CreateRandomData(*outType, imageWidth * imageHeight * 4, d);
417
418 // Create a vanilla input buffer
419 inStream = clCreateBuffer( context, CL_MEM_COPY_HOST_PTR, channelSize * 4 * imageWidth * imageHeight, *outSourceBuffer, &error );
420 test_error( error, "Unable to create output buffer" );
421
422 /* Assign streams and execute */
423 clSamplerWrapper sampler = clCreateSampler( context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
424 test_error( error, "Unable to create sampler" );
425
426 error = clSetKernelArg( kernel, 0, sizeof( inStream ), &inStream );
427 test_error( error, "Unable to set kernel arguments" );
428 error = clSetKernelArg( kernel, 1, sizeof( clImage ), &clImage );
429 test_error( error, "Unable to set kernel arguments" );
430
431 glFlush();
432
433 error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &clImage, 0, NULL, NULL);
434 test_error( error, "Unable to acquire GL obejcts");
435
436 /* Run the kernel */
437 threads[ 0 ] = imageWidth;
438 threads[ 1 ] = imageHeight;
439
440 error = get_max_common_2D_work_group_size( context, kernel, threads, localThreads );
441 test_error( error, "Unable to get work group size to use" );
442
443 error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, localThreads, 0, NULL, NULL );
444 test_error( error, "Unable to execute test kernel" );
445
446 clEventWrapper event;
447 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &clImage, 0, NULL, &event );
448 test_error(error, "clEnqueueReleaseGLObjects failed");
449
450 error = clWaitForEvents( 1, &event );
451 test_error(error, "clWaitForEvents failed");
452
453 #ifdef GLES_DEBUG
454 int i;
455 size_t origin[] = {0, 0, 0,};
456 size_t region[] = {imageWidth, imageHeight, 1 };
457 void* cldata = malloc( channelSize * 4 * imageWidth * imageHeight );
458 clEnqueueReadImage( queue, clImage, 1, origin, region, 0, 0, cldata, 0, 0, 0);
459 log_info("- start CL Image Data -- \n");
460 DumpGLBuffer(GetGLTypeForExplicitType(*outType), imageWidth, imageHeight, cldata);
461 log_info("- end CL Image Data -- \n");
462 free(cldata);
463 #endif
464
465 // All done!
466 return 0;
467 }
468
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)469 int test_image_write( cl_context context, cl_command_queue queue, GLenum glTarget, GLuint glTexture,
470 size_t imageWidth, size_t imageHeight, cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer, MTdata d )
471 {
472 int error;
473
474 // Create a CL image from the supplied GL texture
475 clMemWrapper image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_WRITE_ONLY, glTarget, 0, glTexture, &error );
476 if( error != CL_SUCCESS )
477 {
478 print_error( error, "Unable to create CL image from GL texture" );
479 #ifndef GL_ES_VERSION_2_0
480 GLint fmt;
481 glGetTexLevelParameteriv( glTarget, 0, GL_TEXTURE_INTERNAL_FORMAT, &fmt );
482 log_error( " Supplied GL texture was baseformat %s and internalformat %s\n", GetGLBaseFormatName( fmt ), GetGLFormatName( fmt ) );
483 #endif
484 return error;
485 }
486
487 return test_cl_image_write( context, queue, image, imageWidth, imageHeight, outFormat, outType, outSourceBuffer, d );
488 }
489
490
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)491 int test_image_format_write( cl_context context, cl_command_queue queue,
492 size_t width, size_t height, GLenum target,
493 GLenum format, GLenum internalFormat,
494 GLenum glType, ExplicitType type, MTdata d )
495 {
496 int error;
497
498 // Create the GL texture
499 glTextureWrapper glTexture;
500 void *tmp = CreateGLTexture2D( width, height, target, format, internalFormat, glType, type, &glTexture, &error, true, d );
501 BufferOwningPtr<char> inputBuffer(tmp);
502 if( error != 0 )
503 {
504 return error;
505 }
506
507 /* skip formats not supported by OpenGL */
508 if(!tmp)
509 {
510 return 0;
511 }
512
513 // Run and get the results
514 cl_image_format clFormat;
515 ExplicitType sourceType;
516 void *outSourceBuffer;
517 error = test_image_write( context, queue, target, glTexture, width, height, &clFormat, &sourceType, (void **)&outSourceBuffer, d );
518 if( error != 0 )
519 return error;
520
521 BufferOwningPtr<char> actualSource(outSourceBuffer);
522
523 log_info( "- Write [%4d x %4d] : GL Texture : %s : %s : %s => CL Image : %s : %s \n", (int)width, (int)height,
524 GetGLFormatName( format ), GetGLFormatName( internalFormat ), GetGLTypeName( glType),
525 GetChannelOrderName( clFormat.image_channel_order ), GetChannelTypeName( clFormat.image_channel_data_type ));
526
527 // Now read the results from the GL texture
528 ExplicitType readType = type;
529 BufferOwningPtr<char> glResults( ReadGLTexture( target, glTexture, format, internalFormat, glType, readType, width, height ) );
530
531 // We have to convert our input buffer to the returned type, so we can validate.
532 BufferOwningPtr<char> convertedGLResults( convert_to_expected( glResults, width * height, readType, sourceType ) );
533
534 #ifdef GLES_DEBUG
535 log_info("- start read GL data -- \n");
536 DumpGLBuffer(glType, width, height, glResults);
537 log_info("- end read GL data -- \n");
538
539 log_info("- start converted data -- \n");
540 DumpGLBuffer(glType, width, height, convertedGLResults);
541 log_info("- end converted data -- \n");
542 #endif
543
544 // Now we validate
545 int valid = 0;
546 if(convertedGLResults) {
547 if( sourceType == kFloat )
548 valid = validate_float_results( actualSource, convertedGLResults, width, height );
549 else
550 valid = validate_integer_results( actualSource, convertedGLResults, width, height, get_explicit_type_size( readType ) );
551 }
552
553 return valid;
554 }
555
test_images_write(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)556 int test_images_write( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
557 {
558 GLenum targets[] =
559 #ifdef GL_ES_VERSION_2_0
560 { GL_TEXTURE_2D };
561 #else // GL_ES_VERSION_2_0
562 { GL_TEXTURE_2D, GL_TEXTURE_RECTANGLE_EXT };
563 #endif
564
565 struct {
566 GLenum internal;
567 GLenum format;
568 GLenum datatype;
569 ExplicitType type;
570
571 } formats[] = {
572 #ifdef GL_ES_VERSION_2_0
573 { GL_RGBA, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
574 { GL_RGBA, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
575 // XXX add others
576 #else // GL_ES_VERSION_2_0
577 { GL_RGBA, GL_BGRA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
578 { GL_RGBA, GL_RGBA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
579 { GL_RGBA8, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
580 { GL_RGBA16, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
581 { GL_RGBA8I_EXT, GL_RGBA_INTEGER_EXT, GL_BYTE, kChar },
582 { GL_RGBA16I_EXT, GL_RGBA_INTEGER_EXT, GL_SHORT, kShort },
583 { GL_RGBA32I_EXT, GL_RGBA_INTEGER_EXT, GL_INT, kInt },
584 { GL_RGBA8UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, kUChar },
585 { GL_RGBA16UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_SHORT, kUShort },
586 { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT, kUInt },
587 { GL_RGBA32F_ARB, GL_RGBA, GL_FLOAT, kFloat }
588 #endif
589 };
590
591 size_t fmtIdx, tgtIdx;
592 int error = 0;
593 size_t iter = 6;
594 RandomSeed seed(gRandomSeed);
595
596 // Check if images are supported
597 if (checkForImageSupport(device)) {
598 log_info("Device does not support images. Skipping test.\n");
599 return 0;
600 }
601
602 // Loop through a set of GL formats, testing a set of sizes against each one
603 for( fmtIdx = 0; fmtIdx < sizeof( formats ) / sizeof( formats[ 0 ] ); fmtIdx++ )
604 {
605 for( tgtIdx = 0; tgtIdx < sizeof( targets ) / sizeof( targets[ 0 ] ); tgtIdx++ )
606 {
607 log_info( "Testing image write test for %s : %s : %s : %s\n",
608 GetGLTargetName( targets[ tgtIdx ] ),
609 GetGLFormatName( formats[ fmtIdx ].internal ),
610 GetGLBaseFormatName( formats[ fmtIdx ].format ),
611 GetGLTypeName( formats[ fmtIdx ].datatype ) );
612
613 size_t i;
614 for( i = 0; i < iter; i++ )
615 {
616 size_t width = random_in_range( 16, 512, seed );
617 size_t height = random_in_range( 16, 512, seed );
618
619 if( targets[ tgtIdx ] == GL_TEXTURE_2D )
620 width = height;
621
622 if( test_image_format_write( context, queue, width, height,
623 targets[ tgtIdx ],
624 formats[ fmtIdx ].format,
625 formats[ fmtIdx ].internal,
626 formats[ fmtIdx ].datatype,
627 formats[ fmtIdx ].type, seed ) )
628 {
629 log_error( "ERROR: Image write test failed for %s : %s : %s : %s\n\n",
630 GetGLTargetName( targets[ tgtIdx ] ),
631 GetGLFormatName( formats[ fmtIdx ].internal ),
632 GetGLBaseFormatName( formats[ fmtIdx ].format ),
633 GetGLTypeName( formats[ fmtIdx ].datatype ) );
634
635 error++;
636 break; // Skip other sizes for this combination
637 }
638 }
639 if( i == 6 )
640 {
641 log_info( "passed: Image write for GL format %s : %s : %s : %s\n\n",
642 GetGLTargetName( targets[ tgtIdx ] ),
643 GetGLFormatName( formats[ fmtIdx ].internal ),
644 GetGLBaseFormatName( formats[ fmtIdx ].format ),
645 GetGLTypeName( formats[ fmtIdx ].datatype ) );
646
647 }
648 }
649 }
650
651 return error;
652 }
653
test_images_write_cube(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)654 int test_images_write_cube( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
655 {
656 GLenum targets[] = {
657 GL_TEXTURE_CUBE_MAP_POSITIVE_X,
658 GL_TEXTURE_CUBE_MAP_POSITIVE_Y,
659 GL_TEXTURE_CUBE_MAP_POSITIVE_Z,
660 GL_TEXTURE_CUBE_MAP_NEGATIVE_X,
661 GL_TEXTURE_CUBE_MAP_NEGATIVE_Y,
662 GL_TEXTURE_CUBE_MAP_NEGATIVE_Z };
663
664 struct {
665 GLenum internal;
666 GLenum format;
667 GLenum datatype;
668 ExplicitType type;
669
670 } formats[] = {
671 #ifdef GL_ES_VERSION_2_0
672 { GL_RGBA, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
673 { GL_RGBA, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
674 // XXX add others
675 #else // GL_ES_VERSION_2_0
676 { GL_RGBA, GL_BGRA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
677 { GL_RGBA, GL_RGBA, GL_UNSIGNED_INT_8_8_8_8_REV, kUChar },
678 { GL_RGBA8, GL_RGBA, GL_UNSIGNED_BYTE, kUChar },
679 { GL_RGBA16, GL_RGBA, GL_UNSIGNED_SHORT, kUShort },
680 { GL_RGBA8I_EXT, GL_RGBA_INTEGER_EXT, GL_BYTE, kChar },
681 { GL_RGBA16I_EXT, GL_RGBA_INTEGER_EXT, GL_SHORT, kShort },
682 { GL_RGBA32I_EXT, GL_RGBA_INTEGER_EXT, GL_INT, kInt },
683 { GL_RGBA8UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_BYTE, kUChar },
684 { GL_RGBA16UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_SHORT, kUShort },
685 { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT, kUInt },
686 { GL_RGBA32F_ARB, GL_RGBA, GL_FLOAT, kFloat }
687 #endif
688 };
689
690 size_t sizes[] = { 16, 32, 64, 128, 256, 512, 1024, 2048, 4096 };
691
692 size_t fmtIdx, tgtIdx;
693 int error = 0;
694 size_t iter = 6;
695 RandomSeed seed( gRandomSeed );
696
697 // Check if images are supported
698 if (checkForImageSupport(device)) {
699 log_info("Device does not support images. Skipping test.\n");
700 return 0;
701 }
702
703 // Loop through a set of GL formats, testing a set of sizes against each one
704 for( fmtIdx = 0; fmtIdx < sizeof( formats ) / sizeof( formats[ 0 ] ); fmtIdx++ )
705 {
706 for( tgtIdx = 0; tgtIdx < sizeof( targets ) / sizeof( targets[ 0 ] ); tgtIdx++ )
707 {
708 size_t i;
709 log_info( "Testing image write cubemap test for %s : %s : %s : %s\n",
710 GetGLTargetName( targets[ tgtIdx ] ),
711 GetGLFormatName( formats[ fmtIdx ].internal ),
712 GetGLBaseFormatName( formats[ fmtIdx ].format ),
713 GetGLTypeName( formats[ fmtIdx ].datatype ) );
714
715 for( i = 0; i < iter; i++ )
716 {
717 if( test_image_format_write( context, queue, sizes[i], sizes[i],
718 targets[ tgtIdx ],
719 formats[ fmtIdx ].format,
720 formats[ fmtIdx ].internal,
721 formats[ fmtIdx ].datatype,
722 formats[ fmtIdx ].type, seed ) )
723 {
724 log_error( "ERROR: Image write cubemap test failed for %s : %s : %s : %s\n\n",
725 GetGLTargetName( targets[ tgtIdx ] ),
726 GetGLFormatName( formats[ fmtIdx ].internal ),
727 GetGLBaseFormatName( formats[ fmtIdx ].format ),
728 GetGLTypeName( formats[ fmtIdx ].datatype ) );
729
730
731 error++;
732 break; // Skip other sizes for this combination
733 }
734 }
735 if( i == iter )
736 {
737 log_info( "passed: Image write cubemap for GL format %s : %s : %s : %s\n\n",
738 GetGLTargetName( targets[ tgtIdx ] ),
739 GetGLFormatName( formats[ fmtIdx ].internal ),
740 GetGLBaseFormatName( formats[ fmtIdx ].format ),
741 GetGLTypeName( formats[ fmtIdx ].datatype ) );
742 }
743 else
744 break; // Skip other cube map targets; they're unlikely to pass either
745 }
746 }
747
748 return error;
749 }
750