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