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(cmdq, allocator, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_int), 0, NULL,NULL, &error);
102 test_error2(error, pAllocator, "clEnqueueMapBuffer failed");
103 // reset allocator index
104 *pAllocator = numLists; // the first numLists elements of the nodes array are already allocated (they hold the head of each list).
105 error = clEnqueueUnmapMemObject(cmdq, allocator, pAllocator, 0,NULL,NULL);
106 test_error(error, " clEnqueueUnmapMemObject failed.");
107
108 error = clEnqueueNDRangeKernel(cmdq, kernel_create_lists, 1, NULL, &numLists, NULL, 0, NULL, NULL);
109 test_error(error, "clEnqueueNDRange failed.");
110 error = clFinish(cmdq);
111 test_error(error, "clFinish failed.");
112
113 return error;
114 }
115
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)116 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 )
117 {
118 cl_int error = CL_SUCCESS;
119
120 log_info(" and verifying on device: %d ", vi);
121
122 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);
123 test_error2(error, pNumCorrect, "clEnqueueMapBuffer failed");
124
125 *pNumCorrect = 0; // reset numCorrect to zero
126
127 error = clEnqueueUnmapMemObject(cmdq, num_correct, pNumCorrect, 0,NULL,NULL);
128 test_error(error, "clEnqueueUnmapMemObject failed.");
129
130 error = clEnqueueNDRangeKernel(cmdq, kernel_verify_lists, 1, NULL, &numLists, NULL, 0, NULL, NULL);
131 test_error(error,"clEnqueueNDRangeKernel failed");
132
133 pNumCorrect = (cl_int*) clEnqueueMapBuffer(cmdq, num_correct, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_int), 0, NULL,NULL, &error);
134 test_error2(error, pNumCorrect, "clEnqueueMapBuffer failed");
135 cl_int correct_count = *pNumCorrect;
136 error = clEnqueueUnmapMemObject(cmdq, num_correct, pNumCorrect, 0,NULL,NULL);
137 test_error(error, "clEnqueueUnmapMemObject failed");
138 clFinish(cmdq);
139 test_error(error,"clFinish failed");
140
141 if(correct_count != ListLength * (cl_uint)numLists)
142 {
143 error = -1;
144 log_info("Failed\n");
145 }
146 else
147 log_info("Passed\n");
148
149 return error;
150 }
151
152 // This tests that all devices and the host share a common address space; using only the coarse-grain features.
153 // This is done by creating a linked list on a device and then verifying the correctness of the list
154 // on another device or the host. This basic test is performed for all combinations of devices and the host that exist within
155 // 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)156 int shared_address_space_coarse_grain(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements, cl_bool useNewAPI)
157 {
158 clContextWrapper context = NULL;
159 clProgramWrapper program = NULL;
160 cl_uint num_devices = 0;
161 cl_int error = CL_SUCCESS;
162 clCommandQueueWrapper queues[MAXQ];
163
164 error = create_cl_objects(deviceID, &linked_list_create_and_verify_kernels[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
165 if(error) return -1;
166
167 size_t numLists = num_elements;
168 cl_int ListLength = 32;
169
170 clKernelWrapper kernel_create_lists = clCreateKernel(program, "create_linked_lists", &error);
171 test_error(error, "clCreateKernel failed");
172
173 clKernelWrapper kernel_verify_lists = clCreateKernel(program, "verify_linked_lists", &error);
174 test_error(error, "clCreateKernel failed");
175
176 // this buffer holds the linked list nodes.
177 Node* pNodes = (Node*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(Node)*ListLength*numLists, 0);
178
179 {
180 cl_bool usesSVMpointer = CL_FALSE;
181 clMemWrapper nodes;
182 if (useNewAPI == CL_FALSE)
183 {
184 nodes = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(Node)*ListLength*numLists, pNodes, &error);
185 test_error(error, "clCreateBuffer failed.");
186
187 // verify if buffer uses SVM pointer
188 size_t paramSize = 0;
189 error = clGetMemObjectInfo(nodes, CL_MEM_USES_SVM_POINTER, 0, 0, ¶mSize);
190 test_error(error, "clGetMemObjectInfo failed.");
191
192 if (paramSize != sizeof(cl_bool))
193 {
194 log_error("clGetMemObjectInfo(CL_MEM_USES_SVM_POINTER) returned wrong size.");
195 return -1;
196 }
197
198 error = clGetMemObjectInfo(nodes, CL_MEM_USES_SVM_POINTER, sizeof(cl_bool), &usesSVMpointer, 0);
199 test_error(error, "clGetMemObjectInfo failed.");
200
201 if (usesSVMpointer != CL_TRUE)
202 {
203 log_error("clGetMemObjectInfo(CL_MEM_USES_SVM_POINTER) returned CL_FALSE for buffer created from SVM pointer.");
204 return -1;
205 }
206 }
207
208 // this buffer holds an index into the nodes buffer, it is used for node allocation
209 clMemWrapper allocator = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, &error);
210 test_error(error, "clCreateBuffer failed.");
211
212 error = clGetMemObjectInfo(allocator, CL_MEM_USES_SVM_POINTER, sizeof(cl_bool), &usesSVMpointer, 0);
213 test_error(error, "clGetMemObjectInfo failed.");
214
215 if (usesSVMpointer != CL_FALSE)
216 {
217 log_error("clGetMemObjectInfo(CL_MEM_USES_SVM_POINTER) returned CL_TRUE for non-SVM buffer.");
218 return -1;
219 }
220
221 // this buffer holds the count of correct nodes, which is computed by the verify kernel.
222 clMemWrapper num_correct = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, &error);
223 test_error(error, "clCreateBuffer failed.");
224
225 if (useNewAPI == CL_TRUE)
226 error |= clSetKernelArgSVMPointer(kernel_create_lists, 0, pNodes);
227 else
228 error |= clSetKernelArg(kernel_create_lists, 0, sizeof(void*), (void *) &nodes);
229
230 error |= clSetKernelArg(kernel_create_lists, 1, sizeof(void*), (void *) &allocator);
231 error |= clSetKernelArg(kernel_create_lists, 2, sizeof(cl_int), (void *) &ListLength);
232
233 error |= clSetKernelArgSVMPointer(kernel_verify_lists, 0, pNodes);
234 error |= clSetKernelArg(kernel_verify_lists, 1, sizeof(void*), (void *) &num_correct);
235 error |= clSetKernelArg(kernel_verify_lists, 2, sizeof(cl_int), (void *) &ListLength);
236 test_error(error, "clSetKernelArg failed");
237
238 // Create linked list on one device and verify on another device (or the host).
239 // Do this for all possible combinations of devices and host within the platform.
240 for (int ci=0; ci<(int)num_devices+1; ci++) // ci is CreationIndex, index of device/q to create linked list on
241 {
242 for (int vi=0; vi<(int)num_devices+1; vi++) // vi is VerificationIndex, index of device/q to verify linked list on
243 {
244 if(ci == num_devices) // last device index represents the host, note the num_device+1 above.
245 {
246 error = create_linked_lists_on_host(queues[0], nodes, pNodes, ListLength, numLists, useNewAPI);
247 if(error) return -1;
248 }
249 else
250 {
251 error = create_linked_lists_on_device(ci, queues[ci], allocator, kernel_create_lists, numLists);
252 if(error) return -1;
253 }
254
255 if(vi == num_devices)
256 {
257 error = verify_linked_lists_on_host(vi, queues[0], nodes, pNodes, ListLength, numLists, useNewAPI);
258 if(error) return -1;
259 }
260 else
261 {
262 error = verify_linked_lists_on_device(vi, queues[vi], num_correct, kernel_verify_lists, ListLength, numLists);
263 if(error) return -1;
264 }
265 }
266 }
267 }
268
269 clSVMFree(context, pNodes);
270
271 return 0;
272 }
273
test_svm_shared_address_space_coarse_grain_old_api(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements)274 int test_svm_shared_address_space_coarse_grain_old_api(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
275 {
276 return shared_address_space_coarse_grain(deviceID, context2, queue, num_elements, CL_FALSE);
277 }
278
test_svm_shared_address_space_coarse_grain_new_api(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements)279 int test_svm_shared_address_space_coarse_grain_new_api(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
280 {
281 return shared_address_space_coarse_grain(deviceID, context2, queue, num_elements, CL_TRUE);
282 }
283