xref: /aosp_15_r20/external/OpenCL-CTS/test_conformance/SVM/test_migrate.cpp (revision 6467f958c7de8070b317fc65bcb0f6472e388d82) 
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_event evs[20];
82 
83     const size_t global_size = GLOBAL_SIZE;
84 
85     RandomSeed seed(0);
86 
87     clContextWrapper context = NULL;
88     clCommandQueueWrapper queues[MAXQ];
89     cl_uint num_devices = 0;
90     clProgramWrapper program;
91     cl_int error;
92 
93     error = create_cl_objects(deviceID, &sources[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
94     if (error)
95         return -1;
96 
97     if (num_devices > 1) {
98         log_info("  Running on two devices.\n");
99     } else {
100         // Ensure we have two distinct queues
101         cl_device_id did;
102         error = clGetCommandQueueInfo(queues[0], CL_QUEUE_DEVICE, sizeof(did), (void *)&did, NULL);
103         test_error(error, "clGetCommandQueueInfo failed");
104 
105         cl_command_queue_properties cqp;
106         error = clGetCommandQueueInfo(queues[0], CL_QUEUE_PROPERTIES, sizeof(cqp), &cqp, NULL);
107         test_error(error, "clGetCommandQueueInfo failed");
108 
109         cl_queue_properties qp[3] = { CL_QUEUE_PROPERTIES, cqp, 0 };
110         queues[1] = clCreateCommandQueueWithProperties(context, did, qp, &error);
111         test_error(error, "clCteateCommandQueueWithProperties failed");
112     }
113 
114     clKernelWrapper kernel = clCreateKernel(program, "migrate_kernel", &error);
115     test_error(error, "clCreateKernel failed");
116 
117     char* asvm = (char*)clSVMAlloc(context, CL_MEM_READ_WRITE, global_size*sizeof(cl_uint), 16);
118     if (asvm == NULL) {
119         log_error("ERROR: clSVMAlloc returned NULL at %s:%d\n", __FILE__, __LINE__);
120         return -1;
121     }
122 
123     char* bsvm = (char *)clSVMAlloc(context, CL_MEM_READ_WRITE, global_size*sizeof(cl_uint), 16);
124     if (bsvm == NULL) {
125         log_error("ERROR: clSVMAlloc returned NULL at %s:%d\n", __FILE__, __LINE__);
126         clSVMFree(context, asvm);
127         return -1;
128     }
129 
130     char* csvm = (char *)clSVMAlloc(context, CL_MEM_READ_WRITE, global_size*sizeof(cl_uint), 16);
131     if (csvm == NULL) {
132         log_error("ERROR: clSVMAlloc returned NULL at %s:%d\n", __FILE__, __LINE__);
133         clSVMFree(context, bsvm);
134         clSVMFree(context, asvm);
135         return -1;
136     }
137 
138     error = clSetKernelArgSVMPointer(kernel, 0, (void*)asvm);
139     test_error(error, "clSetKernelArgSVMPointer failed");
140 
141     error = clSetKernelArgSVMPointer(kernel, 1, (void*)bsvm);
142     test_error(error, "clSetKernelArgSVMPointer failed");
143 
144     error = clSetKernelArgSVMPointer(kernel, 2, (void*)csvm);
145     test_error(error, "clSetKernelArgSVMPointer failed");
146 
147     // Initialize host copy of data (and result)
148     fill_buffer(amem, global_size, seed);
149     fill_buffer(bmem, global_size, seed);
150     fill_buffer(cmem, global_size, seed);
151 
152     // Now we're ready to start
153     {
154         // First, fill in the data on device0
155         cl_uint patt[] = { 0, 0, 0, 0};
156         error = clEnqueueSVMMemFill(queues[0], (void *)asvm, patt, sizeof(patt), global_size*sizeof(cl_uint), 0, NULL, &evs[0]);
157         test_error(error, "clEnqueueSVMMemFill failed");
158 
159         error = clEnqueueSVMMemFill(queues[0], (void *)bsvm, patt, sizeof(patt), global_size*sizeof(cl_uint), 0, NULL, &evs[1]);
160         test_error(error, "clEnqueueSVMMemFill failed");
161 
162         error = clEnqueueSVMMemFill(queues[0], (void *)csvm, patt, sizeof(patt), global_size*sizeof(cl_uint), 0, NULL, &evs[2]);
163         test_error(error, "clEnqueueSVMMemFill failed");
164     }
165 
166     {
167         // Now migrate fully to device 1 and discard the data
168         char* ptrs[] = { asvm, bsvm, csvm };
169         error = clEnqueueSVMMigrateMem(queues[1], 3, (const void**)ptrs, NULL, CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED, 1, &evs[2], &evs[3]);
170         test_error(error, "clEnqueueSVMMigrateMem failed");
171     }
172 
173     {
174         // Test host flag
175         char *ptrs[] = { asvm+1, bsvm+3, csvm+5 };
176         const size_t szs[] = { 1, 1, 0 };
177         error = clEnqueueSVMMigrateMem(queues[0], 3, (const void**)ptrs, szs, CL_MIGRATE_MEM_OBJECT_HOST, 1, &evs[3], &evs[4]);
178         test_error(error, "clEnqueueSVMMigrateMem failed");
179     }
180 
181     {
182         // Next fill with known data
183         error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_WRITE, (void*)asvm, global_size*sizeof(cl_uint), 1, &evs[4], &evs[5]);
184         test_error(error, "clEnqueueSVMMap failed");
185 
186         error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_WRITE, (void*)bsvm, global_size*sizeof(cl_uint), 0, NULL, &evs[6]);
187         test_error(error, "clEnqueueSVMMap failed");
188 
189         error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_WRITE, (void*)csvm, global_size*sizeof(cl_uint), 0, NULL, &evs[7]);
190         test_error(error, "clEnqueueSVMMap failed");
191     }
192 
193     error = clFlush(queues[0]);
194     test_error(error, "clFlush failed");
195 
196     error = clFlush(queues[1]);
197     test_error(error, "clFlush failed");
198 
199     // Check the event command type for clEnqueueSVMMigrateMem (OpenCL 3.0 and
200     // newer)
201     Version version = get_device_cl_version(deviceID);
202     if (version >= Version(3, 0))
203     {
204         cl_command_type commandType;
205         error = clGetEventInfo(evs[3], CL_EVENT_COMMAND_TYPE,
206                                sizeof(commandType), &commandType, NULL);
207         test_error(error, "clGetEventInfo failed");
208         if (commandType != CL_COMMAND_SVM_MIGRATE_MEM)
209         {
210             log_error("Invalid command type returned for "
211                       "clEnqueueSVMMigrateMem: %X\n",
212                       commandType);
213             return TEST_FAIL;
214         }
215     }
216 
217     error = wait_and_release("first batch", evs, 8);
218     if (error)
219         return -1;
220 
221     memcpy((void *)asvm, (void *)amem, global_size*sizeof(cl_uint));
222     memcpy((void *)bsvm, (void *)bmem, global_size*sizeof(cl_uint));
223     memcpy((void *)csvm, (void *)cmem, global_size*sizeof(cl_uint));
224 
225     {
226         error = clEnqueueSVMUnmap(queues[1], (void *)asvm, 0, NULL, &evs[0]);
227         test_error(error, "clEnqueueSVMUnmap failed");
228 
229         error = clEnqueueSVMUnmap(queues[1], (void *)bsvm, 0, NULL, &evs[1]);
230         test_error(error, "clEnqueueSVMUnmap failed");
231 
232         error = clEnqueueSVMUnmap(queues[1], (void *)csvm, 0, NULL, &evs[2]);
233         test_error(error, "clEnqueueSVMUnmap failed");
234     }
235 
236 
237     {
238         // Now try some overlapping regions, and operate on the result
239         char *ptrs[] = { asvm+100, bsvm+17, csvm+1000, asvm+101, bsvm+19, csvm+1017 };
240         const size_t szs[] = { 13, 23, 43, 3, 7, 11 };
241 
242         error = clEnqueueSVMMigrateMem(queues[0], 3, (const void**)ptrs, szs, 0, 1, &evs[2], &evs[3]);
243         test_error(error, "clEnqueueSVMMigrateMem failed");
244 
245         error = clEnqueueNDRangeKernel(queues[0], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[4]);
246         test_error(error, "clEnqueueNDRangeKernel failed");
247     }
248 
249     {
250         // Now another pair
251         char *ptrs[] = { asvm+8, bsvm+17, csvm+31, csvm+83 };
252         const size_t szs[] = { 0, 1, 3, 7 };
253 
254         error = clEnqueueSVMMigrateMem(queues[1], 4, (const void**)ptrs, szs, 0, 1, &evs[4], &evs[5]);
255         test_error(error, "clEnqueueSVMMigrateMem failed");
256 
257         error = clEnqueueNDRangeKernel(queues[1], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[6]);
258         test_error(error, "clEnqueueNDRangeKernel failed");
259     }
260 
261     {
262         // Another pair
263         char *ptrs[] = { asvm+64, asvm+128, bsvm+64, bsvm+128, csvm, csvm+64 };
264         const size_t szs[] = { 64, 64, 64, 64, 64, 64 };
265 
266         error = clEnqueueSVMMigrateMem(queues[0], 6, (const void**)ptrs, szs, 0, 1, &evs[6], &evs[7]);
267         test_error(error, "clEnqueueSVMMigrateMem failed");
268 
269         error = clEnqueueNDRangeKernel(queues[0], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[8]);
270         test_error(error, "clEnqueueNDRangeKernel failed");
271     }
272 
273     {
274         // Final pair
275         char *ptrs[] = { asvm, asvm, bsvm, csvm, csvm };
276         const size_t szs[] = { 0, 1, 0, 1, 0 };
277 
278         error = clEnqueueSVMMigrateMem(queues[1], 5, (const void**)ptrs, szs, 0, 1, &evs[8], &evs[9]);
279         test_error(error, "clEnqueueSVMMigrateMem failed");
280 
281         error = clEnqueueNDRangeKernel(queues[1], kernel, 1, NULL, &global_size, NULL, 0, NULL, &evs[10]);
282         test_error(error, "clEnqueueNDRangeKernel failed");
283     }
284 
285     {
286         error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_READ, (void*)asvm, global_size*sizeof(cl_uint), 0, NULL, &evs[11]);
287         test_error(error, "clEnqueueSVMMap failed");
288 
289         error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_READ, (void*)bsvm, global_size*sizeof(cl_uint), 0, NULL, &evs[12]);
290         test_error(error, "clEnqueueSVMMap failed");
291 
292         error = clEnqueueSVMMap(queues[1], CL_FALSE, CL_MAP_READ, (void*)csvm, global_size*sizeof(cl_uint), 0, NULL, &evs[13]);
293         test_error(error, "clEnqueueSVMMap failed");
294     }
295 
296     error = clFlush(queues[0]);
297     test_error(error, "clFlush failed");
298 
299     error = clFlush(queues[1]);
300     test_error(error, "clFlush failed");
301 
302     error = wait_and_release("batch 2", evs, 14);
303     if (error)
304         return -1;
305 
306     // Check kernel results
307     bool ok = check("memory a", (cl_uint *)asvm, amem, global_size);
308     ok &= check("memory b", (cl_uint *)bsvm, bmem, global_size);
309     ok &= check("memory c", (cl_uint *)csvm, cmem, global_size);
310 
311     {
312         void *ptrs[] = { asvm, bsvm, csvm };
313 
314         error = clEnqueueSVMUnmap(queues[1], (void *)asvm, 0, NULL, &evs[0]);
315         test_error(error, "clEnqueueSVMUnmap failed");
316 
317         error = clEnqueueSVMUnmap(queues[1], (void *)bsvm, 0, NULL, &evs[1]);
318         test_error(error, "clEnqueueSVMUnmap failed");
319 
320         error = clEnqueueSVMUnmap(queues[1], (void *)csvm, 0, NULL, &evs[2]);
321         test_error(error, "clEnqueueSVMUnmap failed");
322 
323         error = clEnqueueSVMFree(queues[1], 3, ptrs, NULL, NULL, 0, NULL, &evs[3]);
324     }
325 
326     error = clFlush(queues[1]);
327     test_error(error, "clFlush failed");
328 
329     error = wait_and_release("batch 3", evs, 4);
330     if (error)
331         return -1;
332 
333     // The wrappers will clean up the rest
334     return ok ? 0 : -1;
335 }
336 
337