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