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 typedef struct {
19 cl_int *pA;
20 cl_int *pB;
21 cl_int *pC;
22 } BufPtrs;
23
24 const char *set_kernel_exec_info_svm_ptrs_kernel[] = {
25 "struct BufPtrs;\n"
26 "\n"
27 "typedef struct {\n"
28 " __global int *pA;\n"
29 " __global int *pB;\n"
30 " __global int *pC;\n"
31 "} BufPtrs;\n"
32 "\n"
33 "__kernel void set_kernel_exec_info_test(__global BufPtrs* pBufs)\n"
34 "{\n"
35 " size_t i;\n"
36 " i = get_global_id(0);\n"
37 " pBufs->pA[i]++;\n"
38 " pBufs->pB[i]++;\n"
39 " pBufs->pC[i]++;\n"
40 "}\n"
41 };
42
43 // Test that clSetKernelExecInfo works correctly with CL_KERNEL_EXEC_INFO_SVM_PTRS flag.
44 //
test_svm_set_kernel_exec_info_svm_ptrs(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements)45 int test_svm_set_kernel_exec_info_svm_ptrs(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
46 {
47 clContextWrapper c = NULL;
48 clProgramWrapper program = NULL;
49 cl_uint num_devices = 0;
50 cl_int error = CL_SUCCESS;
51 clCommandQueueWrapper queues[MAXQ];
52
53 //error = create_cl_objects(deviceID, &set_kernel_exec_info_svm_ptrs_kernel[0], &context, &program, &q, &num_devices, CL_DEVICE_SVM_FINE_GRAIN);
54 error = create_cl_objects(deviceID, &set_kernel_exec_info_svm_ptrs_kernel[0], &c, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
55 if(error == 1) return 0; // no devices capable of requested SVM level, so don't execute but count test as passing.
56 if(error < 0) return -1; // fail test.
57
58
59 clKernelWrapper k = clCreateKernel(program, "set_kernel_exec_info_test", &error);
60 test_error(error, "clCreateKernel failed");
61
62 size_t size = num_elements*sizeof(int);
63 //int* pA = (int*) clSVMalloc(c, CL_MEM_READ_WRITE | CL_DEVICE_SVM_FINE_GRAIN_SYSTEM, sizeof(int)*num_elements, 0);
64 //int* pB = (int*) clSVMalloc(c, CL_MEM_READ_WRITE | CL_DEVICE_SVM_FINE_GRAIN_SYSTEM, sizeof(int)*num_elements, 0);
65 //int* pC = (int*) clSVMalloc(c, CL_MEM_READ_WRITE | CL_DEVICE_SVM_FINE_GRAIN_SYSTEM, sizeof(int)*num_elements, 0);
66 int* pA = (int*) clSVMAlloc(c, CL_MEM_READ_WRITE, size, 0);
67 int* pB = (int*) clSVMAlloc(c, CL_MEM_READ_WRITE, size, 0);
68 int* pC = (int*) clSVMAlloc(c, CL_MEM_READ_WRITE, size, 0);
69 BufPtrs* pBuf = (BufPtrs*) clSVMAlloc(c, CL_MEM_READ_WRITE, sizeof(BufPtrs), 0);
70
71 bool failed = false;
72 {
73 clMemWrapper ba,bb,bc,bBuf;
74 ba = clCreateBuffer(c, CL_MEM_USE_HOST_PTR, size, pA, &error);
75 test_error(error, "clCreateBuffer failed");
76 bb = clCreateBuffer(c, CL_MEM_USE_HOST_PTR, size, pB, &error);
77 test_error(error, "clCreateBuffer failed");
78 bc = clCreateBuffer(c, CL_MEM_USE_HOST_PTR, size, pC, &error);
79 test_error(error, "clCreateBuffer failed");
80 bBuf = clCreateBuffer(c, CL_MEM_USE_HOST_PTR, sizeof(BufPtrs), pBuf, &error);
81 test_error(error, "clCreateBuffer failed");
82
83 clEnqueueMapBuffer(queues[0], ba, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, size, 0, NULL, NULL, &error);
84 test_error(error, "clEnqueueMapBuffer failed");
85 clEnqueueMapBuffer(queues[0], bb, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, size, 0, NULL, NULL, &error);
86 test_error(error, "clEnqueueMapBuffer failed");
87 clEnqueueMapBuffer(queues[0], bc, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, size, 0, NULL, NULL, &error);
88 test_error(error, "clEnqueueMapBuffer failed");
89 clEnqueueMapBuffer(queues[0], bBuf, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(BufPtrs), 0, NULL, NULL, &error);
90 test_error(error, "clEnqueueMapBuffer failed");
91
92 for(int i = 0; i < num_elements; i++) pA[i] = pB[i] = pC[i] = 0;
93
94 pBuf->pA = pA;
95 pBuf->pB = pB;
96 pBuf->pC = pC;
97
98 error = clEnqueueUnmapMemObject(queues[0], ba, pA, 0, NULL, NULL);
99 test_error(error, " clEnqueueUnmapMemObject failed.");
100 error = clEnqueueUnmapMemObject(queues[0], bb, pB, 0, NULL, NULL);
101 test_error(error, " clEnqueueUnmapMemObject failed.");
102 error = clEnqueueUnmapMemObject(queues[0], bc, pC, 0, NULL, NULL);
103 test_error(error, " clEnqueueUnmapMemObject failed.");
104 error = clEnqueueUnmapMemObject(queues[0], bBuf, pBuf, 0, NULL, NULL);
105 test_error(error, " clEnqueueUnmapMemObject failed.");
106
107
108 error = clSetKernelArgSVMPointer(k, 0, pBuf);
109 test_error(error, "clSetKernelArg failed");
110
111 error = clSetKernelExecInfo(k, CL_KERNEL_EXEC_INFO_SVM_PTRS, sizeof(BufPtrs), pBuf);
112 test_error(error, "clSetKernelExecInfo failed");
113
114 size_t range = num_elements;
115 error = clEnqueueNDRangeKernel(queues[0], k, 1, NULL, &range, NULL, 0, NULL, NULL);
116 test_error(error,"clEnqueueNDRangeKernel failed");
117
118 error = clFinish(queues[0]);
119 test_error(error, "clFinish failed.");
120
121 clEnqueueMapBuffer(queues[0], ba, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, size, 0, NULL, NULL, &error);
122 test_error(error, "clEnqueueMapBuffer failed");
123 clEnqueueMapBuffer(queues[0], bb, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, size, 0, NULL, NULL, &error);
124 test_error(error, "clEnqueueMapBuffer failed");
125 clEnqueueMapBuffer(queues[0], bc, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, size, 0, NULL, NULL, &error);
126 test_error(error, "clEnqueueMapBuffer failed");
127
128 for(int i = 0; i < num_elements; i++)
129 {
130 if(pA[i] + pB[i] + pC[i] != 3)
131 failed = true;
132 }
133
134 error = clEnqueueUnmapMemObject(queues[0], ba, pA, 0, NULL, NULL);
135 test_error(error, " clEnqueueUnmapMemObject failed.");
136 error = clEnqueueUnmapMemObject(queues[0], bb, pB, 0, NULL, NULL);
137 test_error(error, " clEnqueueUnmapMemObject failed.");
138 error = clEnqueueUnmapMemObject(queues[0], bc, pC, 0, NULL, NULL);
139 test_error(error, " clEnqueueUnmapMemObject failed.");
140 }
141
142 error = clFinish(queues[0]);
143 test_error(error, " clFinish failed.");
144
145 clSVMFree(c, pA);
146 clSVMFree(c, pB);
147 clSVMFree(c, pC);
148 clSVMFree(c, pBuf);
149
150 if(failed) return -1;
151
152 return 0;
153 }
154