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 break;
121 }
122
123 if( format.image_channel_order == CL_BGRA )
124 {
125 cl_float tmp = expected[ 0 ];
126 expected[ 0 ] = expected[ 2 ];
127 expected[ 2 ] = tmp;
128 }
129
130 // Within an error tolerance, make sure the results match
131 cl_float error1 = fabsf( expected[ 0 ] - actualResults[ 0 ] );
132 cl_float error2 = fabsf( expected[ 1 ] - actualResults[ 1 ] );
133 cl_float error3 = fabsf( expected[ 2 ] - actualResults[ 2 ] );
134 cl_float error4 = fabsf( expected[ 3 ] - actualResults[ 3 ] );
135
136 if( error1 > tolerance || error2 > tolerance || error3 > tolerance || error4 > tolerance )
137 {
138 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,
139 GetChannelOrderName( format.image_channel_order ), GetChannelTypeName( format.image_channel_data_type ) );
140 log_error( " Expected: %f %f %f %f\n", (float)expected[ 0 ], (float)expected[ 1 ], (float)expected[ 2 ], (float)expected[ 3 ] );
141 log_error( " Actual: %f %f %f %f\n", (float)actualResults[ 0 ], (float)actualResults[ 1 ], (float)actualResults[ 2 ], (float)actualResults[ 3 ] );
142
143 // Check real quick a special case error here
144 cl_float error1 = fabsf( expected[ 3 ] - actualResults[ 0 ] );
145 cl_float error2 = fabsf( expected[ 2 ] - actualResults[ 1 ] );
146 cl_float error3 = fabsf( expected[ 1 ] - actualResults[ 2 ] );
147 cl_float error4 = fabsf( expected[ 0 ] - actualResults[ 3 ] );
148 if( error1 <= tolerance && error2 <= tolerance && error3 <= tolerance && error4 <= tolerance )
149 {
150 log_error( "\t(Kernel did not respect change in channel order)\n" );
151 }
152 return -1;
153 }
154
155 // Increment and go
156 actualResults += 4;
157 inputData += get_format_type_size( &format ) * 4;
158 }
159
160 return 0;
161 }
162
test_image_param(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)163 int test_image_param(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
164 {
165 size_t sizes[] = { 64, 100, 128, 250, 512 };
166 cl_image_format formats[] = { { CL_RGBA, CL_UNORM_INT8 }, { CL_RGBA, CL_UNORM_INT16 }, { CL_RGBA, CL_FLOAT }, { CL_BGRA, CL_UNORM_INT8 } };
167 cl_image_format *supported_formats;
168 ExplicitType types[] = { kUChar, kUShort, kFloat, kUChar };
169 int error;
170 size_t i, j, idx;
171 size_t threads[ 2 ];
172 MTdata d;
173 int supportsBGRA = 0;
174 cl_uint numSupportedFormats = 0;
175
176 const size_t numSizes = sizeof( sizes ) / sizeof( sizes[ 0 ] );
177 const size_t numFormats = sizeof( formats ) / sizeof( formats[ 0 ] );
178 const size_t numAttempts = numSizes * numFormats;
179
180
181 clProgramWrapper program;
182 clKernelWrapper kernel;
183 clMemWrapper streams[ numAttempts ][ 2 ];
184 BufferOwningPtr<char> inputs[ numAttempts ];
185
186 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
187
188 if(gIsEmbedded)
189 {
190 /* Get the supported image formats to see if BGRA is supported */
191 clGetSupportedImageFormats (context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &numSupportedFormats);
192 supported_formats = (cl_image_format *) malloc(sizeof(cl_image_format) * numSupportedFormats);
193 clGetSupportedImageFormats (context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, numFormats, supported_formats, NULL);
194
195 for(i = 0; i < numSupportedFormats; i++)
196 {
197 if(supported_formats[i].image_channel_order == CL_BGRA)
198 {
199 supportsBGRA = 1;
200 break;
201 }
202 }
203 }
204 else
205 {
206 supportsBGRA = 1;
207 }
208
209 d = init_genrand( gRandomSeed );
210 for( i = 0, idx = 0; i < numSizes; i++ )
211 {
212 for( j = 0; j < numFormats; j++, idx++ )
213 {
214 if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
215 continue;
216
217 // For each attempt, we create a pair: an input image, whose parameters keep changing, and an output buffer
218 // that we can read values from. The output buffer will remain consistent to ensure that any changes we
219 // witness are due to the image changes
220 inputs[ idx ].reset(create_random_data( types[ j ], d, sizes[ i ] * sizes[ i ] * 4 ));
221
222 streams[ idx ][ 0 ] = create_image_2d( context, CL_MEM_COPY_HOST_PTR, &formats[ j ], sizes[ i ], sizes[ i ], 0, inputs[ idx ], &error );
223 {
224 char err_str[256];
225 sprintf(err_str, "Unable to create input image for format %s order %s" ,
226 GetChannelOrderName( formats[j].image_channel_order ),
227 GetChannelTypeName( formats[j].image_channel_data_type ));
228 test_error( error, err_str);
229 }
230
231 streams[ idx ][ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), NULL, &error );
232 test_error( error, "Unable to create output buffer" );
233 }
234 }
235 free_mtdata(d); d = NULL;
236
237 // Create a single kernel to use for all the tests
238 error = create_single_kernel_helper( context, &program, &kernel, 1, param_kernel, "test_fn" );
239 test_error( error, "Unable to create testing kernel" );
240
241 // Also create a sampler to use for all the runs
242 clSamplerWrapper sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &error );
243 test_error( error, "clCreateSampler failed" );
244
245 // Set up the arguments for each and queue
246 for( i = 0, idx = 0; i < numSizes; i++ )
247 {
248 for( j = 0; j < numFormats; j++, idx++ )
249 {
250 if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
251 continue;
252
253 error = clSetKernelArg( kernel, 0, sizeof( streams[ idx ][ 0 ] ), &streams[ idx ][ 0 ] );
254 error |= clSetKernelArg( kernel, 1, sizeof( sampler ), &sampler );
255 error |= clSetKernelArg( kernel, 2, sizeof( streams[ idx ][ 1 ] ), &streams[ idx ][ 1 ]);
256 test_error( error, "Unable to set kernel arguments" );
257
258 threads[ 0 ] = threads[ 1 ] = (size_t)sizes[ i ];
259
260 error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL );
261 test_error( error, "clEnqueueNDRangeKernel failed" );
262 }
263 }
264
265 // Now go through each combo and validate the results
266 for( i = 0, idx = 0; i < numSizes; i++ )
267 {
268 for( j = 0; j < numFormats; j++, idx++ )
269 {
270 if(formats[j].image_channel_order == CL_BGRA && !supportsBGRA)
271 continue;
272
273 BufferOwningPtr<cl_float> output(malloc(sizeof(cl_float) * sizes[ i ] * sizes[ i ] * 4 ));
274
275 error = clEnqueueReadBuffer( queue, streams[ idx ][ 1 ], CL_TRUE, 0, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), output, 0, NULL, NULL );
276 test_error( error, "Unable to read results" );
277
278 error = validate_results( sizes[ i ], sizes[ i ], formats[ j ], inputs[ idx ], output );
279 if( error )
280 return -1;
281 }
282 }
283
284 return 0;
285 }
286