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 "harness/conversions.h"
18
19 static const int vector_sizes[] = {1, 2, 3, 4, 8, 16};
20 #define NUM_VECTOR_SIZES 6
21
22 const char *permute_2_param_kernel_pattern =
23 "__kernel void test_upsample(__global %s *sourceA, __global %s *sourceB, __global %s *destValues)\n"
24 "{\n"
25 " int tid = get_global_id(0);\n"
26 " destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n"
27 "\n"
28 "}\n";
29
30
31 const char *permute_2_param_kernel_pattern_v3srcdst =
32 "__kernel void test_upsample(__global %s *sourceA, __global %s *sourceB, __global %s *destValues)\n"
33 "{\n"
34 " int tid = get_global_id(0);\n"
35 " vstore3( %s( vload3(tid,sourceA), vload3(tid, sourceB) ), tid, destValues);\n"
36 "\n"
37 "}\n";
38
test_upsample_2_param_fn(cl_command_queue queue,cl_context context,const char * fnName,ExplicitType sourceAType,ExplicitType sourceBType,ExplicitType outType,size_t sourceAVecSize,size_t sourceBVecSize,size_t outVecSize,size_t count,void * sourceA,void * sourceB,void * expectedResults)39 int test_upsample_2_param_fn(cl_command_queue queue, cl_context context, const char *fnName, ExplicitType sourceAType, ExplicitType sourceBType, ExplicitType outType,
40 size_t sourceAVecSize, size_t sourceBVecSize, size_t outVecSize, size_t count,
41 void *sourceA, void *sourceB, void *expectedResults )
42 {
43 cl_program program;
44 cl_kernel kernel;
45 int error, retCode = 0;
46 cl_mem streams[3];
47 void *outData;
48 size_t threadSize, groupSize, i;
49 unsigned char *expectedPtr, *outPtr;
50 size_t sourceATypeSize, sourceBTypeSize, outTypeSize, outStride;
51 char programSource[ 10240 ], aType[ 64 ], bType[ 64 ], tType[ 64 ];
52 const char *progPtr;
53
54
55 sourceATypeSize = get_explicit_type_size( sourceAType );
56 sourceBTypeSize = get_explicit_type_size( sourceBType );
57 outTypeSize = get_explicit_type_size( outType );
58
59 outStride = outTypeSize * outVecSize;
60 outData = malloc( outStride * count );
61
62 /* Construct the program */
63 strcpy( aType, get_explicit_type_name( sourceAType ) );
64 strcpy( bType, get_explicit_type_name( sourceBType ) );
65 strcpy( tType, get_explicit_type_name( outType ) );
66 if( sourceAVecSize > 1 && sourceAVecSize != 3)
67 sprintf( aType + strlen( aType ), "%d", (int)sourceAVecSize );
68 if( sourceBVecSize > 1 && sourceBVecSize != 3)
69 sprintf( bType + strlen( bType ), "%d", (int)sourceBVecSize );
70 if( outVecSize > 1 && outVecSize != 3)
71 sprintf( tType + strlen( tType ), "%d", (int)outVecSize );
72
73 if(sourceAVecSize == 3 && sourceBVecSize == 3 && outVecSize == 3)
74 {
75 // permute_2_param_kernel_pattern_v3srcdst
76 sprintf( programSource, permute_2_param_kernel_pattern_v3srcdst, aType, bType, tType, fnName );
77 }
78 else if(sourceAVecSize != 3 && sourceBVecSize != 3 && outVecSize != 3)
79 {
80 sprintf( programSource, permute_2_param_kernel_pattern, aType, bType, tType, fnName );
81 } else {
82 vlog_error("Not implemented for %d,%d -> %d\n",
83 (int)sourceAVecSize, (int)sourceBVecSize, (int)outVecSize);
84 return -1;
85 }
86
87 progPtr = (const char *)programSource;
88 if( create_single_kernel_helper( context, &program, &kernel, 1, &progPtr, "test_upsample" ) )
89 {
90 free( outData );
91 return -1;
92 }
93
94 /* Set up parameters */
95 streams[0] =
96 clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
97 sourceATypeSize * sourceAVecSize * count, sourceA, NULL);
98 if (!streams[0])
99 {
100 log_error("ERROR: Creating input array A failed!\n");
101 return -1;
102 }
103 streams[1] =
104 clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
105 sourceBTypeSize * sourceBVecSize * count, sourceB, NULL);
106 if (!streams[1])
107 {
108 log_error("ERROR: Creating input array B failed!\n");
109 return -1;
110 }
111 streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, outStride * count,
112 NULL, NULL);
113 if (!streams[2])
114 {
115 log_error("ERROR: Creating output array failed!\n");
116 return -1;
117 }
118
119 /* Set the arguments */
120 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0] );
121 test_error( error, "Unable to set kernel arguments" );
122 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1] );
123 test_error( error, "Unable to set kernel arguments" );
124 error = clSetKernelArg(kernel, 2, sizeof( streams[2] ), &streams[2] );
125 test_error( error, "Unable to set kernel arguments" );
126
127 /* Run the kernel */
128 threadSize = count;
129
130 error = get_max_common_work_group_size( context, kernel, threadSize, &groupSize );
131 test_error( error, "Unable to get work group size to use" );
132
133 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, &threadSize, &groupSize, 0, NULL, NULL );
134 test_error( error, "Unable to execute test kernel" );
135
136 /* Now verify the results. Each value should have been duplicated four times, and we should be able to just
137 do a memcpy instead of relying on the actual type of data */
138 error = clEnqueueReadBuffer( queue, streams[2], CL_TRUE, 0, outStride * count, outData, 0, NULL, NULL );
139 test_error( error, "Unable to read output values!" );
140
141 expectedPtr = (unsigned char *)expectedResults;
142 outPtr = (unsigned char *)outData;
143
144 for( i = 0; i < count; i++ )
145 {
146 if( memcmp( outPtr, expectedPtr, outTypeSize * outVecSize ) != 0 )
147 {
148 log_error( "ERROR: Output value %d does not validate!\n", (int)i );
149 retCode = -1;
150 break;
151 }
152 expectedPtr += outTypeSize * outVecSize;
153 outPtr += outStride;
154 }
155
156 clReleaseMemObject( streams[0] );
157 clReleaseMemObject( streams[1] );
158 clReleaseMemObject( streams[2] );
159 clReleaseKernel( kernel );
160 clReleaseProgram( program );
161 free( outData );
162
163 return retCode;
164 }
165
create_upsample_data(ExplicitType type,void * sourceA,void * sourceB,size_t count)166 void * create_upsample_data( ExplicitType type, void *sourceA, void *sourceB, size_t count )
167 {
168 void *outData;
169 size_t i, tSize;
170
171 tSize = get_explicit_type_size( type );
172 outData = malloc( tSize * count * 2 );
173
174 switch( tSize )
175 {
176 case 1:
177 {
178 const cl_uchar *aPtr = (const cl_uchar *) sourceA;
179 const cl_uchar *bPtr = (const cl_uchar *) sourceB;
180 cl_ushort *dPtr = (cl_ushort*) outData;
181 for( i = 0; i < count; i++ )
182 {
183 cl_ushort u = *bPtr++;
184 u |= ((cl_ushort) *aPtr++) << 8;
185 *dPtr++ = u;
186 }
187 }
188 break;
189 case 2:
190 {
191 const cl_ushort *aPtr = (const cl_ushort *) sourceA;
192 const cl_ushort *bPtr = (const cl_ushort *) sourceB;
193 cl_uint *dPtr = (cl_uint*) outData;
194 for( i = 0; i < count; i++ )
195 {
196 cl_uint u = *bPtr++;
197 u |= ((cl_uint) *aPtr++) << 16;
198 *dPtr++ = u;
199 }
200 }
201 break;
202 case 4:
203 {
204 const cl_uint *aPtr = (const cl_uint *) sourceA;
205 const cl_uint *bPtr = (const cl_uint *) sourceB;
206 cl_ulong *dPtr = (cl_ulong*) outData;
207 for( i = 0; i < count; i++ )
208 {
209 cl_ulong u = *bPtr++;
210 u |= ((cl_ulong) *aPtr++) << 32;
211 *dPtr++ = u;
212 }
213 }
214 break;
215 default:
216 log_error( "ERROR: unknown type size: %ld\n", tSize );
217 return NULL;
218 }
219
220 return outData;
221 }
222
test_integer_upsample(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)223 int test_integer_upsample(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
224 {
225 ExplicitType typesToTest[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kNumExplicitTypes };
226 ExplicitType baseTypes[] = { kUChar, kUChar, kUShort, kUShort, kUInt, kUInt, kNumExplicitTypes };
227 ExplicitType outTypes[] = { kShort, kUShort, kInt, kUInt, kLong, kULong, kNumExplicitTypes };
228 int i, err = 0;
229 int sizeIndex;
230 size_t size;
231 void *sourceA, *sourceB, *expected;
232 RandomSeed seed(gRandomSeed );
233
234 for( i = 0; typesToTest[ i ] != kNumExplicitTypes; i++ )
235 {
236 if ((outTypes[i] == kLong || outTypes[i] == kULong) && !gHasLong)
237 {
238 log_info( "Longs unsupported on this device. Skipping...\n");
239 continue;
240 }
241
242 for( sizeIndex = 0; sizeIndex < NUM_VECTOR_SIZES; sizeIndex++)
243 {
244 size = (size_t)vector_sizes[sizeIndex];
245 log_info("running upsample test for %s %s vector size %d\n", get_explicit_type_name(typesToTest[i]), get_explicit_type_name(baseTypes[i]), (int)size);
246 sourceA = create_random_data( typesToTest[ i ], seed, 256 );
247 sourceB = create_random_data( baseTypes[ i ], seed, 256 );
248 expected = create_upsample_data( typesToTest[ i ], sourceA, sourceB, 256 );
249
250 if( test_upsample_2_param_fn( queue, context, "upsample",
251 typesToTest[ i ], baseTypes[ i ],
252 outTypes[ i ],
253 size, size, size,
254 256 / size,
255 sourceA, sourceB, expected ) != 0 )
256 {
257 log_error( "TEST FAILED: %s for %s%d\n", "upsample", get_explicit_type_name( typesToTest[ i ] ), (int)size );
258 err = -1;
259 }
260 free( sourceA );
261 free( sourceB );
262 free( expected );
263 }
264 }
265 return err;
266 }
267
268
269