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
18 // Creates linked list using host code
create_linked_lists_on_host(cl_command_queue cmdq,cl_mem nodes,Node * pNodes2,cl_int ListLength,size_t numLists,cl_bool useNewAPI)19 cl_int create_linked_lists_on_host(cl_command_queue cmdq, cl_mem nodes, Node *pNodes2, cl_int ListLength, size_t numLists, cl_bool useNewAPI )
20 {
21 cl_int error = CL_SUCCESS;
22
23 log_info("SVM: creating linked list on host ");
24
25 Node *pNodes;
26 if (useNewAPI == CL_FALSE)
27 {
28 pNodes = (Node*) clEnqueueMapBuffer(cmdq, nodes, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(Node)*ListLength*numLists, 0, NULL,NULL, &error);
29 test_error2(error, pNodes, "clEnqMapBuffer failed");
30 }
31 else
32 {
33 pNodes = pNodes2;
34 error = clEnqueueSVMMap(cmdq, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, pNodes2, sizeof(Node)*ListLength*numLists, 0, NULL,NULL);
35 test_error2(error, pNodes, "clEnqueueSVMMap failed");
36 }
37
38 create_linked_lists(pNodes, numLists, ListLength);
39
40 if (useNewAPI == CL_FALSE)
41 {
42 error = clEnqueueUnmapMemObject(cmdq, nodes, pNodes, 0,NULL,NULL);
43 test_error(error, "clEnqueueUnmapMemObject failed.");
44 }
45 else
46 {
47 error = clEnqueueSVMUnmap(cmdq, pNodes2, 0, NULL, NULL);
48 test_error(error, "clEnqueueSVMUnmap failed.");
49 }
50
51 error = clFinish(cmdq);
52 test_error(error, "clFinish failed.");
53 return error;
54 }
55
56 // Purpose: uses host code to verify correctness of the linked list
verify_linked_lists_on_host(int ci,cl_command_queue cmdq,cl_mem nodes,Node * pNodes2,cl_int ListLength,size_t numLists,cl_bool useNewAPI)57 cl_int verify_linked_lists_on_host(int ci, cl_command_queue cmdq, cl_mem nodes, Node *pNodes2, cl_int ListLength, size_t numLists, cl_bool useNewAPI )
58 {
59 cl_int error = CL_SUCCESS;
60 cl_int correct_count;
61
62 Node *pNodes;
63 if (useNewAPI == CL_FALSE)
64 {
65 pNodes = (Node*) clEnqueueMapBuffer(cmdq, nodes, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(Node)*ListLength * numLists, 0, NULL,NULL, &error);
66 test_error2(error, pNodes, "clEnqueueMapBuffer failed");
67 }
68 else
69 {
70 pNodes = pNodes2;
71 error = clEnqueueSVMMap(cmdq, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, pNodes2, sizeof(Node)*ListLength * numLists, 0, NULL,NULL);
72 test_error2(error, pNodes, "clEnqueueSVMMap failed");
73 }
74
75 correct_count = 0;
76
77 error = verify_linked_lists(pNodes, numLists, ListLength);
78 if(error) return -1;
79
80 if (useNewAPI == CL_FALSE)
81 {
82 error = clEnqueueUnmapMemObject(cmdq, nodes, pNodes, 0,NULL,NULL);
83 test_error(error, "clEnqueueUnmapMemObject failed.");
84 }
85 else
86 {
87 error = clEnqueueSVMUnmap(cmdq, pNodes2, 0,NULL,NULL);
88 test_error(error, "clEnqueueSVMUnmap failed.");
89 }
90
91 error = clFinish(cmdq);
92 test_error(error, "clFinish failed.");
93 return error;
94 }
95
create_linked_lists_on_device(int ci,cl_command_queue cmdq,cl_mem allocator,cl_kernel kernel_create_lists,size_t numLists)96 cl_int create_linked_lists_on_device(int ci, cl_command_queue cmdq, cl_mem allocator, cl_kernel kernel_create_lists, size_t numLists )
97 {
98 cl_int error = CL_SUCCESS;
99 log_info("SVM: creating linked list on device: %d ", ci);
100
101 size_t *pAllocator = (size_t *)clEnqueueMapBuffer(
102 cmdq, allocator, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(size_t),
103 0, NULL, NULL, &error);
104 test_error2(error, pAllocator, "clEnqueueMapBuffer failed");
105 // reset allocator index
106 *pAllocator = numLists; // the first numLists elements of the nodes array are already allocated (they hold the head of each list).
107 error = clEnqueueUnmapMemObject(cmdq, allocator, pAllocator, 0,NULL,NULL);
108 test_error(error, " clEnqueueUnmapMemObject failed.");
109
110 error = clEnqueueNDRangeKernel(cmdq, kernel_create_lists, 1, NULL, &numLists, NULL, 0, NULL, NULL);
111 test_error(error, "clEnqueueNDRange failed.");
112 error = clFinish(cmdq);
113 test_error(error, "clFinish failed.");
114
115 return error;
116 }
117
verify_linked_lists_on_device(int vi,cl_command_queue cmdq,cl_mem num_correct,cl_kernel kernel_verify_lists,cl_int ListLength,size_t numLists)118 cl_int verify_linked_lists_on_device(int vi, cl_command_queue cmdq,cl_mem num_correct, cl_kernel kernel_verify_lists, cl_int ListLength, size_t numLists )
119 {
120 cl_int error = CL_SUCCESS;
121
122 log_info(" and verifying on device: %d ", vi);
123
124 cl_int *pNumCorrect = (cl_int*) clEnqueueMapBuffer(cmdq, num_correct, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_int), 0, NULL,NULL, &error);
125 test_error2(error, pNumCorrect, "clEnqueueMapBuffer failed");
126
127 *pNumCorrect = 0; // reset numCorrect to zero
128
129 error = clEnqueueUnmapMemObject(cmdq, num_correct, pNumCorrect, 0,NULL,NULL);
130 test_error(error, "clEnqueueUnmapMemObject failed.");
131
132 error = clEnqueueNDRangeKernel(cmdq, kernel_verify_lists, 1, NULL, &numLists, NULL, 0, NULL, NULL);
133 test_error(error,"clEnqueueNDRangeKernel failed");
134
135 pNumCorrect = (cl_int*) clEnqueueMapBuffer(cmdq, num_correct, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_int), 0, NULL,NULL, &error);
136 test_error2(error, pNumCorrect, "clEnqueueMapBuffer failed");
137 cl_int correct_count = *pNumCorrect;
138 error = clEnqueueUnmapMemObject(cmdq, num_correct, pNumCorrect, 0,NULL,NULL);
139 test_error(error, "clEnqueueUnmapMemObject failed");
140 clFinish(cmdq);
141 test_error(error,"clFinish failed");
142
143 if(correct_count != ListLength * (cl_uint)numLists)
144 {
145 error = -1;
146 log_info("Failed\n");
147 }
148 else
149 log_info("Passed\n");
150
151 return error;
152 }
153
154 // This tests that all devices and the host share a common address space; using only the coarse-grain features.
155 // This is done by creating a linked list on a device and then verifying the correctness of the list
156 // on another device or the host. This basic test is performed for all combinations of devices and the host that exist within
157 // the platform. The test passes only if every combination passes.
shared_address_space_coarse_grain(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements,cl_bool useNewAPI)158 int shared_address_space_coarse_grain(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements, cl_bool useNewAPI)
159 {
160 clContextWrapper context = NULL;
161 clProgramWrapper program = NULL;
162 cl_uint num_devices = 0;
163 cl_int error = CL_SUCCESS;
164 clCommandQueueWrapper queues[MAXQ];
165
166 error = create_cl_objects(deviceID, &linked_list_create_and_verify_kernels[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
167 if(error) return -1;
168
169 size_t numLists = num_elements;
170 cl_int ListLength = 32;
171
172 clKernelWrapper kernel_create_lists = clCreateKernel(program, "create_linked_lists", &error);
173 test_error(error, "clCreateKernel failed");
174
175 clKernelWrapper kernel_verify_lists = clCreateKernel(program, "verify_linked_lists", &error);
176 test_error(error, "clCreateKernel failed");
177
178 // this buffer holds the linked list nodes.
179 Node* pNodes = (Node*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(Node)*ListLength*numLists, 0);
180
181 {
182 cl_bool usesSVMpointer = CL_FALSE;
183 clMemWrapper nodes;
184 if (useNewAPI == CL_FALSE)
185 {
186 nodes = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(Node)*ListLength*numLists, pNodes, &error);
187 test_error(error, "clCreateBuffer failed.");
188
189 // verify if buffer uses SVM pointer
190 size_t paramSize = 0;
191 error = clGetMemObjectInfo(nodes, CL_MEM_USES_SVM_POINTER, 0, 0, ¶mSize);
192 test_error(error, "clGetMemObjectInfo failed.");
193
194 if (paramSize != sizeof(cl_bool))
195 {
196 log_error("clGetMemObjectInfo(CL_MEM_USES_SVM_POINTER) returned wrong size.");
197 return -1;
198 }
199
200 error = clGetMemObjectInfo(nodes, CL_MEM_USES_SVM_POINTER, sizeof(cl_bool), &usesSVMpointer, 0);
201 test_error(error, "clGetMemObjectInfo failed.");
202
203 if (usesSVMpointer != CL_TRUE)
204 {
205 log_error("clGetMemObjectInfo(CL_MEM_USES_SVM_POINTER) returned CL_FALSE for buffer created from SVM pointer.");
206 return -1;
207 }
208 }
209
210 // this buffer holds an index into the nodes buffer, it is used for node allocation
211 clMemWrapper allocator = clCreateBuffer(context, CL_MEM_READ_WRITE,
212 sizeof(size_t), NULL, &error);
213
214 test_error(error, "clCreateBuffer failed.");
215
216 error = clGetMemObjectInfo(allocator, CL_MEM_USES_SVM_POINTER, sizeof(cl_bool), &usesSVMpointer, 0);
217 test_error(error, "clGetMemObjectInfo failed.");
218
219 if (usesSVMpointer != CL_FALSE)
220 {
221 log_error("clGetMemObjectInfo(CL_MEM_USES_SVM_POINTER) returned CL_TRUE for non-SVM buffer.");
222 return -1;
223 }
224
225 // this buffer holds the count of correct nodes, which is computed by the verify kernel.
226 clMemWrapper num_correct = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, &error);
227 test_error(error, "clCreateBuffer failed.");
228
229 if (useNewAPI == CL_TRUE)
230 error |= clSetKernelArgSVMPointer(kernel_create_lists, 0, pNodes);
231 else
232 error |= clSetKernelArg(kernel_create_lists, 0, sizeof(void*), (void *) &nodes);
233
234 error |= clSetKernelArg(kernel_create_lists, 1, sizeof(void*), (void *) &allocator);
235 error |= clSetKernelArg(kernel_create_lists, 2, sizeof(cl_int), (void *) &ListLength);
236
237 error |= clSetKernelArgSVMPointer(kernel_verify_lists, 0, pNodes);
238 error |= clSetKernelArg(kernel_verify_lists, 1, sizeof(void*), (void *) &num_correct);
239 error |= clSetKernelArg(kernel_verify_lists, 2, sizeof(cl_int), (void *) &ListLength);
240 test_error(error, "clSetKernelArg failed");
241
242 // Create linked list on one device and verify on another device (or the host).
243 // Do this for all possible combinations of devices and host within the platform.
244 for (int ci=0; ci<(int)num_devices+1; ci++) // ci is CreationIndex, index of device/q to create linked list on
245 {
246 for (int vi=0; vi<(int)num_devices+1; vi++) // vi is VerificationIndex, index of device/q to verify linked list on
247 {
248 if(ci == num_devices) // last device index represents the host, note the num_device+1 above.
249 {
250 error = create_linked_lists_on_host(queues[0], nodes, pNodes, ListLength, numLists, useNewAPI);
251 if(error) return -1;
252 }
253 else
254 {
255 error = create_linked_lists_on_device(ci, queues[ci], allocator, kernel_create_lists, numLists);
256 if(error) return -1;
257 }
258
259 if(vi == num_devices)
260 {
261 error = verify_linked_lists_on_host(vi, queues[0], nodes, pNodes, ListLength, numLists, useNewAPI);
262 if(error) return -1;
263 }
264 else
265 {
266 error = verify_linked_lists_on_device(vi, queues[vi], num_correct, kernel_verify_lists, ListLength, numLists);
267 if(error) return -1;
268 }
269 }
270 }
271 }
272
273 clSVMFree(context, pNodes);
274
275 return 0;
276 }
277
test_svm_shared_address_space_coarse_grain_old_api(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements)278 int test_svm_shared_address_space_coarse_grain_old_api(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
279 {
280 return shared_address_space_coarse_grain(deviceID, context2, queue, num_elements, CL_FALSE);
281 }
282
test_svm_shared_address_space_coarse_grain_new_api(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements)283 int test_svm_shared_address_space_coarse_grain_new_api(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
284 {
285 return shared_address_space_coarse_grain(deviceID, context2, queue, num_elements, CL_TRUE);
286 }
287