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
26 #include "procs.h"
27 #include "harness/conversions.h"
28
29 static const char *async_global_to_local_kernel =
30 "%s\n" // optional pragma string
31 "__kernel void test_fn( const __global %s *src, __global %s *dst, __local %s *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )\n"
32 "{\n"
33 " int i;\n"
34 // Zero the local storage first
35 " for(i=0; i<copiesPerWorkItem; i++)\n"
36 " localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (%s)(%s)0;\n"
37 // Do this to verify all kernels are done zeroing the local buffer before we try the copy
38 " barrier( CLK_LOCAL_MEM_FENCE );\n"
39 " event_t event;\n"
40 " event = async_work_group_copy( (__local %s*)localBuffer, (__global const %s*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, 0 );\n"
41 // Wait for the copy to complete, then verify by manually copying to the dest
42 " wait_group_events( 1, &event );\n"
43 " for(i=0; i<copiesPerWorkItem; i++)\n"
44 " dst[ get_global_id( 0 )*copiesPerWorkItem+i ] = localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ];\n"
45 "}\n" ;
46
47 static const char *async_local_to_global_kernel =
48 "%s\n" // optional pragma string
49 "__kernel void test_fn( const __global %s *src, __global %s *dst, __local %s *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )\n"
50 "{\n"
51 " int i;\n"
52 // Zero the local storage first
53 " for(i=0; i<copiesPerWorkItem; i++)\n"
54 " localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (%s)(%s)0;\n"
55 // Do this to verify all kernels are done zeroing the local buffer before we try the copy
56 " barrier( CLK_LOCAL_MEM_FENCE );\n"
57 " for(i=0; i<copiesPerWorkItem; i++)\n"
58 " localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = src[ get_global_id( 0 )*copiesPerWorkItem+i ];\n"
59 // Do this to verify all kernels are done copying to the local buffer before we try the copy
60 " barrier( CLK_LOCAL_MEM_FENCE );\n"
61 " event_t event;\n"
62 " event = async_work_group_copy((__global %s*)(dst+copiesPerWorkgroup*get_group_id(0)), (__local const %s*)localBuffer, (size_t)copiesPerWorkgroup, 0 );\n"
63 " wait_group_events( 1, &event );\n"
64 "}\n" ;
65
66
67 static const char *prefetch_kernel =
68 "%s\n" // optional pragma string
69 "__kernel void test_fn( const __global %s *src, __global %s *dst, __local %s *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )\n"
70 "{\n"
71 " // Ignore this: %s%s%s\n"
72 " int i;\n"
73 " prefetch( (const __global %s*)(src+copiesPerWorkItem*get_global_id(0)), copiesPerWorkItem);\n"
74 " for(i=0; i<copiesPerWorkItem; i++)\n"
75 " dst[ get_global_id( 0 )*copiesPerWorkItem+i ] = src[ get_global_id( 0 )*copiesPerWorkItem+i ];\n"
76 "}\n" ;
77
78
79
test_copy(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * kernelCode,ExplicitType vecType,int vecSize)80 int test_copy(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *kernelCode,
81 ExplicitType vecType, int vecSize
82 )
83 {
84 int error;
85 clProgramWrapper program;
86 clKernelWrapper kernel;
87 clMemWrapper streams[ 2 ];
88 size_t threads[ 1 ], localThreads[ 1 ];
89 void *inBuffer, *outBuffer;
90 MTdata d;
91 char vecNameString[64]; vecNameString[0] = 0;
92 if (vecSize == 1)
93 sprintf(vecNameString, "%s", get_explicit_type_name(vecType));
94 else
95 sprintf(vecNameString, "%s%d", get_explicit_type_name(vecType), vecSize);
96
97
98 size_t elementSize = get_explicit_type_size(vecType)*vecSize;
99 log_info("Testing %s\n", vecNameString);
100
101 cl_long max_local_mem_size;
102 error = clGetDeviceInfo(deviceID, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(max_local_mem_size), &max_local_mem_size, NULL);
103 test_error( error, "clGetDeviceInfo for CL_DEVICE_LOCAL_MEM_SIZE failed.");
104
105 unsigned int num_of_compute_devices;
106 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(num_of_compute_devices), &num_of_compute_devices, NULL);
107 test_error( error, "clGetDeviceInfo for CL_DEVICE_MAX_COMPUTE_UNITS failed.");
108
109 char programSource[4096]; programSource[0]=0;
110 char *programPtr;
111
112 sprintf(programSource, kernelCode,
113 vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "",
114 vecNameString, vecNameString, vecNameString, vecNameString, get_explicit_type_name(vecType), vecNameString, vecNameString);
115 //log_info("program: %s\n", programSource);
116 programPtr = programSource;
117
118 error = create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "test_fn" );
119 test_error( error, "Unable to create testing kernel" );
120
121 size_t max_workgroup_size;
122 error = clGetKernelWorkGroupInfo(kernel, deviceID, CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_workgroup_size), &max_workgroup_size, NULL);
123 test_error (error, "clGetKernelWorkGroupInfo failed for CL_KERNEL_WORK_GROUP_SIZE.");
124
125 size_t max_local_workgroup_size[3];
126 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_local_workgroup_size), max_local_workgroup_size, NULL);
127 test_error (error, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES");
128
129 // Pick the minimum of the device and the kernel
130 if (max_workgroup_size > max_local_workgroup_size[0])
131 max_workgroup_size = max_local_workgroup_size[0];
132
133 size_t numberOfCopiesPerWorkitem = 13;
134 elementSize = get_explicit_type_size(vecType)* ((vecSize == 3) ? 4 : vecSize);
135 size_t localStorageSpacePerWorkitem = numberOfCopiesPerWorkitem*elementSize;
136 size_t maxLocalWorkgroupSize = (((int)max_local_mem_size/2)/localStorageSpacePerWorkitem);
137
138 // Calculation can return 0 on embedded devices due to 1KB local mem limit
139 if(maxLocalWorkgroupSize == 0)
140 {
141 maxLocalWorkgroupSize = 1;
142 }
143
144 size_t localWorkgroupSize = maxLocalWorkgroupSize;
145 if (maxLocalWorkgroupSize > max_workgroup_size)
146 localWorkgroupSize = max_workgroup_size;
147
148 size_t localBufferSize = localWorkgroupSize*elementSize*numberOfCopiesPerWorkitem;
149 size_t numberOfLocalWorkgroups = 1111;
150 size_t globalBufferSize = numberOfLocalWorkgroups*localBufferSize;
151 size_t globalWorkgroupSize = numberOfLocalWorkgroups*localWorkgroupSize;
152
153 inBuffer = (void*)malloc(globalBufferSize);
154 outBuffer = (void*)malloc(globalBufferSize);
155 memset(outBuffer, 0, globalBufferSize);
156
157 cl_int copiesPerWorkItemInt, copiesPerWorkgroup;
158 copiesPerWorkItemInt = (int)numberOfCopiesPerWorkitem;
159 copiesPerWorkgroup = (int)(numberOfCopiesPerWorkitem*localWorkgroupSize);
160
161 log_info("Global: %d, local %d, local buffer %db, global buffer %db, each work group will copy %d elements and each work item item will copy %d elements.\n",
162 (int) globalWorkgroupSize, (int)localWorkgroupSize, (int)localBufferSize, (int)globalBufferSize, copiesPerWorkgroup, copiesPerWorkItemInt);
163
164 threads[0] = globalWorkgroupSize;
165 localThreads[0] = localWorkgroupSize;
166
167 d = init_genrand( gRandomSeed );
168 generate_random_data( vecType, globalBufferSize/get_explicit_type_size(vecType), d, inBuffer );
169 free_mtdata(d); d = NULL;
170
171 streams[ 0 ] = clCreateBuffer( context, CL_MEM_COPY_HOST_PTR, globalBufferSize, inBuffer, &error );
172 test_error( error, "Unable to create input buffer" );
173 streams[ 1 ] = clCreateBuffer( context, CL_MEM_COPY_HOST_PTR, globalBufferSize, outBuffer, &error );
174 test_error( error, "Unable to create output buffer" );
175
176 error = clSetKernelArg( kernel, 0, sizeof( streams[ 0 ] ), &streams[ 0 ] );
177 test_error( error, "Unable to set kernel argument" );
178 error = clSetKernelArg( kernel, 1, sizeof( streams[ 1 ] ), &streams[ 1 ] );
179 test_error( error, "Unable to set kernel argument" );
180 error = clSetKernelArg( kernel, 2, localBufferSize, NULL );
181 test_error( error, "Unable to set kernel argument" );
182 error = clSetKernelArg( kernel, 3, sizeof(copiesPerWorkgroup), &copiesPerWorkgroup );
183 test_error( error, "Unable to set kernel argument" );
184 error = clSetKernelArg( kernel, 4, sizeof(copiesPerWorkItemInt), &copiesPerWorkItemInt );
185 test_error( error, "Unable to set kernel argument" );
186
187 // Enqueue
188 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
189 test_error( error, "Unable to queue kernel" );
190
191 // Read
192 error = clEnqueueReadBuffer( queue, streams[ 1 ], CL_TRUE, 0, globalBufferSize, outBuffer, 0, NULL, NULL );
193 test_error( error, "Unable to read results" );
194
195 // Verify
196 int failuresPrinted = 0;
197 if( memcmp( inBuffer, outBuffer, globalBufferSize ) != 0 )
198 {
199 size_t typeSize = get_explicit_type_size(vecType)* vecSize;
200 unsigned char * inchar = (unsigned char*)inBuffer;
201 unsigned char * outchar = (unsigned char*)outBuffer;
202 for (int i=0; i< (int)globalBufferSize; i+=(int)elementSize) {
203 if (memcmp( ((char *)inchar)+i, ((char *)outchar)+i, typeSize) != 0 )
204 {
205 char values[4096];
206 values[0] = 0;
207 if ( failuresPrinted == 0 ) {
208 // Print first failure message
209 log_error( "ERROR: Results of copy did not validate!\n" );
210 }
211 sprintf(values + strlen( values), "%d -> [", i);
212 for (int j=0; j<(int)elementSize; j++)
213 sprintf(values + strlen( values), "%2x ", inchar[i+j]);
214 sprintf(values + strlen(values), "] != [");
215 for (int j=0; j<(int)elementSize; j++)
216 sprintf(values + strlen( values), "%2x ", outchar[i+j]);
217 sprintf(values + strlen(values), "]");
218 log_error("%s\n", values);
219 failuresPrinted++;
220 }
221
222 if (failuresPrinted > 5) {
223 log_error("Not printing further failures...\n");
224 break;
225 }
226 }
227 }
228
229 free(inBuffer);
230 free(outBuffer);
231
232 return failuresPrinted ? -1 : 0;
233 }
234
test_copy_all_types(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * kernelCode)235 int test_copy_all_types(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *kernelCode) {
236 ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble, kNumExplicitTypes };
237 unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
238 unsigned int size, typeIndex;
239
240 int errors = 0;
241
242 for( typeIndex = 0; vecType[ typeIndex ] != kNumExplicitTypes; typeIndex++ )
243 {
244 if( vecType[ typeIndex ] == kDouble && !is_extension_available( deviceID, "cl_khr_fp64" ) )
245 continue;
246
247 if (( vecType[ typeIndex ] == kLong || vecType[ typeIndex ] == kULong ) && !gHasLong )
248 continue;
249
250 for( size = 0; vecSizes[ size ] != 0; size++ )
251 {
252 if (test_copy( deviceID, context, queue, kernelCode, vecType[typeIndex],vecSizes[size] )) {
253 errors++;
254 }
255 }
256 }
257 if (errors)
258 return -1;
259 return 0;
260 }
261
262
263
264
test_async_copy_global_to_local(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)265 int test_async_copy_global_to_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
266 {
267 return test_copy_all_types( deviceID, context, queue, async_global_to_local_kernel );
268 }
269
test_async_copy_local_to_global(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)270 int test_async_copy_local_to_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
271 {
272 return test_copy_all_types( deviceID, context, queue, async_local_to_global_kernel );
273 }
274
test_prefetch(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)275 int test_prefetch(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
276 {
277 return test_copy_all_types( deviceID, context, queue, prefetch_kernel );
278 }
279
280