1 //
2 // Copyright (c) 2022 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
17 #include "basic_command_buffer.h"
18 #include "procs.h"
19
20 #include <vector>
21
22 namespace {
23
24 ////////////////////////////////////////////////////////////////////////////////
25 // clSetKernelArg tests for cl_khr_command_buffer which handles below cases:
26 // -test interactions of clSetKernelArg with command-buffers
27 // -test interactions of clSetKernelArg on a command-buffer pending execution
28
29 template <bool simul_use>
30 struct CommandBufferSetKernelArg : public BasicCommandBufferTest
31 {
CommandBufferSetKernelArg__anon523ff78a0111::CommandBufferSetKernelArg32 CommandBufferSetKernelArg(cl_device_id device, cl_context context,
33 cl_command_queue queue)
34 : BasicCommandBufferTest(device, context, queue), trigger_event(nullptr)
35 {
36 simultaneous_use_requested = simul_use;
37 if (simul_use) buffer_size_multiplier = 2;
38 }
39
40 //--------------------------------------------------------------------------
SetUpKernel__anon523ff78a0111::CommandBufferSetKernelArg41 cl_int SetUpKernel() override
42 {
43 cl_int error = CL_SUCCESS;
44
45 const char* kernel_str =
46 R"(
47 __kernel void copy(int in, __global int* out, __global int* offset)
48 {
49 size_t id = get_global_id(0);
50 size_t ind = offset[0] + id;
51 out[ind] = in;
52 })";
53
54 error = create_single_kernel_helper_create_program(context, &program, 1,
55 &kernel_str);
56 test_error(error, "Failed to create program with source");
57
58 error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
59 test_error(error, "Failed to build program");
60
61 kernel = clCreateKernel(program, "copy", &error);
62 test_error(error, "Failed to create copy kernel");
63
64 return CL_SUCCESS;
65 }
66
67 //--------------------------------------------------------------------------
SetUpKernelArgs__anon523ff78a0111::CommandBufferSetKernelArg68 cl_int SetUpKernelArgs() override
69 {
70 cl_int error = CL_SUCCESS;
71 out_mem = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
72 num_elements * buffer_size_multiplier
73 * sizeof(cl_int),
74 nullptr, &error);
75 test_error(error, "clCreateBuffer failed");
76
77 // create secondary output buffer to test kernel args substitution
78 out_mem_k2 = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
79 num_elements * buffer_size_multiplier
80 * sizeof(cl_int),
81 nullptr, &error);
82 test_error(error, "clCreateBuffer failed");
83
84 cl_int offset = 0;
85 off_mem =
86 clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
87 sizeof(cl_int), &offset, &error);
88 test_error(error, "clCreateBuffer failed");
89
90 cl_int in_arg = pattern_pri;
91 error = clSetKernelArg(kernel, 0, sizeof(cl_int), &in_arg);
92 test_error(error, "clSetKernelArg failed");
93
94 error = clSetKernelArg(kernel, 1, sizeof(out_mem), &out_mem);
95 test_error(error, "clSetKernelArg failed");
96
97 error = clSetKernelArg(kernel, 2, sizeof(off_mem), &off_mem);
98 test_error(error, "clSetKernelArg failed");
99
100 return CL_SUCCESS;
101 }
102
103 //--------------------------------------------------------------------------
Run__anon523ff78a0111::CommandBufferSetKernelArg104 cl_int Run() override
105 {
106 cl_int error = CL_SUCCESS;
107 if (simultaneous_use_requested)
108 {
109 // enqueue simultaneous command-buffers with clSetKernelArg calls
110 error = RunSimultaneous();
111 test_error(error, "RunSimultaneous failed");
112 }
113 else
114 {
115 // enqueue single command-buffer with clSetKernelArg calls
116 error = RunSingle();
117 test_error(error, "RunSingle failed");
118 }
119
120 return CL_SUCCESS;
121 }
122
123 //--------------------------------------------------------------------------
RecordCommandBuffer__anon523ff78a0111::CommandBufferSetKernelArg124 cl_int RecordCommandBuffer()
125 {
126 cl_int error = CL_SUCCESS;
127
128 error = clCommandNDRangeKernelKHR(
129 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
130 nullptr, 0, nullptr, nullptr, nullptr);
131 test_error(error, "clCommandNDRangeKernelKHR failed");
132
133 // changing kernel args at this point should have no effect,
134 // test will verify if clSetKernelArg didn't affect the first command
135 cl_int in_arg = pattern_sec;
136 error = clSetKernelArg(kernel, 0, sizeof(cl_int), &in_arg);
137 test_error(error, "clSetKernelArg failed");
138
139 error = clSetKernelArg(kernel, 1, sizeof(out_mem_k2), &out_mem_k2);
140 test_error(error, "clSetKernelArg failed");
141
142 error = clCommandNDRangeKernelKHR(
143 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
144 nullptr, 0, nullptr, nullptr, nullptr);
145 test_error(error, "clCommandNDRangeKernelKHR failed");
146
147 error = clFinalizeCommandBufferKHR(command_buffer);
148 test_error(error, "clFinalizeCommandBufferKHR failed");
149 return CL_SUCCESS;
150 }
151
152 //--------------------------------------------------------------------------
RunSingle__anon523ff78a0111::CommandBufferSetKernelArg153 cl_int RunSingle()
154 {
155 cl_int error = CL_SUCCESS;
156 std::vector<cl_int> output_data(num_elements);
157
158 // record command buffer
159 error = RecordCommandBuffer();
160 test_error(error, "RecordCommandBuffer failed");
161
162 const cl_int pattern_base = 0;
163 error =
164 clEnqueueFillBuffer(queue, out_mem, &pattern_base, sizeof(cl_int),
165 0, data_size(), 0, nullptr, nullptr);
166 test_error(error, "clEnqueueFillBuffer failed");
167
168 const cl_int pattern_base_k2 = 1;
169 error = clEnqueueFillBuffer(queue, out_mem_k2, &pattern_base_k2,
170 sizeof(cl_int), 0, data_size(), 0, nullptr,
171 nullptr);
172 test_error(error, "clEnqueueFillBuffer failed");
173
174 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
175 nullptr, nullptr);
176 test_error(error, "clEnqueueCommandBufferKHR failed");
177
178 error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
179 output_data.data(), 0, nullptr, nullptr);
180 test_error(error, "clEnqueueReadBuffer failed");
181
182 // verify the result - result buffer must contain initial pattern
183 for (size_t i = 0; i < num_elements; i++)
184 {
185 CHECK_VERIFICATION_ERROR(pattern_pri, output_data[i], i);
186 }
187
188 return CL_SUCCESS;
189 }
190
191 //--------------------------------------------------------------------------
192 struct SimulPassData
193 {
194 cl_int pattern;
195 cl_int offset;
196 std::vector<cl_int> output_buffer;
197 };
198
199 //--------------------------------------------------------------------------
RecordSimultaneousCommandBuffer__anon523ff78a0111::CommandBufferSetKernelArg200 cl_int RecordSimultaneousCommandBuffer() const
201 {
202 cl_int error = CL_SUCCESS;
203
204 error = clCommandNDRangeKernelKHR(
205 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
206 nullptr, 0, nullptr, nullptr, nullptr);
207 test_error(error, "clCommandNDRangeKernelKHR failed");
208
209 error = clFinalizeCommandBufferKHR(command_buffer);
210 test_error(error, "clFinalizeCommandBufferKHR failed");
211 return CL_SUCCESS;
212 }
213
214 //--------------------------------------------------------------------------
EnqueueSimultaneousPass__anon523ff78a0111::CommandBufferSetKernelArg215 cl_int EnqueueSimultaneousPass(SimulPassData& pd)
216 {
217 cl_int error = clEnqueueFillBuffer(
218 queue, out_mem, &pd.pattern, sizeof(cl_int),
219 pd.offset * sizeof(cl_int), data_size(), 0, nullptr, nullptr);
220 test_error(error, "clEnqueueFillBuffer failed");
221
222 error = clEnqueueFillBuffer(queue, off_mem, &pd.offset, sizeof(cl_int),
223 0, sizeof(cl_int), 0, nullptr, nullptr);
224 test_error(error, "clEnqueueFillBuffer failed");
225
226 if (!trigger_event)
227 {
228 trigger_event = clCreateUserEvent(context, &error);
229 test_error(error, "clCreateUserEvent failed");
230 }
231
232 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 1,
233 &trigger_event, nullptr);
234 test_error(error, "clEnqueueCommandBufferKHR failed");
235
236 error = clEnqueueReadBuffer(
237 queue, out_mem, CL_FALSE, pd.offset * sizeof(cl_int), data_size(),
238 pd.output_buffer.data(), 0, nullptr, nullptr);
239 test_error(error, "clEnqueueReadBuffer failed");
240
241 return CL_SUCCESS;
242 }
243
244 //--------------------------------------------------------------------------
RunSimultaneous__anon523ff78a0111::CommandBufferSetKernelArg245 cl_int RunSimultaneous()
246 {
247 cl_int error = CL_SUCCESS;
248
249 // record command buffer with primary queue
250 error = RecordSimultaneousCommandBuffer();
251 test_error(error, "RecordSimultaneousCommandBuffer failed");
252
253 std::vector<SimulPassData> simul_passes = {
254 { 0, 0, std::vector<cl_int>(num_elements) }
255 };
256
257 error = EnqueueSimultaneousPass(simul_passes.front());
258 test_error(error, "EnqueueSimultaneousPass 1 failed");
259
260 // changing kernel args at this point should have no effect,
261 // test will verify if clSetKernelArg didn't affect command-buffer
262 cl_int in_arg = pattern_sec;
263 error = clSetKernelArg(kernel, 0, sizeof(cl_int), &in_arg);
264 test_error(error, "clSetKernelArg failed");
265
266 error = clSetKernelArg(kernel, 1, sizeof(out_mem_k2), &out_mem_k2);
267 test_error(error, "clSetKernelArg failed");
268
269 if (simultaneous_use_support)
270 {
271 cl_int offset = static_cast<cl_int>(num_elements);
272 simul_passes.push_back(
273 { 1, offset, std::vector<cl_int>(num_elements) });
274
275 error = EnqueueSimultaneousPass(simul_passes.back());
276 test_error(error, "EnqueueSimultaneousPass 2 failed");
277 }
278
279 error = clSetUserEventStatus(trigger_event, CL_COMPLETE);
280 test_error(error, "clSetUserEventStatus failed");
281
282 error = clFinish(queue);
283 test_error(error, "clFinish failed");
284
285 // verify the result buffer
286 for (auto&& pass : simul_passes)
287 {
288 auto& res_data = pass.output_buffer;
289 for (size_t i = 0; i < num_elements; i++)
290 {
291 CHECK_VERIFICATION_ERROR(pattern_pri, res_data[i], i);
292 }
293 }
294
295 return CL_SUCCESS;
296 }
297
298 //--------------------------------------------------------------------------
299 clEventWrapper trigger_event = nullptr;
300
301 const cl_int pattern_pri = 2;
302 const cl_int pattern_sec = 3;
303
304 clMemWrapper out_mem_k2 = nullptr;
305 };
306
307 } // anonymous namespace
308
test_basic_set_kernel_arg(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)309 int test_basic_set_kernel_arg(cl_device_id device, cl_context context,
310 cl_command_queue queue, int num_elements)
311 {
312 return MakeAndRunTest<CommandBufferSetKernelArg<false>>(
313 device, context, queue, num_elements);
314 }
315
test_pending_set_kernel_arg(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)316 int test_pending_set_kernel_arg(cl_device_id device, cl_context context,
317 cl_command_queue queue, int num_elements)
318 {
319 return MakeAndRunTest<CommandBufferSetKernelArg<true>>(device, context,
320 queue, num_elements);
321 }
322