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 // This test is designed to stress passing multiple vector parameters to kernels and verifying access between them all
20 
21 const char *multi_arg_kernel_source_pattern =
22 "__kernel void sample_test(__global %s *src1, __global %s *src2, __global %s *src3, __global %s *dst1, __global %s *dst2, __global %s *dst3 )\n"
23 "{\n"
24 "    int tid = get_global_id(0);\n"
25 "    dst1[tid] = src1[tid];\n"
26 "    dst2[tid] = src2[tid];\n"
27 "    dst3[tid] = src3[tid];\n"
28 "}\n";
29 
30 #define MAX_ERROR_TOLERANCE 0.0005f
31 
test_multi_arg_set(cl_device_id device,cl_context context,cl_command_queue queue,ExplicitType vec1Type,int vec1Size,ExplicitType vec2Type,int vec2Size,ExplicitType vec3Type,int vec3Size,MTdata d)32 int test_multi_arg_set(cl_device_id device, cl_context context, cl_command_queue queue,
33                        ExplicitType vec1Type, int vec1Size,
34                        ExplicitType vec2Type, int vec2Size,
35                        ExplicitType vec3Type, int vec3Size, MTdata d)
36 {
37     clProgramWrapper program;
38     clKernelWrapper kernel;
39     int error, i, j;
40     clMemWrapper streams[ 6 ];
41     size_t threads[1], localThreads[1];
42     char programSrc[ 10248 ], vec1Name[ 64 ], vec2Name[ 64 ], vec3Name[ 64 ];
43     char sizeNames[][ 4 ] = { "", "2", "3", "4", "", "", "", "8" };
44     const char *ptr;
45     void *initData[3], *resultData[3];
46 
47 
48     // Create the program source
49     sprintf( vec1Name, "%s%s", get_explicit_type_name( vec1Type ), sizeNames[ vec1Size - 1 ] );
50     sprintf( vec2Name, "%s%s", get_explicit_type_name( vec2Type ), sizeNames[ vec2Size - 1 ] );
51     sprintf( vec3Name, "%s%s", get_explicit_type_name( vec3Type ), sizeNames[ vec3Size - 1 ] );
52 
53     sprintf( programSrc, multi_arg_kernel_source_pattern,
54             vec1Name, vec2Name, vec3Name, vec1Name, vec2Name, vec3Name,
55             vec1Size, vec1Size, vec2Size, vec2Size, vec3Size, vec3Size );
56     ptr = programSrc;
57 
58     // Create our testing kernel
59     error = create_single_kernel_helper( context, &program, &kernel, 1, &ptr, "sample_test" );
60     test_error( error, "Unable to create testing kernel" );
61 
62     // Get thread dimensions
63     threads[0] = 1024;
64     error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
65     test_error( error, "Unable to get work group size for kernel" );
66 
67     // Create input streams
68     initData[ 0 ] = create_random_data( vec1Type, d, (unsigned int)threads[ 0 ] * vec1Size );
69     streams[0] =
70         clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
71                        get_explicit_type_size(vec1Type) * threads[0] * vec1Size,
72                        initData[0], &error);
73     test_error( error, "Unable to create testing stream" );
74 
75     initData[ 1 ] = create_random_data( vec2Type, d, (unsigned int)threads[ 0 ] * vec2Size );
76     streams[1] =
77         clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
78                        get_explicit_type_size(vec2Type) * threads[0] * vec2Size,
79                        initData[1], &error);
80     test_error( error, "Unable to create testing stream" );
81 
82     initData[ 2 ] = create_random_data( vec3Type, d, (unsigned int)threads[ 0 ] * vec3Size );
83     streams[2] =
84         clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
85                        get_explicit_type_size(vec3Type) * threads[0] * vec3Size,
86                        initData[2], &error);
87     test_error( error, "Unable to create testing stream" );
88 
89     streams[3] = clCreateBuffer(
90         context, CL_MEM_READ_WRITE,
91         get_explicit_type_size(vec1Type) * threads[0] * vec1Size, NULL, &error);
92     test_error( error, "Unable to create testing stream" );
93 
94     streams[4] = clCreateBuffer(
95         context, CL_MEM_READ_WRITE,
96         get_explicit_type_size(vec2Type) * threads[0] * vec2Size, NULL, &error);
97     test_error( error, "Unable to create testing stream" );
98 
99     streams[5] = clCreateBuffer(
100         context, CL_MEM_READ_WRITE,
101         get_explicit_type_size(vec3Type) * threads[0] * vec3Size, NULL, &error);
102     test_error( error, "Unable to create testing stream" );
103 
104     // Set the arguments
105     error = 0;
106     for( i = 0; i < 6; i++ )
107         error |= clSetKernelArg( kernel, i, sizeof( cl_mem ), &streams[ i ] );
108     test_error( error, "Unable to set arguments for kernel" );
109 
110     // Execute!
111     error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
112     test_error( error, "Unable to execute kernel" );
113 
114     // Read results
115     resultData[0] = malloc( get_explicit_type_size( vec1Type ) * vec1Size * threads[0] );
116     resultData[1] = malloc( get_explicit_type_size( vec2Type ) * vec2Size * threads[0] );
117     resultData[2] = malloc( get_explicit_type_size( vec3Type ) * vec3Size * threads[0] );
118     error = clEnqueueReadBuffer( queue, streams[ 3 ], CL_TRUE, 0, get_explicit_type_size( vec1Type ) * vec1Size * threads[ 0 ], resultData[0], 0, NULL, NULL );
119     error |= clEnqueueReadBuffer( queue, streams[ 4 ], CL_TRUE, 0, get_explicit_type_size( vec2Type ) * vec2Size * threads[ 0 ], resultData[1], 0, NULL, NULL );
120     error |= clEnqueueReadBuffer( queue, streams[ 5 ], CL_TRUE, 0, get_explicit_type_size( vec3Type ) * vec3Size * threads[ 0 ], resultData[2], 0, NULL, NULL );
121     test_error( error, "Unable to read result stream" );
122 
123     // Verify
124     char *ptr1 = (char *)initData[ 0 ], *ptr2 = (char *)resultData[ 0 ];
125     size_t span = get_explicit_type_size( vec1Type );
126     for( i = 0; i < (int)threads[0]; i++ )
127     {
128         for( j = 0; j < vec1Size; j++ )
129         {
130             if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
131             {
132                 log_error( "ERROR: Value did not validate for component %d of item %d of stream 0!\n", j, i );
133                 free( initData[ 0 ] );
134                 free( initData[ 1 ] );
135                 free( initData[ 2 ] );
136                 free( resultData[ 0 ] );
137                 free( resultData[ 1 ] );
138                 free( resultData[ 2 ] );
139                 return -1;
140             }
141         }
142         ptr1 += span * vec1Size;
143         ptr2 += span * vec1Size;
144     }
145 
146     ptr1 = (char *)initData[ 1 ];
147     ptr2 = (char *)resultData[ 1 ];
148     span = get_explicit_type_size( vec2Type );
149     for( i = 0; i < (int)threads[0]; i++ )
150     {
151         for( j = 0; j < vec2Size; j++ )
152         {
153             if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
154             {
155                 log_error( "ERROR: Value did not validate for component %d of item %d of stream 1!\n", j, i );
156                 free( initData[ 0 ] );
157                 free( initData[ 1 ] );
158                 free( initData[ 2 ] );
159                 free( resultData[ 0 ] );
160                 free( resultData[ 1 ] );
161                 free( resultData[ 2 ] );
162                 return -1;
163             }
164         }
165         ptr1 += span * vec2Size;
166         ptr2 += span * vec2Size;
167     }
168 
169     ptr1 = (char *)initData[ 2 ];
170     ptr2 = (char *)resultData[ 2 ];
171     span = get_explicit_type_size( vec3Type );
172     for( i = 0; i < (int)threads[0]; i++ )
173     {
174         for( j = 0; j < vec3Size; j++ )
175         {
176             if( memcmp( ptr1 + span * j , ptr2 + span * j, span ) != 0 )
177             {
178                 log_error( "ERROR: Value did not validate for component %d of item %d of stream 2!\n", j, i );
179                 free( initData[ 0 ] );
180                 free( initData[ 1 ] );
181                 free( initData[ 2 ] );
182                 free( resultData[ 0 ] );
183                 free( resultData[ 1 ] );
184                 free( resultData[ 2 ] );
185                 return -1;
186             }
187         }
188         ptr1 += span * vec3Size;
189         ptr2 += span * vec3Size;
190     }
191 
192     // If we got here, everything verified successfully
193     free( initData[ 0 ] );
194     free( initData[ 1 ] );
195     free( initData[ 2 ] );
196     free( resultData[ 0 ] );
197     free( resultData[ 1 ] );
198     free( resultData[ 2 ] );
199 
200     return 0;
201 }
202 
test_kernel_arg_multi_setup_exhaustive(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)203 int test_kernel_arg_multi_setup_exhaustive(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
204 {
205     // Loop through every combination of input and output types
206     ExplicitType types[] = { kChar, kShort, kInt, kFloat, kNumExplicitTypes };
207     int type1, type2, type3;
208     int size1, size2, size3;
209     RandomSeed seed( gRandomSeed );
210 
211     log_info( "\n" ); // for formatting
212 
213     for( type1 = 0; types[ type1 ] != kNumExplicitTypes; type1++ )
214     {
215         for( type2 = 0; types[ type2 ] != kNumExplicitTypes; type2++ )
216         {
217             for( type3 = 0; types[ type3 ] != kNumExplicitTypes; type3++ )
218             {
219                 log_info( "\n\ttesting %s, %s, %s...", get_explicit_type_name( types[ type1 ] ), get_explicit_type_name( types[ type2 ] ), get_explicit_type_name( types[ type3 ] ) );
220 
221                 // Loop through every combination of vector size
222                 for( size1 = 2; size1 <= 8; size1 <<= 1 )
223                 {
224                     for( size2 = 2; size2 <= 8; size2 <<= 1 )
225                     {
226                         for( size3 = 2; size3 <= 8; size3 <<= 1 )
227                         {
228                             log_info(".");
229                             fflush( stdout);
230                             if( test_multi_arg_set( device, context, queue,
231                                                    types[ type1 ], size1,
232                                                    types[ type2 ], size2,
233                                                    types[ type3 ], size3, seed ) )
234                                 return -1;
235                         }
236                     }
237                 }
238             }
239         }
240     }
241     log_info( "\n" );
242     return 0;
243 }
244 
test_kernel_arg_multi_setup_random(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)245 int test_kernel_arg_multi_setup_random(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
246 {
247     // Loop through a selection of combinations
248     ExplicitType types[] = { kChar, kShort, kInt, kFloat, kNumExplicitTypes };
249     int type1, type2, type3;
250     int size1, size2, size3;
251     RandomSeed seed( gRandomSeed );
252 
253     num_elements = 3*3*3*4;
254     log_info( "Testing %d random configurations\n", num_elements );
255 
256     // Loop through every combination of vector size
257     for( size1 = 2; size1 <= 8; size1 <<= 1 )
258     {
259         for( size2 = 2; size2 <= 8; size2 <<= 1 )
260         {
261             for( size3 = 2; size3 <= 8; size3 <<= 1 )
262             {
263                 // Loop through 4 type combinations for each size combination
264                 int n;
265                 for (n=0; n<4; n++) {
266                     type1 = (int)get_random_float(0,4, seed);
267                     type2 = (int)get_random_float(0,4, seed);
268                     type3 = (int)get_random_float(0,4, seed);
269 
270 
271                     log_info( "\ttesting %s%d, %s%d, %s%d...\n",
272                              get_explicit_type_name( types[ type1 ] ), size1,
273                              get_explicit_type_name( types[ type2 ] ), size2,
274                              get_explicit_type_name( types[ type3 ] ), size3 );
275 
276                     if( test_multi_arg_set( device, context, queue,
277                                            types[ type1 ], size1,
278                                            types[ type2 ], size2,
279                                            types[ type3 ], size3, seed ) )
280                         return -1;
281                 }
282             }
283         }
284     }
285     return 0;
286 }
287 
288 
289 
290 
291