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 #if !defined(_WIN32)
19 #include <sys/mman.h>
20 #endif
21
22 extern cl_mem_flags gMemFlagsToUse;
23 extern int gtestTypesToRun;
24
25 extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
26 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
27
28 const char *readwrite1DArrayKernelSourcePattern =
29 "__kernel void sample_kernel( __global %s4 *input, read_write image1d_array_t output %s)\n"
30 "{\n"
31 " int tidX = get_global_id(0), tidY = get_global_id(1);\n"
32 "%s"
33 " write_image%s( output, (int2)( tidX, tidY )%s, input[ offset ]);\n"
34 "}";
35
36 const char *write1DArrayKernelSourcePattern =
37 "__kernel void sample_kernel( __global %s4 *input, write_only image1d_array_t output %s)\n"
38 "{\n"
39 " int tidX = get_global_id(0), tidY = get_global_id(1);\n"
40 "%s"
41 " write_image%s( output, (int2)( tidX, tidY ) %s, input[ offset ]);\n"
42 "}";
43
44 const char *offset1DArraySource =
45 " int offset = tidY*get_image_width(output) + tidX;\n";
46
47 const char *offset1DArrayLodSource =
48 " int width_lod = ( get_image_width(output) >> lod ) ? ( get_image_width(output) >> lod ) : 1;\n"
49 " int offset = tidY*width_lod + tidX;\n";
50
test_write_image_1D_array(cl_device_id device,cl_context context,cl_command_queue queue,cl_kernel kernel,image_descriptor * imageInfo,ExplicitType inputType,MTdata d)51 int test_write_image_1D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_kernel kernel,
52 image_descriptor *imageInfo, ExplicitType inputType, MTdata d )
53 {
54 int totalErrors = 0;
55 size_t num_flags = 0;
56 const cl_mem_flags *mem_flag_types = NULL;
57 const char * *mem_flag_names = NULL;
58 const cl_mem_flags write_only_mem_flag_types[2] = { CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE };
59 const char * write_only_mem_flag_names[2] = { "CL_MEM_WRITE_ONLY", "CL_MEM_READ_WRITE" };
60 const cl_mem_flags read_write_mem_flag_types[1] = { CL_MEM_READ_WRITE};
61 const char * read_write_mem_flag_names[1] = { "CL_MEM_READ_WRITE"};
62
63 if(gtestTypesToRun & kWriteTests)
64 {
65 mem_flag_types = write_only_mem_flag_types;
66 mem_flag_names = write_only_mem_flag_names;
67 num_flags = sizeof( write_only_mem_flag_types ) / sizeof( write_only_mem_flag_types[0] );
68 }
69 else
70 {
71 mem_flag_types = read_write_mem_flag_types;
72 mem_flag_names = read_write_mem_flag_names;
73 num_flags = sizeof( read_write_mem_flag_types ) / sizeof( read_write_mem_flag_types[0] );
74 }
75
76 size_t pixelSize = get_pixel_size( imageInfo->format );
77
78 for( size_t mem_flag_index = 0; mem_flag_index < num_flags; mem_flag_index++ )
79 {
80 int error;
81 size_t threads[2];
82 bool verifyRounding = false;
83 int forceCorrectlyRoundedWrites = 0;
84
85 #if defined( __APPLE__ )
86 // Require Apple's CPU implementation to be correctly rounded, not just within 0.6
87 if( GetDeviceType(device) == CL_DEVICE_TYPE_CPU )
88 forceCorrectlyRoundedWrites = 1;
89 #endif
90
91 if( imageInfo->format->image_channel_data_type == CL_HALF_FLOAT )
92 if( DetectFloatToHalfRoundingMode(queue) )
93 return 1;
94
95 BufferOwningPtr<char> maxImageUseHostPtrBackingStore, imageValues;
96
97 create_random_image_data( inputType, imageInfo, imageValues, d );
98
99 if(!gTestMipmaps)
100 {
101 if( inputType == kFloat && imageInfo->format->image_channel_data_type != CL_FLOAT && imageInfo->format->image_channel_data_type != CL_HALF_FLOAT )
102 {
103 /* Pilot data for sRGB images */
104 if(is_sRGBA_order(imageInfo->format->image_channel_order))
105 {
106 // We want to generate ints (mostly) in range of the target format which should be [0,255]
107 // However the range chosen here is [-test_range_ext, 255 + test_range_ext] so that
108 // it can test some out-of-range data points
109 const unsigned int test_range_ext = 16;
110 int formatMin = 0 - test_range_ext;
111 int formatMax = 255 + test_range_ext;
112 int pixel_value = 0;
113
114 // First, fill with arbitrary floats
115 for( size_t y = 0; y < imageInfo->arraySize; y++ )
116 {
117 float *inputValues = (float *)(char*)imageValues + y * imageInfo->width * 4;
118 for( size_t i = 0; i < imageInfo->width * 4; i++ )
119 {
120 pixel_value = random_in_range( formatMin, (int)formatMax, d );
121 inputValues[ i ] = (float)(pixel_value/255.0f);
122 }
123 }
124
125 // Throw a few extra test values in there
126 float *inputValues = (float *)(char*)imageValues;
127 size_t i = 0;
128
129 // Piloting some debug inputs.
130 inputValues[ i++ ] = -0.5f;
131 inputValues[ i++ ] = 0.5f;
132 inputValues[ i++ ] = 2.f;
133 inputValues[ i++ ] = 0.5f;
134
135 // Also fill in the first few vectors with some deliberate tests to determine the rounding mode
136 // is correct
137 if( imageInfo->width > 12 )
138 {
139 float formatMax = (float)get_format_max_int( imageInfo->format );
140 inputValues[ i++ ] = 4.0f / formatMax;
141 inputValues[ i++ ] = 4.3f / formatMax;
142 inputValues[ i++ ] = 4.5f / formatMax;
143 inputValues[ i++ ] = 4.7f / formatMax;
144 inputValues[ i++ ] = 5.0f / formatMax;
145 inputValues[ i++ ] = 5.3f / formatMax;
146 inputValues[ i++ ] = 5.5f / formatMax;
147 inputValues[ i++ ] = 5.7f / formatMax;
148 }
149 }
150 else
151 {
152 // First, fill with arbitrary floats
153 for( size_t y = 0; y < imageInfo->arraySize; y++ )
154 {
155 float *inputValues = (float *)(char*)imageValues + y * imageInfo->width * 4;
156 for( size_t i = 0; i < imageInfo->width * 4; i++ )
157 inputValues[ i ] = get_random_float( -0.1f, 1.1f, d );
158 }
159
160 // Throw a few extra test values in there
161 float *inputValues = (float *)(char*)imageValues;
162 size_t i = 0;
163 inputValues[ i++ ] = -0.0000000000009f;
164 inputValues[ i++ ] = 1.f;
165 inputValues[ i++ ] = -1.f;
166 inputValues[ i++ ] = 2.f;
167
168 // Also fill in the first few vectors with some deliberate tests to determine the rounding mode
169 // is correct
170 if( imageInfo->width > 12 )
171 {
172 float formatMax = (float)get_format_max_int( imageInfo->format );
173 inputValues[ i++ ] = 4.0f / formatMax;
174 inputValues[ i++ ] = 4.3f / formatMax;
175 inputValues[ i++ ] = 4.5f / formatMax;
176 inputValues[ i++ ] = 4.7f / formatMax;
177 inputValues[ i++ ] = 5.0f / formatMax;
178 inputValues[ i++ ] = 5.3f / formatMax;
179 inputValues[ i++ ] = 5.5f / formatMax;
180 inputValues[ i++ ] = 5.7f / formatMax;
181 verifyRounding = true;
182 }
183 }
184 }
185 else if( inputType == kUInt )
186 {
187 unsigned int *inputValues = (unsigned int*)(char*)imageValues;
188 size_t i = 0;
189 inputValues[ i++ ] = 0;
190 inputValues[ i++ ] = 65535;
191 inputValues[ i++ ] = 7271820;
192 inputValues[ i++ ] = 0;
193 }
194 }
195
196 // Construct testing sources
197 clProtectedImage protImage;
198 clMemWrapper unprotImage;
199 cl_mem image;
200
201 if( gMemFlagsToUse == CL_MEM_USE_HOST_PTR )
202 {
203 // clProtectedImage uses USE_HOST_PTR, so just rely on that for the testing (via Ian)
204 // Do not use protected images for max image size test since it rounds the row size to a page size
205 if (gTestMaxImages) {
206 create_random_image_data( inputType, imageInfo, maxImageUseHostPtrBackingStore, d );
207
208 unprotImage = create_image_1d_array( context, mem_flag_types[mem_flag_index] | CL_MEM_USE_HOST_PTR, imageInfo->format,
209 imageInfo->width, imageInfo->arraySize, 0, 0,
210 maxImageUseHostPtrBackingStore, &error );
211 } else {
212 error = protImage.Create( context, (cl_mem_object_type)CL_MEM_OBJECT_IMAGE1D_ARRAY, mem_flag_types[mem_flag_index], imageInfo->format, imageInfo->width, 1, 1, imageInfo->arraySize );
213 }
214 if( error != CL_SUCCESS )
215 {
216 log_error( "ERROR: Unable to create 1D image array of size %ld x %ld pitch %ld (%s, %s)\n", imageInfo->width, imageInfo->arraySize,
217 imageInfo->rowPitch, IGetErrorString( error ), mem_flag_names[mem_flag_index] );
218 return error;
219 }
220
221 if (gTestMaxImages)
222 image = (cl_mem)unprotImage;
223 else
224 image = (cl_mem)protImage;
225 }
226 else // Either CL_MEM_ALLOC_HOST_PTR, CL_MEM_COPY_HOST_PTR or none
227 {
228 // Note: if ALLOC_HOST_PTR is used, the driver allocates memory that can be accessed by the host, but otherwise
229 // it works just as if no flag is specified, so we just do the same thing either way
230 // Note: if the flags is really CL_MEM_COPY_HOST_PTR, we want to remove it, because we don't want to copy any incoming data
231 if( gTestMipmaps )
232 {
233 cl_image_desc image_desc = {0};
234 image_desc.image_type = imageInfo->type;
235 image_desc.num_mip_levels = imageInfo->num_mip_levels;
236 image_desc.image_width = imageInfo->width;
237 image_desc.image_array_size = imageInfo->arraySize;
238
239 unprotImage = clCreateImage( context, mem_flag_types[mem_flag_index] | ( gMemFlagsToUse & ~(CL_MEM_COPY_HOST_PTR) ),
240 imageInfo->format, &image_desc, NULL, &error);
241 if( error != CL_SUCCESS )
242 {
243 log_error( "ERROR: Unable to create %d level 1D image array of size %ld x %ld (%s, %s)\n", imageInfo->num_mip_levels, imageInfo->width, imageInfo->arraySize,
244 IGetErrorString( error ), mem_flag_names[mem_flag_index] );
245 return error;
246 }
247 }
248 else
249 {
250 unprotImage = create_image_1d_array( context, mem_flag_types[mem_flag_index] | ( gMemFlagsToUse & ~(CL_MEM_COPY_HOST_PTR) ), imageInfo->format,
251 imageInfo->width, imageInfo->arraySize, 0, 0,
252 imageValues, &error );
253 if( error != CL_SUCCESS )
254 {
255 log_error( "ERROR: Unable to create 1D image array of size %ld x %ld pitch %ld (%s, %s)\n", imageInfo->width, imageInfo->arraySize,
256 imageInfo->rowPitch, IGetErrorString( error ), mem_flag_names[mem_flag_index] );
257 return error;
258 }
259 }
260 image = unprotImage;
261 }
262
263 error = clSetKernelArg( kernel, 1, sizeof( cl_mem ), &image );
264 test_error( error, "Unable to set kernel arguments" );
265
266 size_t width_lod = imageInfo->width, nextLevelOffset = 0;
267 size_t origin[ 3 ] = { 0, 0, 0 };
268 size_t region[ 3 ] = { imageInfo->width, imageInfo->arraySize, 1 };
269 size_t resultSize;
270
271 for( int lod = 0; (gTestMipmaps && lod < imageInfo->num_mip_levels) || (!gTestMipmaps && lod < 1); lod++)
272 {
273 if(gTestMipmaps)
274 {
275 error = clSetKernelArg( kernel, 2, sizeof( int ), &lod );
276
277 }
278 // Run the kernel
279 threads[0] = (size_t)width_lod;
280 threads[1] = (size_t)imageInfo->arraySize;
281
282 clMemWrapper inputStream;
283
284 char *imagePtrOffset = imageValues + nextLevelOffset;
285 inputStream = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
286 get_explicit_type_size(inputType) * 4
287 * width_lod * imageInfo->arraySize,
288 imagePtrOffset, &error);
289 test_error( error, "Unable to create input buffer" );
290
291 // Set arguments
292 error = clSetKernelArg( kernel, 0, sizeof( cl_mem ), &inputStream );
293 test_error( error, "Unable to set kernel arguments" );
294
295 error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL );
296 test_error( error, "Unable to run kernel" );
297
298 // Get results
299 if( gTestMipmaps )
300 resultSize = width_lod * get_pixel_size(imageInfo->format) * imageInfo->arraySize;
301 else
302 resultSize = imageInfo->rowPitch * imageInfo->arraySize;
303
304 clProtectedArray PA(resultSize);
305 char *resultValues = (char *)((void *)PA);
306
307 if( gDebugTrace )
308 log_info( " reading results, %ld kbytes\n", (unsigned long)( resultSize / 1024 ) );
309
310
311 origin[2] = lod;
312 region[0] = width_lod;
313 error = clEnqueueReadImage( queue, image, CL_TRUE, origin, region,
314 gEnablePitch ? imageInfo->rowPitch : 0, gEnablePitch ? imageInfo->slicePitch : 0, resultValues, 0, NULL, NULL );
315 test_error( error, "Unable to read results from kernel" );
316 if( gDebugTrace )
317 log_info( " results read\n" );
318
319 // Validate results element by element
320 char *imagePtr = imageValues + nextLevelOffset;
321 int numTries = 5;
322 for( size_t y = 0, i = 0; y < imageInfo->arraySize; y++ )
323 {
324 char *resultPtr;
325 if( gTestMipmaps )
326 resultPtr = (char *)resultValues + y * width_lod * pixelSize;
327 else
328 resultPtr = (char*)resultValues + y * imageInfo->rowPitch;
329 for( size_t x = 0; x < width_lod; x++, i++ )
330 {
331 char resultBuffer[ 16 ]; // Largest format would be 4 channels * 4 bytes (32 bits) each
332
333 // Convert this pixel
334 if( inputType == kFloat )
335 pack_image_pixel( (float *)imagePtr, imageInfo->format, resultBuffer );
336 else if( inputType == kInt )
337 pack_image_pixel( (int *)imagePtr, imageInfo->format, resultBuffer );
338 else // if( inputType == kUInt )
339 pack_image_pixel( (unsigned int *)imagePtr, imageInfo->format, resultBuffer );
340
341 // Compare against the results
342 if(is_sRGBA_order(imageInfo->format->image_channel_order))
343 {
344 // Compare sRGB-mapped values
345 cl_float expected[4] = {0};
346 cl_float* input_values = (float*)imagePtr;
347 cl_uchar *actual = (cl_uchar*)resultPtr;
348 float max_err = MAX_lRGB_TO_sRGB_CONVERSION_ERROR;
349 float err[4] = {0.0f};
350
351 for( unsigned int j = 0; j < get_format_channel_count( imageInfo->format ); j++ )
352 {
353 if(j < 3)
354 {
355 expected[j] = sRGBmap(input_values[j]);
356 }
357 else // there is no sRGB conversion for alpha component if it exists
358 {
359 expected[j] = NORMALIZE(input_values[j], 255.0f);
360 }
361
362 err[j] = fabsf( expected[ j ] - actual[ j ] );
363 }
364
365 if ((err[0] > max_err) ||
366 (err[1] > max_err) ||
367 (err[2] > max_err) ||
368 (err[3] > 0)) // there is no conversion for alpha so the error should be zero
369 {
370 log_error( " Error: %g %g %g %g\n", err[0], err[1], err[2], err[3]);
371 log_error( " Input: %g %g %g %g\n", *((float *)imagePtr), *((float *)imagePtr + 1), *((float *)imagePtr + 2), *((float *)imagePtr + 3));
372 log_error( " Expected: %g %g %g %g\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] );
373 log_error( " Actual: %d %d %d %d\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] );
374 return 1;
375 }
376 }
377 else if( imageInfo->format->image_channel_data_type == CL_FLOAT )
378 {
379 float *expected = (float *)resultBuffer;
380 float *actual = (float *)resultPtr;
381
382 if( !validate_float_write_results( expected, actual, imageInfo ) )
383 {
384 unsigned int *e = (unsigned int *)resultBuffer;
385 unsigned int *a = (unsigned int *)resultPtr;
386 log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[ mem_flag_index ] );
387 log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] );
388 log_error( " Expected: %08x %08x %08x %08x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] );
389 log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] );
390 log_error( " Actual: %08x %08x %08x %08x\n", a[ 0 ], a[ 1 ], a[ 2 ], a[ 3 ] );
391 totalErrors++;
392 if( ( --numTries ) == 0 )
393 return 1;
394 }
395 }
396 else if( imageInfo->format->image_channel_data_type == CL_HALF_FLOAT )
397 {
398 cl_half *e = (cl_half *)resultBuffer;
399 cl_half *a = (cl_half *)resultPtr;
400 if( !validate_half_write_results( e, a, imageInfo ) )
401 {
402 totalErrors++;
403 log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[ mem_flag_index ] );
404 log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] );
405 log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[ 0 ], a[ 1 ], a[ 2 ], a[ 3 ] );
406 if( inputType == kFloat )
407 {
408 float *p = (float *)imagePtr;
409 log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
410 log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
411 }
412 if( ( --numTries ) == 0 )
413 return 1;
414 }
415 }
416 else
417 {
418 // Exact result passes every time
419 if( memcmp( resultBuffer, resultPtr, pixelSize ) != 0 )
420 {
421 // result is inexact. Calculate error
422 int failure = 1;
423 float errors[4] = {NAN, NAN, NAN, NAN};
424 pack_image_pixel_error( (float *)imagePtr, imageInfo->format, resultBuffer, errors );
425
426 // We are allowed 0.6 absolute error vs. infinitely precise for some normalized formats
427 if( 0 == forceCorrectlyRoundedWrites &&
428 (
429 imageInfo->format->image_channel_data_type == CL_UNORM_INT8 ||
430 imageInfo->format->image_channel_data_type == CL_UNORM_INT_101010 ||
431 imageInfo->format->image_channel_data_type == CL_UNORM_INT16 ||
432 imageInfo->format->image_channel_data_type == CL_SNORM_INT8 ||
433 imageInfo->format->image_channel_data_type == CL_SNORM_INT16
434 ))
435 {
436 if( ! (fabsf( errors[0] ) > 0.6f) && ! (fabsf( errors[1] ) > 0.6f) &&
437 ! (fabsf( errors[2] ) > 0.6f) && ! (fabsf( errors[3] ) > 0.6f) )
438 failure = 0;
439 }
440
441
442 if( failure )
443 {
444 totalErrors++;
445 // Is it our special rounding test?
446 if( verifyRounding && i >= 1 && i <= 2 )
447 {
448 // Try to guess what the rounding mode of the device really is based on what it returned
449 const char *deviceRounding = "unknown";
450 unsigned int deviceResults[8];
451 read_image_pixel<unsigned int>( resultPtr, imageInfo, 0, 0, 0, deviceResults, lod );
452 read_image_pixel<unsigned int>( resultPtr, imageInfo, 1, 0, 0, &deviceResults[ 4 ], lod );
453
454 if( deviceResults[ 0 ] == 4 && deviceResults[ 1 ] == 4 && deviceResults[ 2 ] == 4 && deviceResults[ 3 ] == 4 &&
455 deviceResults[ 4 ] == 5 && deviceResults[ 5 ] == 5 && deviceResults[ 6 ] == 5 && deviceResults[ 7 ] == 5 )
456 deviceRounding = "truncate";
457 else if( deviceResults[ 0 ] == 4 && deviceResults[ 1 ] == 4 && deviceResults[ 2 ] == 5 && deviceResults[ 3 ] == 5 &&
458 deviceResults[ 4 ] == 5 && deviceResults[ 5 ] == 5 && deviceResults[ 6 ] == 6 && deviceResults[ 7 ] == 6 )
459 deviceRounding = "round to nearest";
460 else if( deviceResults[ 0 ] == 4 && deviceResults[ 1 ] == 4 && deviceResults[ 2 ] == 4 && deviceResults[ 3 ] == 5 &&
461 deviceResults[ 4 ] == 5 && deviceResults[ 5 ] == 5 && deviceResults[ 6 ] == 6 && deviceResults[ 7 ] == 6 )
462 deviceRounding = "round to even";
463
464 log_error( "ERROR: Rounding mode sample (%ld) did not validate, probably due to the device's rounding mode being wrong (%s)\n", i, mem_flag_names[mem_flag_index] );
465 log_error( " Actual values rounded by device: %x %x %x %x %x %x %x %x\n", deviceResults[ 0 ], deviceResults[ 1 ], deviceResults[ 2 ], deviceResults[ 3 ],
466 deviceResults[ 4 ], deviceResults[ 5 ], deviceResults[ 6 ], deviceResults[ 7 ] );
467 log_error( " Rounding mode of device appears to be %s\n", deviceRounding );
468 return 1;
469 }
470 log_error( "ERROR: Sample %d (%d,%d) did not validate!\n", (int)i, (int)x, (int)y );
471 switch(imageInfo->format->image_channel_data_type)
472 {
473 case CL_UNORM_INT8:
474 case CL_SNORM_INT8:
475 case CL_UNSIGNED_INT8:
476 case CL_SIGNED_INT8:
477 log_error( " Expected: 0x%2.2x 0x%2.2x 0x%2.2x 0x%2.2x\n", ((cl_uchar*)resultBuffer)[0], ((cl_uchar*)resultBuffer)[1], ((cl_uchar*)resultBuffer)[2], ((cl_uchar*)resultBuffer)[3] );
478 log_error( " Actual: 0x%2.2x 0x%2.2x 0x%2.2x 0x%2.2x\n", ((cl_uchar*)resultPtr)[0], ((cl_uchar*)resultPtr)[1], ((cl_uchar*)resultPtr)[2], ((cl_uchar*)resultPtr)[3] );
479 log_error( " Error: %f %f %f %f\n", errors[0], errors[1], errors[2], errors[3] );
480 break;
481 case CL_UNORM_INT16:
482 case CL_SNORM_INT16:
483 case CL_UNSIGNED_INT16:
484 case CL_SIGNED_INT16:
485 #ifdef CL_SFIXED14_APPLE
486 case CL_SFIXED14_APPLE:
487 #endif
488 log_error( " Expected: 0x%4.4x 0x%4.4x 0x%4.4x 0x%4.4x\n", ((cl_ushort*)resultBuffer)[0], ((cl_ushort*)resultBuffer)[1], ((cl_ushort*)resultBuffer)[2], ((cl_ushort*)resultBuffer)[3] );
489 log_error( " Actual: 0x%4.4x 0x%4.4x 0x%4.4x 0x%4.4x\n", ((cl_ushort*)resultPtr)[0], ((cl_ushort*)resultPtr)[1], ((cl_ushort*)resultPtr)[2], ((cl_ushort*)resultPtr)[3] );
490 log_error( " Error: %f %f %f %f\n", errors[0], errors[1], errors[2], errors[3] );
491 break;
492 case CL_HALF_FLOAT:
493 log_error(" Expected: 0x%4.4x "
494 "0x%4.4x 0x%4.4x 0x%4.4x\n",
495 ((cl_half *)resultBuffer)[0],
496 ((cl_half *)resultBuffer)[1],
497 ((cl_half *)resultBuffer)[2],
498 ((cl_half *)resultBuffer)[3]);
499 log_error(" Actual: 0x%4.4x "
500 "0x%4.4x 0x%4.4x 0x%4.4x\n",
501 ((cl_half *)resultPtr)[0],
502 ((cl_half *)resultPtr)[1],
503 ((cl_half *)resultPtr)[2],
504 ((cl_half *)resultPtr)[3]);
505 log_error( " Ulps: %f %f %f %f\n", errors[0], errors[1], errors[2], errors[3] );
506 break;
507 case CL_UNSIGNED_INT32:
508 case CL_SIGNED_INT32:
509 log_error( " Expected: 0x%8.8x 0x%8.8x 0x%8.8x 0x%8.8x\n", ((cl_uint*)resultBuffer)[0], ((cl_uint*)resultBuffer)[1], ((cl_uint*)resultBuffer)[2], ((cl_uint*)resultBuffer)[3] );
510 log_error( " Actual: 0x%8.8x 0x%8.8x 0x%8.8x 0x%8.8x\n", ((cl_uint*)resultPtr)[0], ((cl_uint*)resultPtr)[1], ((cl_uint*)resultPtr)[2], ((cl_uint*)resultPtr)[3] );
511 break;
512 case CL_FLOAT:
513 log_error( " Expected: %a %a %a %a\n", ((cl_float*)resultBuffer)[0], ((cl_float*)resultBuffer)[1], ((cl_float*)resultBuffer)[2], ((cl_float*)resultBuffer)[3] );
514 log_error( " Actual: %a %a %a %a\n", ((cl_float*)resultPtr)[0], ((cl_float*)resultPtr)[1], ((cl_float*)resultPtr)[2], ((cl_float*)resultPtr)[3] );
515 log_error( " Ulps: %f %f %f %f\n", errors[0], errors[1], errors[2], errors[3] );
516 break;
517 }
518
519 float *v = (float *)imagePtr;
520 log_error( " src: %g %g %g %g\n", v[ 0 ], v[ 1], v[ 2 ], v[ 3 ] );
521 log_error( " : %a %a %a %a\n", v[ 0 ], v[ 1], v[ 2 ], v[ 3 ] );
522 log_error( " src: %12.24f %12.24f %12.24f %12.24f\n", v[0 ], v[ 1], v[ 2 ], v[ 3 ] );
523
524 if( ( --numTries ) == 0 )
525 return 1;
526 }
527 }
528 }
529 imagePtr += get_explicit_type_size( inputType ) * 4;
530 resultPtr += pixelSize;
531 }
532 }
533 {
534 nextLevelOffset += width_lod * imageInfo->arraySize * get_pixel_size(imageInfo->format);
535 width_lod = (width_lod >> 1) ? (width_lod >> 1) : 1;
536 }
537 }
538 }
539
540 // All done!
541 return totalErrors;
542 }
543
544
test_write_image_1D_array_set(cl_device_id device,cl_context context,cl_command_queue queue,const cl_image_format * format,ExplicitType inputType,MTdata d)545 int test_write_image_1D_array_set(cl_device_id device, cl_context context,
546 cl_command_queue queue,
547 const cl_image_format *format,
548 ExplicitType inputType, MTdata d)
549 {
550 char programSrc[10240];
551 const char *ptr;
552 const char *readFormat;
553 clProgramWrapper program;
554 clKernelWrapper kernel;
555 const char *KernelSourcePattern = NULL;
556 int error;
557
558 // Get our operating parameters
559 size_t maxWidth, maxArraySize;
560 cl_ulong maxAllocSize, memSize;
561 size_t pixelSize;
562
563 image_descriptor imageInfo = { 0x0 };
564
565 imageInfo.format = format;
566 imageInfo.slicePitch = 0;
567 imageInfo.height = imageInfo.depth = 1;
568 imageInfo.type = CL_MEM_OBJECT_IMAGE1D_ARRAY;
569 pixelSize = get_pixel_size( imageInfo.format );
570
571 error = clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof( maxWidth ), &maxWidth, NULL );
572 error |= clGetDeviceInfo( device, CL_DEVICE_IMAGE_MAX_ARRAY_SIZE, sizeof( maxArraySize ), &maxArraySize, NULL );
573 error |= clGetDeviceInfo( device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof( maxAllocSize ), &maxAllocSize, NULL );
574 error |= clGetDeviceInfo( device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof( memSize ), &memSize, NULL );
575 test_error( error, "Unable to get max image 2D size from device" );
576
577 if (memSize > (cl_ulong)SIZE_MAX) {
578 memSize = (cl_ulong)SIZE_MAX;
579 }
580
581 // Determine types
582 if( inputType == kInt )
583 readFormat = "i";
584 else if( inputType == kUInt )
585 readFormat = "ui";
586 else // kFloat
587 readFormat = "f";
588
589 if(gtestTypesToRun & kWriteTests)
590 {
591 KernelSourcePattern = write1DArrayKernelSourcePattern;
592 }
593 else
594 {
595 KernelSourcePattern = readwrite1DArrayKernelSourcePattern;
596 }
597 // Construct the source
598 // Construct the source
599 sprintf( programSrc,
600 KernelSourcePattern,
601 get_explicit_type_name( inputType ),
602 gTestMipmaps ? ", int lod" : "",
603 gTestMipmaps ? offset1DArrayLodSource : offset1DArraySource,
604 readFormat,
605 gTestMipmaps ? ", lod" :"" );
606
607 ptr = programSrc;
608 error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
609 "sample_kernel");
610 test_error( error, "Unable to create testing kernel" );
611
612 // Run tests
613 if( gTestSmallImages )
614 {
615 for( imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++ )
616 {
617 imageInfo.rowPitch = imageInfo.width * pixelSize;
618 imageInfo.slicePitch = imageInfo.rowPitch;
619 for( imageInfo.arraySize = 2; imageInfo.arraySize < 9; imageInfo.arraySize++ )
620 {
621 if(gTestMipmaps)
622 imageInfo.num_mip_levels = (size_t)random_in_range(2, (compute_max_mip_levels(imageInfo.width, 0, 0)-1), d);
623
624 if( gDebugTrace )
625 log_info( " at size %d,%d\n", (int)imageInfo.width, (int)imageInfo.arraySize );
626 int retCode = test_write_image_1D_array( device, context, queue, kernel, &imageInfo, inputType, d );
627 if( retCode )
628 return retCode;
629 }
630 }
631 }
632 else if( gTestMaxImages )
633 {
634 // Try a specific set of maximum sizes
635 size_t numbeOfSizes;
636 size_t sizes[100][3];
637
638 get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, maxArraySize, maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_ARRAY, imageInfo.format, CL_TRUE);
639
640 for( size_t idx = 0; idx < numbeOfSizes; idx++ )
641 {
642 imageInfo.width = sizes[ idx ][ 0 ];
643 imageInfo.arraySize = sizes[ idx ][ 2 ];
644 imageInfo.rowPitch = imageInfo.width * pixelSize;
645 imageInfo.slicePitch = imageInfo.rowPitch;
646 if(gTestMipmaps)
647 imageInfo.num_mip_levels = (size_t)random_in_range(2, (compute_max_mip_levels(imageInfo.width, 0, 0)-1), d);
648 log_info("Testing %d x %d\n", (int)imageInfo.width, (int)imageInfo.arraySize);
649 int retCode = test_write_image_1D_array( device, context, queue, kernel, &imageInfo, inputType, d );
650 if( retCode )
651 return retCode;
652 }
653 }
654 else if( gTestRounding )
655 {
656 size_t typeRange = 1 << ( get_format_type_size( imageInfo.format ) * 8 );
657 imageInfo.arraySize = typeRange / 256;
658 imageInfo.width = (size_t)( typeRange / (cl_ulong)imageInfo.arraySize );
659
660 imageInfo.rowPitch = imageInfo.width * pixelSize;
661 imageInfo.slicePitch = imageInfo.rowPitch;
662 int retCode = test_write_image_1D_array( device, context, queue, kernel, &imageInfo, inputType, d );
663 if( retCode )
664 return retCode;
665 }
666 else
667 {
668 for( int i = 0; i < NUM_IMAGE_ITERATIONS; i++ )
669 {
670 cl_ulong size;
671 // Loop until we get a size that a) will fit in the max alloc size and b) that an allocation of that
672 // image, the result array, plus offset arrays, will fit in the global ram space
673 do
674 {
675 imageInfo.width = (size_t)random_log_in_range( 16, (int)maxWidth / 32, d );
676 imageInfo.arraySize = (size_t)random_log_in_range( 16, (int)maxArraySize / 32, d );
677
678 if( gTestMipmaps)
679 {
680 imageInfo.num_mip_levels = (size_t)random_in_range(2, (compute_max_mip_levels(imageInfo.width, 0, 0)-1), d);
681 size = (cl_ulong) compute_mipmapped_image_size(imageInfo) * 4;
682 }
683 else
684 {
685 imageInfo.rowPitch = imageInfo.width * pixelSize;
686 if( gEnablePitch )
687 {
688 size_t extraWidth = (int)random_log_in_range( 0, 64, d );
689 imageInfo.rowPitch += extraWidth * pixelSize;
690 }
691 imageInfo.slicePitch = imageInfo.rowPitch;
692
693 size = (size_t)imageInfo.rowPitch * (size_t)imageInfo.arraySize * 4;
694 }
695 } while( size > maxAllocSize || ( size * 3 ) > memSize );
696
697 if( gDebugTrace )
698 log_info( " at size %d,%d (pitch %d) out of %d,%d\n", (int)imageInfo.width, (int)imageInfo.arraySize, (int)imageInfo.rowPitch, (int)maxWidth, (int)maxArraySize );
699
700 int retCode = test_write_image_1D_array( device, context, queue, kernel, &imageInfo, inputType, d );
701 if( retCode )
702 return retCode;
703 }
704 }
705
706 return 0;
707 }
708