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
19
20
create_linked_lists_on_device_no_map(int ci,cl_command_queue cmdq,size_t * pAllocator,cl_kernel kernel_create_lists,size_t numLists)21 cl_int create_linked_lists_on_device_no_map(int ci, cl_command_queue cmdq, size_t* pAllocator, cl_kernel kernel_create_lists, size_t numLists )
22 {
23 cl_int error = CL_SUCCESS;
24 log_info("SVM: creating linked list on device: %d ", ci);
25
26 // reset allocator index
27 *pAllocator = numLists; // the first numLists elements of the nodes array are already allocated (they hold the head of each list).
28 error = clEnqueueNDRangeKernel(cmdq, kernel_create_lists, 1, NULL, &numLists, NULL, 0, NULL, NULL);
29 test_error(error, "clEnqueueNDRange failed.");
30 error = clFinish(cmdq);
31 test_error(error, "clFinish failed.");
32 return error;
33 }
34
verify_linked_lists_on_device_no_map(int vi,cl_command_queue cmdq,cl_int * pNumCorrect,cl_kernel kernel_verify_lists,cl_int ListLength,size_t numLists)35 cl_int verify_linked_lists_on_device_no_map(int vi, cl_command_queue cmdq,cl_int* pNumCorrect, cl_kernel kernel_verify_lists, cl_int ListLength, size_t numLists )
36 {
37 cl_int error = CL_SUCCESS;
38
39 log_info(" and verifying on device: %d ", vi);
40
41 *pNumCorrect = 0; // reset numCorrect to zero
42
43 error = clEnqueueNDRangeKernel(cmdq, kernel_verify_lists, 1, NULL, &numLists, NULL, 0, NULL, NULL);
44 test_error(error,"clEnqueueNDRangeKernel failed");
45 clFinish(cmdq);
46 test_error(error,"clFinish failed");
47
48 cl_int correct_count = *pNumCorrect;
49 if(correct_count != ListLength * (cl_uint)numLists)
50 {
51 error = -1;
52 log_info("Failed\n");
53 }
54 else
55 log_info("Passed\n");
56
57 return error;
58 }
59
60 // This tests that all devices and the host share a common address space; using only the fine-grain with buffers mode.
61 // This is done by creating a linked list on a device and then verifying the correctness of the list
62 // on another device or the host. This basic test is performed for all combinations of devices and the host that exist within
63 // the platform. The test passes only if every combination passes.
test_svm_shared_address_space_fine_grain_buffers(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements)64 int test_svm_shared_address_space_fine_grain_buffers(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
65 {
66 clContextWrapper context = NULL;
67 clProgramWrapper program = NULL;
68 cl_uint num_devices = 0;
69 cl_int error = CL_SUCCESS;
70 clCommandQueueWrapper queues[MAXQ];
71
72 error = create_cl_objects(deviceID, &linked_list_create_and_verify_kernels[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_FINE_GRAIN_BUFFER);
73 if(error == 1) return 0; // no devices capable of requested SVM level, so don't execute but count test as passing.
74 if(error < 0) return -1; // fail test.
75
76 size_t numLists = num_elements;
77 cl_int ListLength = 32;
78
79 clKernelWrapper kernel_create_lists = clCreateKernel(program, "create_linked_lists", &error);
80 test_error(error, "clCreateKernel failed");
81
82 clKernelWrapper kernel_verify_lists = clCreateKernel(program, "verify_linked_lists", &error);
83 test_error(error, "clCreateKernel failed");
84
85 // this buffer holds the linked list nodes.
86 Node* pNodes = (Node*) clSVMAlloc(context, CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeof(Node)*ListLength*numLists, 0);
87
88 // this buffer holds an index into the nodes buffer, it is used for node allocation
89 size_t *pAllocator = (size_t*) clSVMAlloc(context, CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeof(size_t), 0);
90
91 // this buffer holds the count of correct nodes, which is computed by the verify kernel.
92 cl_int *pNumCorrect = (cl_int*) clSVMAlloc(context, CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER, sizeof(cl_int), 0);
93
94 error |= clSetKernelArgSVMPointer(kernel_create_lists, 0, pNodes);
95 error |= clSetKernelArgSVMPointer(kernel_create_lists, 1, pAllocator);
96 error |= clSetKernelArg(kernel_create_lists, 2, sizeof(cl_int), (void *) &ListLength);
97
98 error |= clSetKernelArgSVMPointer(kernel_verify_lists, 0, pNodes);
99 error |= clSetKernelArgSVMPointer(kernel_verify_lists, 1, pNumCorrect);
100 error |= clSetKernelArg(kernel_verify_lists, 2, sizeof(cl_int), (void *) &ListLength);
101 test_error(error, "clSetKernelArg failed");
102
103 // Create linked list on one device and verify on another device (or the host).
104 // Do this for all possible combinations of devices and host within the platform.
105 for (int ci=0; ci<(int)num_devices+1; ci++) // ci is CreationIndex, index of device/q to create linked list on
106 {
107 for (int vi=0; vi<(int)num_devices+1; vi++) // vi is VerificationIndex, index of device/q to verify linked list on
108 {
109 if(ci == num_devices) // last device index represents the host, note the num_device+1 above.
110 {
111 log_info("SVM: creating linked list on host ");
112 create_linked_lists(pNodes, numLists, ListLength);
113 }
114 else
115 {
116 error = create_linked_lists_on_device_no_map(ci, queues[ci], pAllocator, kernel_create_lists, numLists);
117 if(error) return -1;
118 }
119
120 if(vi == num_devices)
121 {
122 error = verify_linked_lists(pNodes, numLists, ListLength);
123 if(error) return -1;
124 }
125 else
126 {
127 error = verify_linked_lists_on_device_no_map(vi, queues[vi], pNumCorrect, kernel_verify_lists, ListLength, numLists);
128 if(error) return -1;
129 }
130 }
131 }
132
133 clSVMFree(context, pNodes);
134 clSVMFree(context, pAllocator);
135 clSVMFree(context, pNumCorrect);
136
137 return 0;
138 }
139