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 "common.h"
17 #include "harness/mt19937.h"
18
19 #define GLOBAL_SIZE 65536
20
21 static const char *sources[] = {
22 "__kernel void migrate_kernel(__global uint * restrict a, __global uint * restrict b, __global uint * restrict c)\n"
23 "{\n"
24 " size_t i = get_global_id(0);\n"
25 " a[i] ^= 0x13579bdf;\n"
26 " b[i] ^= 0x2468ace0;\n"
27 " c[i] ^= 0x731fec8f;\n"
28 "}\n"
29 };
30
31 static void
fill_buffer(cl_uint * p,size_t n,MTdata seed)32 fill_buffer(cl_uint* p, size_t n, MTdata seed)
33 {
34 for (size_t i=0; i<n; ++i)
35 p[i] = (cl_uint)genrand_int32(seed);
36 }
37
38 static bool
check(const char * s,cl_uint * a,cl_uint * e,size_t n)39 check(const char* s, cl_uint* a, cl_uint* e, size_t n)
40 {
41 bool ok = true;
42 for (size_t i=0; ok && i<n; ++i) {
43 if (a[i] != e[i]) {
44 log_error("ERROR: %s mismatch at word %u, *%08x vs %08x\n", s, (unsigned int)i, e[i], a[i]);
45 ok = false;
46 }
47 }
48 return ok;
49 }
50
51 static int
wait_and_release(const char * s,cl_event * evs,int n)52 wait_and_release(const char* s, cl_event* evs, int n)
53 {
54 cl_int error = clWaitForEvents(n, evs);
55 if (error == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST) {
56 for (int i=0; i<n; ++i) {
57 cl_int e;
58 error = clGetEventInfo(evs[i], CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(cl_int), &e, NULL);
59 test_error(error, "clGetEventInfo failed");
60 if (e != CL_COMPLETE) {
61 log_error("ERROR: %s event %d execution status was %s\n", s, i, IGetErrorString(e));
62 return e;
63 }
64 }
65 } else
66 test_error(error, "clWaitForEvents failed");
67
68 for (int i=0; i<n; ++i) {
69 error = clReleaseEvent(evs[i]);
70 test_error(error, "clReleaseEvent failed");
71 }
72
73 return 0;
74 }
75
test_svm_migrate(cl_device_id deviceID,cl_context c,cl_command_queue queue,int num_elements)76 int test_svm_migrate(cl_device_id deviceID, cl_context c, cl_command_queue queue, int num_elements)
77 {
78 cl_uint amem[GLOBAL_SIZE];
79 cl_uint bmem[GLOBAL_SIZE];
80 cl_uint cmem[GLOBAL_SIZE];
81 cl_uint ramem[GLOBAL_SIZE];
82 cl_uint rbmem[GLOBAL_SIZE];
83 cl_uint rcmem[GLOBAL_SIZE];
84 cl_event evs[20];
85
86 const size_t global_size = GLOBAL_SIZE;
87
88 RandomSeed seed(0);
89
90 clContextWrapper context = NULL;
91 clCommandQueueWrapper queues[MAXQ];
92 cl_uint num_devices = 0;
93 clProgramWrapper program;
94 cl_int error;
95
96 error = create_cl_objects(deviceID, &sources[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
97 if (error)
98 return -1;
99
100 if (num_devices > 1) {
101 log_info(" Running on two devices.\n");
102 } else {
103 // Ensure we have two distinct queues
104 cl_device_id did;
105 error = clGetCommandQueueInfo(queues[0], CL_QUEUE_DEVICE, sizeof(did), (void *)&did, NULL);
106 test_error(error, "clGetCommandQueueInfo failed");
107
108 cl_command_queue_properties cqp;
109 error = clGetCommandQueueInfo(queues[0], CL_QUEUE_PROPERTIES, sizeof(cqp), &cqp, NULL);
110 test_error(error, "clGetCommandQueueInfo failed");
111
112 cl_queue_properties qp[3] = { CL_QUEUE_PROPERTIES, cqp, 0 };
113 queues[1] = clCreateCommandQueueWithProperties(context, did, qp, &error);
114 test_error(error, "clCteateCommandQueueWithProperties failed");
115 }
116
117 clKernelWrapper kernel = clCreateKernel(program, "migrate_kernel", &error);
118 test_error(error, "clCreateKernel failed");
119
120 char* asvm = (char*)clSVMAlloc(context, CL_MEM_READ_WRITE, global_size*sizeof(cl_uint), 16);
121 if (asvm == NULL) {
122 log_error("ERROR: clSVMAlloc returned NULL at %s:%d\n", __FILE__, __LINE__);
123 return -1;
124 }
125
126 char* bsvm = (char *)clSVMAlloc(context, CL_MEM_READ_WRITE, global_size*sizeof(cl_uint), 16);
127 if (bsvm == NULL) {
128 log_error("ERROR: clSVMAlloc returned NULL at %s:%d\n", __FILE__, __LINE__);
129 clSVMFree(context, asvm);
130 return -1;
131 }
132
133 char* csvm = (char *)clSVMAlloc(context, CL_MEM_READ_WRITE, global_size*sizeof(cl_uint), 16);
134 if (csvm == NULL) {
135 log_error("ERROR: clSVMAlloc returned NULL at %s:%d\n", __FILE__, __LINE__);
136 clSVMFree(context, bsvm);
137 clSVMFree(context, asvm);
138 return -1;
139 }
140
141 error = clSetKernelArgSVMPointer(kernel, 0, (void*)asvm);
142 test_error(error, "clSetKernelArgSVMPointer failed");
143
144 error = clSetKernelArgSVMPointer(kernel, 1, (void*)bsvm);
145 test_error(error, "clSetKernelArgSVMPointer failed");
146
147 error = clSetKernelArgSVMPointer(kernel, 2, (void*)csvm);
148 test_error(error, "clSetKernelArgSVMPointer failed");
149
150 // Initialize host copy of data (and result)
151 fill_buffer(amem, global_size, seed);
152 fill_buffer(bmem, global_size, seed);
153 fill_buffer(cmem, global_size, seed);
154
155 // Now we're ready to start
156 {
157 // First, fill in the data on device0
158 cl_uint patt[] = { 0, 0, 0, 0};
159 error = clEnqueueSVMMemFill(queues[0], (void *)asvm, patt, sizeof(patt), global_size*sizeof(cl_uint), 0, NULL, &evs[0]);
160 test_error(error, "clEnqueueSVMMemFill failed");
161
162 error = clEnqueueSVMMemFill(queues[0], (void *)bsvm, patt, sizeof(patt), global_size*sizeof(cl_uint), 0, NULL, &evs[1]);
163 test_error(error, "clEnqueueSVMMemFill failed");
164
165 error = clEnqueueSVMMemFill(queues[0], (void *)csvm, patt, sizeof(patt), global_size*sizeof(cl_uint), 0, NULL, &evs[2]);
166 test_error(error, "clEnqueueSVMMemFill failed");
167 }
168
169 {
170 // Now migrate fully to device 1 and discard the data
171 char* ptrs[] = { asvm, bsvm, csvm };
172 error = clEnqueueSVMMigrateMem(queues[1], 3, (const void**)ptrs, NULL, CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED, 1, &evs[2], &evs[3]);
173 test_error(error, "clEnqueueSVMMigrateMem failed");
174 }
175
176 {
177 // Test host flag
178 char *ptrs[] = { asvm+1, bsvm+3, csvm+5 };
179 const size_t szs[] = { 1, 1, 0 };
180 error = clEnqueueSVMMigrateMem(queues[0], 3, (const void**)ptrs, szs, CL_MIGRATE_MEM_OBJECT_HOST, 1, &evs[3], &evs[4]);
181 test_error(error, "clEnqueueSVMMigrateMem failed");
182 }
183
184 {
185 // Next fill with known data
186 error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_WRITE, (void*)asvm, global_size*sizeof(cl_uint), 1, &evs[4], &evs[5]);
187 test_error(error, "clEnqueueSVMMap failed");
188
189 error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_WRITE, (void*)bsvm, global_size*sizeof(cl_uint), 0, NULL, &evs[6]);
190 test_error(error, "clEnqueueSVMMap failed");
191
192 error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_WRITE, (void*)csvm, global_size*sizeof(cl_uint), 0, NULL, &evs[7]);
193 test_error(error, "clEnqueueSVMMap failed");
194 }
195
196 error = clFlush(queues[0]);
197 test_error(error, "clFlush failed");
198
199 error = clFlush(queues[1]);
200 test_error(error, "clFlush failed");
201
202 // Check the event command type for clEnqueueSVMMigrateMem (OpenCL 3.0 and
203 // newer)
204 Version version = get_device_cl_version(deviceID);
205 if (version >= Version(3, 0))
206 {
207 cl_command_type commandType;
208 error = clGetEventInfo(evs[3], CL_EVENT_COMMAND_TYPE,
209 sizeof(commandType), &commandType, NULL);
210 test_error(error, "clGetEventInfo failed");
211 if (commandType != CL_COMMAND_SVM_MIGRATE_MEM)
212 {
213 log_error("Invalid command type returned for "
214 "clEnqueueSVMMigrateMem: %X\n",
215 commandType);
216 return TEST_FAIL;
217 }
218 }
219
220 error = wait_and_release("first batch", evs, 8);
221 if (error)
222 return -1;
223
224 memcpy((void *)asvm, (void *)amem, global_size*sizeof(cl_uint));
225 memcpy((void *)bsvm, (void *)bmem, global_size*sizeof(cl_uint));
226 memcpy((void *)csvm, (void *)cmem, global_size*sizeof(cl_uint));
227
228 {
229 error = clEnqueueSVMUnmap(queues[1], (void *)asvm, 0, NULL, &evs[0]);
230 test_error(error, "clEnqueueSVMUnmap failed");
231
232 error = clEnqueueSVMUnmap(queues[1], (void *)bsvm, 0, NULL, &evs[1]);
233 test_error(error, "clEnqueueSVMUnmap failed");
234
235 error = clEnqueueSVMUnmap(queues[1], (void *)csvm, 0, NULL, &evs[2]);
236 test_error(error, "clEnqueueSVMUnmap failed");
237 }
238
239
240 {
241 // Now try some overlapping regions, and operate on the result
242 char *ptrs[] = { asvm+100, bsvm+17, csvm+1000, asvm+101, bsvm+19, csvm+1017 };
243 const size_t szs[] = { 13, 23, 43, 3, 7, 11 };
244
245 error = clEnqueueSVMMigrateMem(queues[0], 3, (const void**)ptrs, szs, 0, 1, &evs[2], &evs[3]);
246 test_error(error, "clEnqueueSVMMigrateMem failed");
247
248 error = clEnqueueNDRangeKernel(queues[0], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[4]);
249 test_error(error, "clEnqueueNDRangeKernel failed");
250 }
251
252 {
253 // Now another pair
254 char *ptrs[] = { asvm+8, bsvm+17, csvm+31, csvm+83 };
255 const size_t szs[] = { 0, 1, 3, 7 };
256
257 error = clEnqueueSVMMigrateMem(queues[1], 4, (const void**)ptrs, szs, 0, 1, &evs[4], &evs[5]);
258 test_error(error, "clEnqueueSVMMigrateMem failed");
259
260 error = clEnqueueNDRangeKernel(queues[1], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[6]);
261 test_error(error, "clEnqueueNDRangeKernel failed");
262 }
263
264 {
265 // Another pair
266 char *ptrs[] = { asvm+64, asvm+128, bsvm+64, bsvm+128, csvm, csvm+64 };
267 const size_t szs[] = { 64, 64, 64, 64, 64, 64 };
268
269 error = clEnqueueSVMMigrateMem(queues[0], 6, (const void**)ptrs, szs, 0, 1, &evs[6], &evs[7]);
270 test_error(error, "clEnqueueSVMMigrateMem failed");
271
272 error = clEnqueueNDRangeKernel(queues[0], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[8]);
273 test_error(error, "clEnqueueNDRangeKernel failed");
274 }
275
276 {
277 // Final pair
278 char *ptrs[] = { asvm, asvm, bsvm, csvm, csvm };
279 const size_t szs[] = { 0, 1, 0, 1, 0 };
280
281 error = clEnqueueSVMMigrateMem(queues[1], 5, (const void**)ptrs, szs, 0, 1, &evs[8], &evs[9]);
282 test_error(error, "clEnqueueSVMMigrateMem failed");
283
284 error = clEnqueueNDRangeKernel(queues[1], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[10]);
285 test_error(error, "clEnqueueNDRangeKernel failed");
286 }
287
288 {
289 error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_READ, (void*)asvm, global_size*sizeof(cl_uint), 0, NULL, &evs[11]);
290 test_error(error, "clEnqueueSVMMap failed");
291
292 error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_READ, (void*)bsvm, global_size*sizeof(cl_uint), 0, NULL, &evs[12]);
293 test_error(error, "clEnqueueSVMMap failed");
294
295 error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_READ, (void*)csvm, global_size*sizeof(cl_uint), 0, NULL, &evs[13]);
296 test_error(error, "clEnqueueSVMMap failed");
297 }
298
299 error = clFlush(queues[0]);
300 test_error(error, "clFlush failed");
301
302 error = clFlush(queues[1]);
303 test_error(error, "clFlush failed");
304
305 error = wait_and_release("batch 2", evs, 14);
306 if (error)
307 return -1;
308
309 // Check kernel results
310 bool ok = check("memory a", (cl_uint *)asvm, amem, global_size);
311 ok &= check("memory b", (cl_uint *)bsvm, bmem, global_size);
312 ok &= check("memory c", (cl_uint *)csvm, cmem, global_size);
313
314 {
315 void *ptrs[] = { asvm, bsvm, csvm };
316
317 error = clEnqueueSVMUnmap(queues[1], (void *)asvm, 0, NULL, &evs[0]);
318 test_error(error, "clEnqueueSVMUnmap failed");
319
320 error = clEnqueueSVMUnmap(queues[1], (void *)bsvm, 0, NULL, &evs[1]);
321 test_error(error, "clEnqueueSVMUnmap failed");
322
323 error = clEnqueueSVMUnmap(queues[1], (void *)csvm, 0, NULL, &evs[2]);
324 test_error(error, "clEnqueueSVMUnmap failed");
325
326 error = clEnqueueSVMFree(queues[1], 3, ptrs, NULL, NULL, 0, NULL, &evs[3]);
327 }
328
329 error = clFlush(queues[1]);
330 test_error(error, "clFlush failed");
331
332 error = wait_and_release("batch 3", evs, 4);
333 if (error)
334 return -1;
335
336 // The wrappers will clean up the rest
337 return ok ? 0 : -1;
338 }
339
340