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 "harness/compat.h"
17
18 #include <stdio.h>
19 #include <stdlib.h>
20 #include <string.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24
25 #include "procs.h"
26 #include "harness/typeWrappers.h"
27 #include "harness/imageHelpers.h"
28 #include "harness/conversions.h"
29
30
31 static const char *param_kernel[] = {
32 "__kernel void test_fn(read_only image2d_t srcimg, sampler_t sampler, __global float4 *results )\n"
33 "{\n"
34 " int tid_x = get_global_id(0);\n"
35 " int tid_y = get_global_id(1);\n"
36 " results[ tid_y * get_image_width( srcimg ) + tid_x ] = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n"
37 "\n"
38 "}\n" };
39
validate_results(size_t width,size_t height,cl_image_format & format,char * inputData,cl_float * actualResults)40 int validate_results( size_t width, size_t height, cl_image_format &format, char *inputData, cl_float *actualResults )
41 {
42 for( size_t i = 0; i < width * height; i++ )
43 {
44 cl_float expected[ 4 ], tolerance;
45
46 switch( format.image_channel_data_type )
47 {
48 case CL_UNORM_INT8:
49 {
50 cl_uchar *p = (cl_uchar *)inputData;
51 expected[ 0 ] = p[ 0 ] / 255.f;
52 expected[ 1 ] = p[ 1 ] / 255.f;
53 expected[ 2 ] = p[ 2 ] / 255.f;
54 expected[ 3 ] = p[ 3 ] / 255.f;
55 tolerance = 1.f / 255.f;
56 break;
57 }
58 case CL_SNORM_INT8:
59 {
60 cl_char *p = (cl_char *)inputData;
61 expected[ 0 ] = fmaxf( p[ 0 ] / 127.f, -1.f );
62 expected[ 1 ] = fmaxf( p[ 1 ] / 127.f, -1.f );
63 expected[ 2 ] = fmaxf( p[ 2 ] / 127.f, -1.f );
64 expected[ 3 ] = fmaxf( p[ 3 ] / 127.f, -1.f );
65 tolerance = 1.f / 127.f;
66 break;
67 }
68 case CL_UNSIGNED_INT8:
69 {
70 cl_uchar *p = (cl_uchar *)inputData;
71 expected[ 0 ] = p[ 0 ];
72 expected[ 1 ] = p[ 1 ];
73 expected[ 2 ] = p[ 2 ];
74 expected[ 3 ] = p[ 3 ];
75 tolerance = 1.f / 127.f;
76 break;
77 }
78 case CL_SIGNED_INT8:
79 {
80 cl_short *p = (cl_short *)inputData;
81 expected[ 0 ] = p[ 0 ];
82 expected[ 1 ] = p[ 1 ];
83 expected[ 2 ] = p[ 2 ];
84 expected[ 3 ] = p[ 3 ];
85 tolerance = 1.f / 127.f;
86 break;
87 }
88 case CL_UNORM_INT16:
89 {
90 cl_ushort *p = (cl_ushort *)inputData;
91 expected[ 0 ] = p[ 0 ] / 65535.f;
92 expected[ 1 ] = p[ 1 ] / 65535.f;
93 expected[ 2 ] = p[ 2 ] / 65535.f;
94 expected[ 3 ] = p[ 3 ] / 65535.f;
95 tolerance = 1.f / 65535.f;
96 break;
97 }
98 case CL_UNSIGNED_INT32:
99 {
100 cl_uint *p = (cl_uint *)inputData;
101 expected[ 0 ] = p[ 0 ];
102 expected[ 1 ] = p[ 1 ];
103 expected[ 2 ] = p[ 2 ];
104 expected[ 3 ] = p[ 3 ];
105 tolerance = 0.0001f;
106 break;
107 }
108 case CL_FLOAT:
109 {
110 cl_float *p = (cl_float *)inputData;
111 expected[ 0 ] = p[ 0 ];
112 expected[ 1 ] = p[ 1 ];
113 expected[ 2 ] = p[ 2 ];
114 expected[ 3 ] = p[ 3 ];
115 tolerance = 0.0001f;
116 break;
117 }
118 default:
119 // Should never get here
120 log_error("Unhandled channel data type\n");
121 return -1;
122 }
123
124 if( format.image_channel_order == CL_BGRA )
125 {
126 cl_float tmp = expected[ 0 ];
127 expected[ 0 ] = expected[ 2 ];
128 expected[ 2 ] = tmp;
129 }
130
131 // Within an error tolerance, make sure the results match
132 cl_float error1 = fabsf( expected[ 0 ] - actualResults[ 0 ] );
133 cl_float error2 = fabsf( expected[ 1 ] - actualResults[ 1 ] );
134 cl_float error3 = fabsf( expected[ 2 ] - actualResults[ 2 ] );
135 cl_float error4 = fabsf( expected[ 3 ] - actualResults[ 3 ] );
136
137 if( error1 > tolerance || error2 > tolerance || error3 > tolerance || error4 > tolerance )
138 {
139 log_error( "ERROR: Sample %d did not validate against expected results for %d x %d %s:%s image\n", (int)i, (int)width, (int)height,
140 GetChannelOrderName( format.image_channel_order ), GetChannelTypeName( format.image_channel_data_type ) );
141 log_error( " Expected: %f %f %f %f\n", (float)expected[ 0 ], (float)expected[ 1 ], (float)expected[ 2 ], (float)expected[ 3 ] );
142 log_error( " Actual: %f %f %f %f\n", (float)actualResults[ 0 ], (float)actualResults[ 1 ], (float)actualResults[ 2 ], (float)actualResults[ 3 ] );
143
144 // Check real quick a special case error here
145 cl_float error1 = fabsf( expected[ 3 ] - actualResults[ 0 ] );
146 cl_float error2 = fabsf( expected[ 2 ] - actualResults[ 1 ] );
147 cl_float error3 = fabsf( expected[ 1 ] - actualResults[ 2 ] );
148 cl_float error4 = fabsf( expected[ 0 ] - actualResults[ 3 ] );
149 if( error1 <= tolerance && error2 <= tolerance && error3 <= tolerance && error4 <= tolerance )
150 {
151 log_error( "\t(Kernel did not respect change in channel order)\n" );
152 }
153 return -1;
154 }
155
156 // Increment and go
157 actualResults += 4;
158 inputData += get_format_type_size( &format ) * 4;
159 }
160
161 return 0;
162 }
163
test_image_param(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)164 int test_image_param(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
165 {
166 size_t sizes[] = { 64, 100, 128, 250, 512 };
167 cl_image_format formats[] = { { CL_RGBA, CL_UNORM_INT8 }, { CL_RGBA, CL_UNORM_INT16 }, { CL_RGBA, CL_FLOAT }, { CL_BGRA, CL_UNORM_INT8 } };
168 cl_image_format *supported_formats;
169 ExplicitType types[] = { kUChar, kUShort, kFloat, kUChar };
170 int error;
171 size_t i, j, idx;
172 size_t threads[ 2 ];
173 MTdata d;
174 int supportsBGRA = 0;
175 cl_uint numSupportedFormats = 0;
176
177 const size_t numSizes = sizeof( sizes ) / sizeof( sizes[ 0 ] );
178 const size_t numFormats = sizeof( formats ) / sizeof( formats[ 0 ] );
179 const size_t numAttempts = numSizes * numFormats;
180
181
182 clProgramWrapper program;
183 clKernelWrapper kernel;
184 clMemWrapper streams[ numAttempts ][ 2 ];
185 BufferOwningPtr<char> inputs[ numAttempts ];
186
187 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
188
189 if(gIsEmbedded)
190 {
191 /* Get the supported image formats to see if BGRA is supported */
192 clGetSupportedImageFormats (context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &numSupportedFormats);
193 supported_formats = (cl_image_format *) malloc(sizeof(cl_image_format) * numSupportedFormats);
194 clGetSupportedImageFormats (context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, numFormats, supported_formats, NULL);
195
196 for(i = 0; i < numSupportedFormats; i++)
197 {
198 if(supported_formats[i].image_channel_order == CL_BGRA)
199 {
200 supportsBGRA = 1;
201 break;
202 }
203 }
204 }
205 else
206 {
207 supportsBGRA = 1;
208 }
209
210 d = init_genrand( gRandomSeed );
211 for( i = 0, idx = 0; i < numSizes; i++ )
212 {
213 for( j = 0; j < numFormats; j++, idx++ )
214 {
215 if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
216 continue;
217
218 // For each attempt, we create a pair: an input image, whose parameters keep changing, and an output buffer
219 // that we can read values from. The output buffer will remain consistent to ensure that any changes we
220 // witness are due to the image changes
221 inputs[ idx ].reset(create_random_data( types[ j ], d, sizes[ i ] * sizes[ i ] * 4 ));
222
223 streams[ idx ][ 0 ] = create_image_2d( context, CL_MEM_COPY_HOST_PTR, &formats[ j ], sizes[ i ], sizes[ i ], 0, inputs[ idx ], &error );
224 {
225 char err_str[256];
226 sprintf(err_str, "Unable to create input image for format %s order %s" ,
227 GetChannelOrderName( formats[j].image_channel_order ),
228 GetChannelTypeName( formats[j].image_channel_data_type ));
229 test_error( error, err_str);
230 }
231
232 streams[ idx ][ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), NULL, &error );
233 test_error( error, "Unable to create output buffer" );
234 }
235 }
236 free_mtdata(d); d = NULL;
237
238 // Create a single kernel to use for all the tests
239 error = create_single_kernel_helper( context, &program, &kernel, 1, param_kernel, "test_fn" );
240 test_error( error, "Unable to create testing kernel" );
241
242 // Also create a sampler to use for all the runs
243 clSamplerWrapper sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &error );
244 test_error( error, "clCreateSampler failed" );
245
246 // Set up the arguments for each and queue
247 for( i = 0, idx = 0; i < numSizes; i++ )
248 {
249 for( j = 0; j < numFormats; j++, idx++ )
250 {
251 if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
252 continue;
253
254 error = clSetKernelArg( kernel, 0, sizeof( streams[ idx ][ 0 ] ), &streams[ idx ][ 0 ] );
255 error |= clSetKernelArg( kernel, 1, sizeof( sampler ), &sampler );
256 error |= clSetKernelArg( kernel, 2, sizeof( streams[ idx ][ 1 ] ), &streams[ idx ][ 1 ]);
257 test_error( error, "Unable to set kernel arguments" );
258
259 threads[ 0 ] = threads[ 1 ] = (size_t)sizes[ i ];
260
261 error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL );
262 test_error( error, "clEnqueueNDRangeKernel failed" );
263 }
264 }
265
266 // Now go through each combo and validate the results
267 for( i = 0, idx = 0; i < numSizes; i++ )
268 {
269 for( j = 0; j < numFormats; j++, idx++ )
270 {
271 if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
272 continue;
273
274 BufferOwningPtr<cl_float> output(malloc(sizeof(cl_float) * sizes[ i ] * sizes[ i ] * 4 ));
275
276 error = clEnqueueReadBuffer( queue, streams[ idx ][ 1 ], CL_TRUE, 0, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), output, 0, NULL, NULL );
277 test_error( error, "Unable to read results" );
278
279 error = validate_results( sizes[ i ], sizes[ i ], formats[ j ], inputs[ idx ], output );
280 if( error )
281 return -1;
282 }
283 }
284
285 return 0;
286 }
287