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