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 <vulkan_interop_common.hpp>
18 #include <vulkan_wrapper.hpp>
19 #include <CL/cl.h>
20 #include <CL/cl_ext.h>
21 #include <assert.h>
22 #include <vector>
23 #include <iostream>
24 #include <string.h>
25 #include "harness/errorHelpers.h"
26
27 #define MAX_BUFFERS 5
28 #define MAX_IMPORTS 5
29 #define BUFFERSIZE 3000
30 static cl_uchar uuid[CL_UUID_SIZE_KHR];
31 static cl_device_id deviceId = NULL;
32
33 namespace {
34 struct Params
35 {
36 uint32_t numBuffers;
37 uint32_t bufferSize;
38 uint32_t interBufferOffset;
39 };
40 }
41
42 const char *kernel_text_numbuffer_1 = " \
43 __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a) { \n\
44 int gid = get_global_id(0); \n\
45 if (gid < bufferSize) { \n\
46 a[gid]++; \n\
47 } \n\
48 }";
49
50 const char *kernel_text_numbuffer_2 = " \
51 __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global unsigned char *b) { \n\
52 int gid = get_global_id(0); \n\
53 if (gid < bufferSize) { \n\
54 a[gid]++; \n\
55 b[gid]++;\n\
56 } \n\
57 }";
58
59 const char *kernel_text_numbuffer_4 = " \
60 __kernel void clUpdateBuffer(int bufferSize, __global unsigned char *a, __global unsigned char *b, __global unsigned char *c, __global unsigned char *d) { \n\
61 int gid = get_global_id(0); \n\
62 if (gid < bufferSize) { \n\
63 a[gid]++;\n\
64 b[gid]++; \n\
65 c[gid]++; \n\
66 d[gid]++; \n\
67 } \n\
68 }";
69
70
71 const char *kernel_text_verify = " \
72 __kernel void checkKernel(__global unsigned char *ptr, int size, int expVal, __global unsigned char *err) \n\
73 { \n\
74 int idx = get_global_id(0); \n\
75 if ((idx < size) && (*err == 0)) { \n\
76 if (ptr[idx] != expVal){ \n\
77 *err = 1; \n\
78 } \n\
79 } \n\
80 }";
81
run_test_with_two_queue(cl_context & context,cl_command_queue & cmd_queue1,cl_command_queue & cmd_queue2,cl_kernel * kernel,cl_kernel & verify_kernel,VulkanDevice & vkDevice,uint32_t numBuffers,uint32_t bufferSize)82 int run_test_with_two_queue(cl_context &context, cl_command_queue &cmd_queue1,
83 cl_command_queue &cmd_queue2, cl_kernel *kernel,
84 cl_kernel &verify_kernel, VulkanDevice &vkDevice,
85 uint32_t numBuffers, uint32_t bufferSize)
86 {
87 int err = CL_SUCCESS;
88 size_t global_work_size[1];
89 uint8_t *error_2;
90 cl_mem error_1;
91 cl_kernel update_buffer_kernel;
92 cl_kernel kernel_cq;
93 clExternalSemaphore *clVk2CLExternalSemaphore = NULL;
94 clExternalSemaphore *clCl2VkExternalSemaphore = NULL;
95 const char *program_source_const = kernel_text_numbuffer_2;
96 size_t program_source_length = strlen(program_source_const);
97 cl_program program = clCreateProgramWithSource(
98 context, 1, &program_source_const, &program_source_length, &err);
99 err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
100 if (err != CL_SUCCESS)
101 {
102 print_error(err, "Error: Failed to build program \n");
103 return err;
104 }
105 // create the kernel
106 kernel_cq = clCreateKernel(program, "clUpdateBuffer", &err);
107 if (err != CL_SUCCESS)
108 {
109 print_error(err, "clCreateKernel failed \n");
110 return err;
111 }
112
113 const std::vector<VulkanExternalMemoryHandleType>
114 vkExternalMemoryHandleTypeList =
115 getSupportedVulkanExternalMemoryHandleTypeList();
116 VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType =
117 getSupportedVulkanExternalSemaphoreHandleTypeList()[0];
118 VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType);
119 VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType);
120
121 VulkanQueue &vkQueue = vkDevice.getQueue();
122
123 std::vector<char> vkBufferShader = readFile("buffer.spv");
124
125 VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader);
126 VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList(
127 MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER);
128 VulkanDescriptorSetLayout vkDescriptorSetLayout(
129 vkDevice, vkDescriptorSetLayoutBindingList);
130 VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout);
131 VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout,
132 vkBufferShaderModule);
133
134 VulkanDescriptorPool vkDescriptorPool(vkDevice,
135 vkDescriptorSetLayoutBindingList);
136 VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool,
137 vkDescriptorSetLayout);
138
139 clVk2CLExternalSemaphore = new clExternalSemaphore(
140 vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
141 clCl2VkExternalSemaphore = new clExternalSemaphore(
142 vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
143
144 const uint32_t maxIter = innerIterations;
145 VulkanCommandPool vkCommandPool(vkDevice);
146 VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool);
147
148 VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params));
149 VulkanDeviceMemory vkParamsDeviceMemory(
150 vkDevice, vkParamsBuffer.getSize(),
151 getVulkanMemoryType(vkDevice,
152 VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT));
153 vkParamsDeviceMemory.bindBuffer(vkParamsBuffer);
154 std::vector<VulkanDeviceMemory *> vkBufferListDeviceMemory;
155 std::vector<clExternalMemory *> externalMemory;
156 for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size();
157 emhtIdx++)
158 {
159 VulkanExternalMemoryHandleType vkExternalMemoryHandleType =
160 vkExternalMemoryHandleTypeList[emhtIdx];
161 log_info("External memory handle type: %d\n",
162 vkExternalMemoryHandleType);
163
164 VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024,
165 vkExternalMemoryHandleType);
166 const VulkanMemoryTypeList &memoryTypeList =
167 vkDummyBuffer.getMemoryTypeList();
168
169 for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++)
170 {
171 const VulkanMemoryType &memoryType = memoryTypeList[mtIdx];
172
173 log_info("Memory type index: %d\n", (uint32_t)memoryType);
174 log_info("Memory type property: %d\n",
175 memoryType.getMemoryTypeProperty());
176
177 VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize,
178 vkExternalMemoryHandleType);
179
180 for (size_t bIdx = 0; bIdx < numBuffers; bIdx++)
181 {
182 vkBufferListDeviceMemory.push_back(
183 new VulkanDeviceMemory(vkDevice, bufferSize, memoryType,
184 vkExternalMemoryHandleType));
185 externalMemory.push_back(new clExternalMemory(
186 vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType,
187 0, bufferSize, context, deviceId));
188 }
189 cl_mem buffers[MAX_BUFFERS];
190 clFinish(cmd_queue1);
191 Params *params = (Params *)vkParamsDeviceMemory.map();
192 params->numBuffers = numBuffers;
193 params->bufferSize = bufferSize;
194 params->interBufferOffset = 0;
195 vkParamsDeviceMemory.unmap();
196 vkDescriptorSet.update(0, vkParamsBuffer);
197 for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++)
198 {
199 size_t buffer_size = vkBufferList[bIdx].getSize();
200 vkBufferListDeviceMemory[bIdx]->bindBuffer(vkBufferList[bIdx],
201 0);
202 buffers[bIdx] = externalMemory[bIdx]->getExternalMemoryBuffer();
203 vkDescriptorSet.update((uint32_t)bIdx + 1, vkBufferList[bIdx]);
204 }
205 vkCommandBuffer.begin();
206 vkCommandBuffer.bindPipeline(vkComputePipeline);
207 vkCommandBuffer.bindDescriptorSets(
208 vkComputePipeline, vkPipelineLayout, vkDescriptorSet);
209 vkCommandBuffer.dispatch(512, 1, 1);
210 vkCommandBuffer.end();
211
212 if (vkBufferList.size() == 2)
213 {
214 update_buffer_kernel = kernel[0];
215 }
216 else if (vkBufferList.size() == 3)
217 {
218 update_buffer_kernel = kernel[1];
219 }
220 else if (vkBufferList.size() == 5)
221 {
222 update_buffer_kernel = kernel[2];
223 }
224 // global work size should be less than or equal to
225 // bufferSizeList[i]
226 global_work_size[0] = bufferSize;
227 for (uint32_t iter = 0; iter < maxIter; iter++)
228 {
229
230 if (iter == 0)
231 {
232 vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore);
233 }
234 else
235 {
236 vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer,
237 vkVk2CLSemaphore);
238 }
239 clVk2CLExternalSemaphore->wait(cmd_queue1);
240
241 err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t),
242 (void *)&bufferSize);
243 err |= clSetKernelArg(kernel_cq, 0, sizeof(uint32_t),
244 (void *)&bufferSize);
245 err |= clSetKernelArg(kernel_cq, 1, sizeof(cl_mem),
246 (void *)&(buffers[0]));
247
248 for (int i = 0; i < vkBufferList.size() - 1; i++)
249 {
250 err |=
251 clSetKernelArg(update_buffer_kernel, i + 1,
252 sizeof(cl_mem), (void *)&(buffers[i]));
253 }
254
255 err |=
256 clSetKernelArg(kernel_cq, 2, sizeof(cl_mem),
257 (void *)&(buffers[vkBufferList.size() - 1]));
258
259 if (err != CL_SUCCESS)
260 {
261 print_error(err,
262 "Error: Failed to set arg values for kernel\n");
263 goto CLEANUP;
264 }
265 cl_event first_launch;
266
267 err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel,
268 1, NULL, global_work_size, NULL, 0,
269 NULL, &first_launch);
270 if (err != CL_SUCCESS)
271 {
272 print_error(err,
273 "Error: Failed to launch update_buffer_kernel,"
274 "error\n");
275 goto CLEANUP;
276 }
277
278 err = clEnqueueNDRangeKernel(cmd_queue2, kernel_cq, 1, NULL,
279 global_work_size, NULL, 1,
280 &first_launch, NULL);
281 if (err != CL_SUCCESS)
282 {
283 print_error(err,
284 "Error: Failed to launch update_buffer_kernel,"
285 "error\n");
286 goto CLEANUP;
287 }
288
289 if (iter != (maxIter - 1))
290 {
291 clCl2VkExternalSemaphore->signal(cmd_queue2);
292 }
293 }
294 error_2 = (uint8_t *)malloc(sizeof(uint8_t));
295 if (NULL == error_2)
296 {
297 log_error("Not able to allocate memory\n");
298 goto CLEANUP;
299 }
300 clFinish(cmd_queue2);
301 error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
302 sizeof(uint8_t), NULL, &err);
303 if (CL_SUCCESS != err)
304 {
305 print_error(err, "Error: clCreateBuffer \n");
306 goto CLEANUP;
307 }
308 uint8_t val = 0;
309 err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0,
310 sizeof(uint8_t), &val, 0, NULL, NULL);
311 if (err != CL_SUCCESS)
312 {
313 print_error(err, "Error: Failed read output, error\n");
314 goto CLEANUP;
315 }
316
317 int calc_max_iter;
318 for (int i = 0; i < vkBufferList.size(); i++)
319 {
320 if (i == 0)
321 calc_max_iter = (maxIter * 3);
322 else
323 calc_max_iter = (maxIter * 2);
324 err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem),
325 (void *)&(buffers[i]));
326 err |=
327 clSetKernelArg(verify_kernel, 1, sizeof(int), &bufferSize);
328 err |= clSetKernelArg(verify_kernel, 2, sizeof(int),
329 &calc_max_iter);
330 err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem),
331 (void *)&error_1);
332 if (err != CL_SUCCESS)
333 {
334 print_error(err,
335 "Error: Failed to set arg values for "
336 "verify_kernel \n");
337 goto CLEANUP;
338 }
339 err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL,
340 global_work_size, NULL, 0, NULL,
341 NULL);
342
343 if (err != CL_SUCCESS)
344 {
345 print_error(err,
346 "Error: Failed to launch verify_kernel,"
347 "error \n");
348 goto CLEANUP;
349 }
350 err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0,
351 sizeof(uint8_t), error_2, 0, NULL,
352 NULL);
353 if (err != CL_SUCCESS)
354 {
355 print_error(err, "Error: Failed read output, error \n ");
356 goto CLEANUP;
357 }
358 if (*error_2 == 1)
359 {
360 log_error("&&&& vulkan_opencl_buffer test FAILED\n");
361 goto CLEANUP;
362 }
363 }
364 for (size_t i = 0; i < vkBufferList.size(); i++)
365 {
366 delete vkBufferListDeviceMemory[i];
367 delete externalMemory[i];
368 }
369 vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(),
370 vkBufferListDeviceMemory.begin()
371 + numBuffers);
372 externalMemory.erase(externalMemory.begin(),
373 externalMemory.begin() + numBuffers);
374 }
375 }
376 CLEANUP:
377 for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++)
378 {
379 if (vkBufferListDeviceMemory[i])
380 {
381 delete vkBufferListDeviceMemory[i];
382 }
383 if (externalMemory[i])
384 {
385 delete externalMemory[i];
386 }
387 }
388 if (program) clReleaseProgram(program);
389 if (kernel_cq) clReleaseKernel(kernel_cq);
390 if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore;
391 if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore;
392 if (error_2) free(error_2);
393 if (error_1) clReleaseMemObject(error_1);
394
395 return err;
396 }
397
run_test_with_one_queue(cl_context & context,cl_command_queue & cmd_queue1,cl_kernel * kernel,cl_kernel & verify_kernel,VulkanDevice & vkDevice,uint32_t numBuffers,uint32_t bufferSize)398 int run_test_with_one_queue(cl_context &context, cl_command_queue &cmd_queue1,
399 cl_kernel *kernel, cl_kernel &verify_kernel,
400 VulkanDevice &vkDevice, uint32_t numBuffers,
401 uint32_t bufferSize)
402 {
403 log_info("RUNNING TEST WITH ONE QUEUE...... \n\n");
404 size_t global_work_size[1];
405 uint8_t *error_2;
406 cl_mem error_1;
407 cl_kernel update_buffer_kernel;
408 clExternalSemaphore *clVk2CLExternalSemaphore = NULL;
409 clExternalSemaphore *clCl2VkExternalSemaphore = NULL;
410 int err = CL_SUCCESS;
411
412 const std::vector<VulkanExternalMemoryHandleType>
413 vkExternalMemoryHandleTypeList =
414 getSupportedVulkanExternalMemoryHandleTypeList();
415 VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType =
416 getSupportedVulkanExternalSemaphoreHandleTypeList()[0];
417 VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType);
418 VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType);
419
420 VulkanQueue &vkQueue = vkDevice.getQueue();
421
422 std::vector<char> vkBufferShader = readFile("buffer.spv");
423 VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader);
424 VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList(
425 MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER);
426 VulkanDescriptorSetLayout vkDescriptorSetLayout(
427 vkDevice, vkDescriptorSetLayoutBindingList);
428 VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout);
429 VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout,
430 vkBufferShaderModule);
431
432 VulkanDescriptorPool vkDescriptorPool(vkDevice,
433 vkDescriptorSetLayoutBindingList);
434 VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool,
435 vkDescriptorSetLayout);
436
437 clVk2CLExternalSemaphore = new clExternalSemaphore(
438 vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
439 clCl2VkExternalSemaphore = new clExternalSemaphore(
440 vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
441 const uint32_t maxIter = innerIterations;
442 VulkanCommandPool vkCommandPool(vkDevice);
443 VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool);
444
445 VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params));
446 VulkanDeviceMemory vkParamsDeviceMemory(
447 vkDevice, vkParamsBuffer.getSize(),
448 getVulkanMemoryType(vkDevice,
449 VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT));
450 vkParamsDeviceMemory.bindBuffer(vkParamsBuffer);
451 std::vector<VulkanDeviceMemory *> vkBufferListDeviceMemory;
452 std::vector<clExternalMemory *> externalMemory;
453
454 for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size();
455 emhtIdx++)
456 {
457 VulkanExternalMemoryHandleType vkExternalMemoryHandleType =
458 vkExternalMemoryHandleTypeList[emhtIdx];
459 log_info("External memory handle type: %d\n",
460 vkExternalMemoryHandleType);
461
462 VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024,
463 vkExternalMemoryHandleType);
464 const VulkanMemoryTypeList &memoryTypeList =
465 vkDummyBuffer.getMemoryTypeList();
466
467 for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++)
468 {
469 const VulkanMemoryType &memoryType = memoryTypeList[mtIdx];
470
471 log_info("Memory type index: %d\n", (uint32_t)memoryType);
472 log_info("Memory type property: %d\n",
473 memoryType.getMemoryTypeProperty());
474
475 VulkanBufferList vkBufferList(numBuffers, vkDevice, bufferSize,
476 vkExternalMemoryHandleType);
477
478 for (size_t bIdx = 0; bIdx < numBuffers; bIdx++)
479 {
480 vkBufferListDeviceMemory.push_back(
481 new VulkanDeviceMemory(vkDevice, bufferSize, memoryType,
482 vkExternalMemoryHandleType));
483 externalMemory.push_back(new clExternalMemory(
484 vkBufferListDeviceMemory[bIdx], vkExternalMemoryHandleType,
485 0, bufferSize, context, deviceId));
486 }
487 cl_mem buffers[4];
488 clFinish(cmd_queue1);
489 Params *params = (Params *)vkParamsDeviceMemory.map();
490 params->numBuffers = numBuffers;
491 params->bufferSize = bufferSize;
492 params->interBufferOffset = 0;
493 vkParamsDeviceMemory.unmap();
494 vkDescriptorSet.update(0, vkParamsBuffer);
495 for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++)
496 {
497 size_t buffer_size = vkBufferList[bIdx].getSize();
498 vkBufferListDeviceMemory[bIdx]->bindBuffer(vkBufferList[bIdx],
499 0);
500 buffers[bIdx] = externalMemory[bIdx]->getExternalMemoryBuffer();
501 vkDescriptorSet.update((uint32_t)bIdx + 1, vkBufferList[bIdx]);
502 }
503 vkCommandBuffer.begin();
504 vkCommandBuffer.bindPipeline(vkComputePipeline);
505 vkCommandBuffer.bindDescriptorSets(
506 vkComputePipeline, vkPipelineLayout, vkDescriptorSet);
507 vkCommandBuffer.dispatch(512, 1, 1);
508 vkCommandBuffer.end();
509
510 if (vkBufferList.size() == 1)
511 {
512 update_buffer_kernel = kernel[0];
513 }
514 else if (vkBufferList.size() == 2)
515 {
516 update_buffer_kernel = kernel[1];
517 }
518 else if (vkBufferList.size() == 4)
519 {
520 update_buffer_kernel = kernel[2];
521 }
522
523 // global work size should be less than or equal to
524 // bufferSizeList[i]
525 global_work_size[0] = bufferSize;
526
527 for (uint32_t iter = 0; iter < maxIter; iter++)
528 {
529 if (iter == 0)
530 {
531 vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore);
532 }
533 else
534 {
535 vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer,
536 vkVk2CLSemaphore);
537 }
538 clVk2CLExternalSemaphore->wait(cmd_queue1);
539
540 err = clSetKernelArg(update_buffer_kernel, 0, sizeof(uint32_t),
541 (void *)&bufferSize);
542 for (int i = 0; i < vkBufferList.size(); i++)
543 {
544 err |=
545 clSetKernelArg(update_buffer_kernel, i + 1,
546 sizeof(cl_mem), (void *)&(buffers[i]));
547 }
548
549 if (err != CL_SUCCESS)
550 {
551 print_error(err,
552 "Error: Failed to set arg values for kernel\n");
553 goto CLEANUP;
554 }
555 err = clEnqueueNDRangeKernel(cmd_queue1, update_buffer_kernel,
556 1, NULL, global_work_size, NULL, 0,
557 NULL, NULL);
558 if (err != CL_SUCCESS)
559 {
560 print_error(err,
561 "Error: Failed to launch update_buffer_kernel,"
562 " error\n");
563 goto CLEANUP;
564 }
565 if (iter != (maxIter - 1))
566 {
567 clCl2VkExternalSemaphore->signal(cmd_queue1);
568 }
569 }
570 error_2 = (uint8_t *)malloc(sizeof(uint8_t));
571 if (NULL == error_2)
572 {
573 log_error("Not able to allocate memory\n");
574 goto CLEANUP;
575 }
576
577 error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
578 sizeof(uint8_t), NULL, &err);
579 if (CL_SUCCESS != err)
580 {
581 print_error(err, "Error: clCreateBuffer \n");
582 goto CLEANUP;
583 }
584 uint8_t val = 0;
585 err = clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0,
586 sizeof(uint8_t), &val, 0, NULL, NULL);
587 if (CL_SUCCESS != err)
588 {
589 print_error(err, "Error: clEnqueueWriteBuffer \n");
590 goto CLEANUP;
591 }
592
593 int calc_max_iter = (maxIter * 2);
594 for (int i = 0; i < vkBufferList.size(); i++)
595 {
596 err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem),
597 (void *)&(buffers[i]));
598 err |=
599 clSetKernelArg(verify_kernel, 1, sizeof(int), &bufferSize);
600 err |= clSetKernelArg(verify_kernel, 2, sizeof(int),
601 &calc_max_iter);
602 err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem),
603 (void *)&error_1);
604 if (err != CL_SUCCESS)
605 {
606 print_error(
607 err,
608 "Error: Failed to set arg values for verify_kernel \n");
609 goto CLEANUP;
610 }
611 err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1, NULL,
612 global_work_size, NULL, 0, NULL,
613 NULL);
614 if (err != CL_SUCCESS)
615 {
616 print_error(
617 err, "Error: Failed to launch verify_kernel, error\n");
618 goto CLEANUP;
619 }
620
621 err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0,
622 sizeof(uint8_t), error_2, 0, NULL,
623 NULL);
624 if (err != CL_SUCCESS)
625 {
626 print_error(err, "Error: Failed read output, error \n");
627 goto CLEANUP;
628 }
629 if (*error_2 == 1)
630 {
631 log_error("&&&& vulkan_opencl_buffer test FAILED\n");
632 goto CLEANUP;
633 }
634 }
635 for (size_t i = 0; i < vkBufferList.size(); i++)
636 {
637 delete vkBufferListDeviceMemory[i];
638 delete externalMemory[i];
639 }
640 vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(),
641 vkBufferListDeviceMemory.begin()
642 + numBuffers);
643 externalMemory.erase(externalMemory.begin(),
644 externalMemory.begin() + numBuffers);
645 }
646 }
647 CLEANUP:
648 for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++)
649 {
650 if (vkBufferListDeviceMemory[i])
651 {
652 delete vkBufferListDeviceMemory[i];
653 }
654 if (externalMemory[i])
655 {
656 delete externalMemory[i];
657 }
658 }
659 if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore;
660 if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore;
661 if (error_2) free(error_2);
662 if (error_1) clReleaseMemObject(error_1);
663 return err;
664 }
665
run_test_with_multi_import_same_ctx(cl_context & context,cl_command_queue & cmd_queue1,cl_kernel * kernel,cl_kernel & verify_kernel,VulkanDevice & vkDevice,uint32_t numBuffers,uint32_t bufferSize,uint32_t bufferSizeForOffset)666 int run_test_with_multi_import_same_ctx(
667 cl_context &context, cl_command_queue &cmd_queue1, cl_kernel *kernel,
668 cl_kernel &verify_kernel, VulkanDevice &vkDevice, uint32_t numBuffers,
669 uint32_t bufferSize, uint32_t bufferSizeForOffset)
670 {
671 size_t global_work_size[1];
672 uint8_t *error_2;
673 cl_mem error_1;
674 int numImports = numBuffers;
675 cl_kernel update_buffer_kernel[MAX_IMPORTS];
676 clExternalSemaphore *clVk2CLExternalSemaphore = NULL;
677 clExternalSemaphore *clCl2VkExternalSemaphore = NULL;
678 int err = CL_SUCCESS;
679 int calc_max_iter;
680 bool withOffset;
681 uint32_t pBufferSize;
682
683 const std::vector<VulkanExternalMemoryHandleType>
684 vkExternalMemoryHandleTypeList =
685 getSupportedVulkanExternalMemoryHandleTypeList();
686 VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType =
687 getSupportedVulkanExternalSemaphoreHandleTypeList()[0];
688 VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType);
689 VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType);
690
691 VulkanQueue &vkQueue = vkDevice.getQueue();
692
693 std::vector<char> vkBufferShader = readFile("buffer.spv");
694
695 VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader);
696 VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList(
697 MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER);
698 VulkanDescriptorSetLayout vkDescriptorSetLayout(
699 vkDevice, vkDescriptorSetLayoutBindingList);
700 VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout);
701 VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout,
702 vkBufferShaderModule);
703
704 VulkanDescriptorPool vkDescriptorPool(vkDevice,
705 vkDescriptorSetLayoutBindingList);
706 VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool,
707 vkDescriptorSetLayout);
708
709 clVk2CLExternalSemaphore = new clExternalSemaphore(
710 vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
711 clCl2VkExternalSemaphore = new clExternalSemaphore(
712 vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
713 const uint32_t maxIter = innerIterations;
714 VulkanCommandPool vkCommandPool(vkDevice);
715 VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool);
716
717 VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params));
718 VulkanDeviceMemory vkParamsDeviceMemory(
719 vkDevice, vkParamsBuffer.getSize(),
720 getVulkanMemoryType(vkDevice,
721 VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT));
722 vkParamsDeviceMemory.bindBuffer(vkParamsBuffer);
723 std::vector<VulkanDeviceMemory *> vkBufferListDeviceMemory;
724 std::vector<std::vector<clExternalMemory *>> externalMemory;
725
726
727 for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size();
728 emhtIdx++)
729 {
730 VulkanExternalMemoryHandleType vkExternalMemoryHandleType =
731 vkExternalMemoryHandleTypeList[emhtIdx];
732 log_info("External memory handle type: %d\n",
733 vkExternalMemoryHandleType);
734
735 VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024,
736 vkExternalMemoryHandleType);
737 const VulkanMemoryTypeList &memoryTypeList =
738 vkDummyBuffer.getMemoryTypeList();
739
740 for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++)
741 {
742 const VulkanMemoryType &memoryType = memoryTypeList[mtIdx];
743
744 log_info("Memory type index: %d\n", (uint32_t)memoryType);
745 log_info("Memory type property: %d\n",
746 memoryType.getMemoryTypeProperty());
747 for (unsigned int withOffset = 0;
748 withOffset <= (unsigned int)enableOffset; withOffset++)
749 {
750 log_info("Running withOffset case %d\n", (uint32_t)withOffset);
751 if (withOffset)
752 {
753 pBufferSize = bufferSizeForOffset;
754 }
755 else
756 {
757 pBufferSize = bufferSize;
758 }
759 cl_mem buffers[MAX_BUFFERS][MAX_IMPORTS];
760 VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize,
761 vkExternalMemoryHandleType);
762 uint32_t interBufferOffset =
763 (uint32_t)(vkBufferList[0].getSize());
764
765 for (size_t bIdx = 0; bIdx < numBuffers; bIdx++)
766 {
767 if (withOffset == 0)
768 {
769 vkBufferListDeviceMemory.push_back(
770 new VulkanDeviceMemory(vkDevice, pBufferSize,
771 memoryType,
772 vkExternalMemoryHandleType));
773 }
774 if (withOffset == 1)
775 {
776 uint32_t totalSize =
777 (uint32_t)(vkBufferList.size() * interBufferOffset);
778 vkBufferListDeviceMemory.push_back(
779 new VulkanDeviceMemory(vkDevice, totalSize,
780 memoryType,
781 vkExternalMemoryHandleType));
782 }
783 std::vector<clExternalMemory *> pExternalMemory;
784 for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++)
785 {
786 pExternalMemory.push_back(new clExternalMemory(
787 vkBufferListDeviceMemory[bIdx],
788 vkExternalMemoryHandleType,
789 withOffset * bIdx * interBufferOffset, pBufferSize,
790 context, deviceId));
791 }
792 externalMemory.push_back(pExternalMemory);
793 }
794
795 clFinish(cmd_queue1);
796 Params *params = (Params *)vkParamsDeviceMemory.map();
797 params->numBuffers = numBuffers;
798 params->bufferSize = pBufferSize;
799 params->interBufferOffset = interBufferOffset * withOffset;
800 vkParamsDeviceMemory.unmap();
801 vkDescriptorSet.update(0, vkParamsBuffer);
802 for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++)
803 {
804 size_t buffer_size = vkBufferList[bIdx].getSize();
805 vkBufferListDeviceMemory[bIdx]->bindBuffer(
806 vkBufferList[bIdx],
807 bIdx * interBufferOffset * withOffset);
808 for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++)
809 {
810 buffers[bIdx][cl_bIdx] =
811 externalMemory[bIdx][cl_bIdx]
812 ->getExternalMemoryBuffer();
813 }
814 vkDescriptorSet.update((uint32_t)bIdx + 1,
815 vkBufferList[bIdx]);
816 }
817 vkCommandBuffer.begin();
818 vkCommandBuffer.bindPipeline(vkComputePipeline);
819 vkCommandBuffer.bindDescriptorSets(
820 vkComputePipeline, vkPipelineLayout, vkDescriptorSet);
821 vkCommandBuffer.dispatch(512, 1, 1);
822 vkCommandBuffer.end();
823 for (int i = 0; i < numImports; i++)
824 {
825 update_buffer_kernel[i] = (numBuffers == 1)
826 ? kernel[0]
827 : ((numBuffers == 2) ? kernel[1] : kernel[2]);
828 }
829 // global work size should be less than or equal to
830 // bufferSizeList[i]
831 global_work_size[0] = pBufferSize;
832
833 for (uint32_t iter = 0; iter < maxIter; iter++)
834 {
835 if (iter == 0)
836 {
837 vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore);
838 }
839 else
840 {
841 vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer,
842 vkVk2CLSemaphore);
843 }
844 clVk2CLExternalSemaphore->wait(cmd_queue1);
845 for (uint8_t launchIter = 0; launchIter < numImports;
846 launchIter++)
847 {
848 err = clSetKernelArg(update_buffer_kernel[launchIter],
849 0, sizeof(uint32_t),
850 (void *)&pBufferSize);
851 for (int i = 0; i < numBuffers; i++)
852 {
853 err |= clSetKernelArg(
854 update_buffer_kernel[launchIter], i + 1,
855 sizeof(cl_mem),
856 (void *)&(buffers[i][launchIter]));
857 }
858
859 if (err != CL_SUCCESS)
860 {
861 print_error(err,
862 "Error: Failed to set arg values for "
863 "kernel\n ");
864 goto CLEANUP;
865 }
866 err = clEnqueueNDRangeKernel(
867 cmd_queue1, update_buffer_kernel[launchIter], 1,
868 NULL, global_work_size, NULL, 0, NULL, NULL);
869 if (err != CL_SUCCESS)
870 {
871 print_error(err,
872 "Error: Failed to launch "
873 "update_buffer_kernel, error\n ");
874 goto CLEANUP;
875 }
876 }
877 if (iter != (maxIter - 1))
878 {
879 clCl2VkExternalSemaphore->signal(cmd_queue1);
880 }
881 }
882 error_2 = (uint8_t *)malloc(sizeof(uint8_t));
883 if (NULL == error_2)
884 {
885 log_error("Not able to allocate memory\n");
886 goto CLEANUP;
887 }
888
889 error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
890 sizeof(uint8_t), NULL, &err);
891 if (CL_SUCCESS != err)
892 {
893 print_error(err, "Error: clCreateBuffer \n");
894 goto CLEANUP;
895 }
896 uint8_t val = 0;
897 err =
898 clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0,
899 sizeof(uint8_t), &val, 0, NULL, NULL);
900 if (CL_SUCCESS != err)
901 {
902 print_error(err, "Error: clEnqueueWriteBuffer \n");
903 goto CLEANUP;
904 }
905 calc_max_iter = maxIter * (numBuffers + 1);
906
907 for (int i = 0; i < vkBufferList.size(); i++)
908 {
909 err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem),
910 (void *)&(buffers[i][0]));
911 err |= clSetKernelArg(verify_kernel, 1, sizeof(int),
912 &pBufferSize);
913 err |= clSetKernelArg(verify_kernel, 2, sizeof(int),
914 &calc_max_iter);
915 err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem),
916 (void *)&error_1);
917 if (err != CL_SUCCESS)
918 {
919 print_error(err,
920 "Error: Failed to set arg values for "
921 "verify_kernel \n");
922 goto CLEANUP;
923 }
924 err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1,
925 NULL, global_work_size, NULL,
926 0, NULL, NULL);
927 if (err != CL_SUCCESS)
928 {
929 print_error(
930 err,
931 "Error: Failed to launch verify_kernel, error\n");
932 goto CLEANUP;
933 }
934
935 err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0,
936 sizeof(uint8_t), error_2, 0, NULL,
937 NULL);
938 if (err != CL_SUCCESS)
939 {
940 print_error(err, "Error: Failed read output, error \n");
941 goto CLEANUP;
942 }
943 if (*error_2 == 1)
944 {
945 log_error("&&&& vulkan_opencl_buffer test FAILED\n");
946 goto CLEANUP;
947 }
948 }
949 for (size_t i = 0; i < vkBufferList.size(); i++)
950 {
951 for (size_t j = 0; j < numImports; j++)
952 {
953 delete externalMemory[i][j];
954 }
955 }
956 for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++)
957 {
958 delete vkBufferListDeviceMemory[i];
959 }
960 vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(),
961 vkBufferListDeviceMemory.end());
962 for (size_t i = 0; i < externalMemory.size(); i++)
963 {
964 externalMemory[i].erase(externalMemory[i].begin(),
965 externalMemory[i].begin()
966 + numBuffers);
967 }
968 externalMemory.clear();
969 }
970 }
971 }
972 CLEANUP:
973 for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++)
974 {
975 if (vkBufferListDeviceMemory[i])
976 {
977 delete vkBufferListDeviceMemory[i];
978 }
979 }
980 for (size_t i = 0; i < externalMemory.size(); i++)
981 {
982 for (size_t j = 0; j < externalMemory[i].size(); j++)
983 {
984 if (externalMemory[i][j])
985 {
986 delete externalMemory[i][j];
987 }
988 }
989 }
990 if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore;
991 if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore;
992 if (error_2) free(error_2);
993 if (error_1) clReleaseMemObject(error_1);
994 return err;
995 }
996
run_test_with_multi_import_diff_ctx(cl_context & context,cl_context & context2,cl_command_queue & cmd_queue1,cl_command_queue & cmd_queue2,cl_kernel * kernel1,cl_kernel * kernel2,cl_kernel & verify_kernel,cl_kernel verify_kernel2,VulkanDevice & vkDevice,uint32_t numBuffers,uint32_t bufferSize,uint32_t bufferSizeForOffset)997 int run_test_with_multi_import_diff_ctx(
998 cl_context &context, cl_context &context2, cl_command_queue &cmd_queue1,
999 cl_command_queue &cmd_queue2, cl_kernel *kernel1, cl_kernel *kernel2,
1000 cl_kernel &verify_kernel, cl_kernel verify_kernel2, VulkanDevice &vkDevice,
1001 uint32_t numBuffers, uint32_t bufferSize, uint32_t bufferSizeForOffset)
1002 {
1003 size_t global_work_size[1];
1004 uint8_t *error_3;
1005 cl_mem error_1;
1006 cl_mem error_2;
1007 int numImports = numBuffers;
1008 cl_kernel update_buffer_kernel1[MAX_IMPORTS];
1009 cl_kernel update_buffer_kernel2[MAX_IMPORTS];
1010 clExternalSemaphore *clVk2CLExternalSemaphore = NULL;
1011 clExternalSemaphore *clCl2VkExternalSemaphore = NULL;
1012 clExternalSemaphore *clVk2CLExternalSemaphore2 = NULL;
1013 clExternalSemaphore *clCl2VkExternalSemaphore2 = NULL;
1014 int err = CL_SUCCESS;
1015 int calc_max_iter;
1016 bool withOffset;
1017 uint32_t pBufferSize;
1018
1019 const std::vector<VulkanExternalMemoryHandleType>
1020 vkExternalMemoryHandleTypeList =
1021 getSupportedVulkanExternalMemoryHandleTypeList();
1022 VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType =
1023 getSupportedVulkanExternalSemaphoreHandleTypeList()[0];
1024 VulkanSemaphore vkVk2CLSemaphore(vkDevice, vkExternalSemaphoreHandleType);
1025 VulkanSemaphore vkCl2VkSemaphore(vkDevice, vkExternalSemaphoreHandleType);
1026
1027 VulkanQueue &vkQueue = vkDevice.getQueue();
1028
1029 std::vector<char> vkBufferShader = readFile("buffer.spv");
1030
1031 VulkanShaderModule vkBufferShaderModule(vkDevice, vkBufferShader);
1032 VulkanDescriptorSetLayoutBindingList vkDescriptorSetLayoutBindingList(
1033 MAX_BUFFERS + 1, VULKAN_DESCRIPTOR_TYPE_STORAGE_BUFFER);
1034 VulkanDescriptorSetLayout vkDescriptorSetLayout(
1035 vkDevice, vkDescriptorSetLayoutBindingList);
1036 VulkanPipelineLayout vkPipelineLayout(vkDevice, vkDescriptorSetLayout);
1037 VulkanComputePipeline vkComputePipeline(vkDevice, vkPipelineLayout,
1038 vkBufferShaderModule);
1039
1040 VulkanDescriptorPool vkDescriptorPool(vkDevice,
1041 vkDescriptorSetLayoutBindingList);
1042 VulkanDescriptorSet vkDescriptorSet(vkDevice, vkDescriptorPool,
1043 vkDescriptorSetLayout);
1044
1045 clVk2CLExternalSemaphore = new clExternalSemaphore(
1046 vkVk2CLSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
1047 clCl2VkExternalSemaphore = new clExternalSemaphore(
1048 vkCl2VkSemaphore, context, vkExternalSemaphoreHandleType, deviceId);
1049
1050 clVk2CLExternalSemaphore2 = new clExternalSemaphore(
1051 vkVk2CLSemaphore, context2, vkExternalSemaphoreHandleType, deviceId);
1052 clCl2VkExternalSemaphore2 = new clExternalSemaphore(
1053 vkCl2VkSemaphore, context2, vkExternalSemaphoreHandleType, deviceId);
1054
1055 const uint32_t maxIter = innerIterations;
1056 VulkanCommandPool vkCommandPool(vkDevice);
1057 VulkanCommandBuffer vkCommandBuffer(vkDevice, vkCommandPool);
1058
1059 VulkanBuffer vkParamsBuffer(vkDevice, sizeof(Params));
1060 VulkanDeviceMemory vkParamsDeviceMemory(
1061 vkDevice, vkParamsBuffer.getSize(),
1062 getVulkanMemoryType(vkDevice,
1063 VULKAN_MEMORY_TYPE_PROPERTY_HOST_VISIBLE_COHERENT));
1064 vkParamsDeviceMemory.bindBuffer(vkParamsBuffer);
1065 std::vector<VulkanDeviceMemory *> vkBufferListDeviceMemory;
1066 std::vector<std::vector<clExternalMemory *>> externalMemory1;
1067 std::vector<std::vector<clExternalMemory *>> externalMemory2;
1068
1069 for (size_t emhtIdx = 0; emhtIdx < vkExternalMemoryHandleTypeList.size();
1070 emhtIdx++)
1071 {
1072 VulkanExternalMemoryHandleType vkExternalMemoryHandleType =
1073 vkExternalMemoryHandleTypeList[emhtIdx];
1074 log_info("External memory handle type:%d\n",
1075 vkExternalMemoryHandleType);
1076
1077 VulkanBuffer vkDummyBuffer(vkDevice, 4 * 1024,
1078 vkExternalMemoryHandleType);
1079 const VulkanMemoryTypeList &memoryTypeList =
1080 vkDummyBuffer.getMemoryTypeList();
1081
1082 for (size_t mtIdx = 0; mtIdx < memoryTypeList.size(); mtIdx++)
1083 {
1084 const VulkanMemoryType &memoryType = memoryTypeList[mtIdx];
1085
1086 log_info("Memory type index: %d\n", (uint32_t)memoryType);
1087 log_info("Memory type property: %d\n",
1088 memoryType.getMemoryTypeProperty());
1089
1090 for (unsigned int withOffset = 0;
1091 withOffset <= (unsigned int)enableOffset; withOffset++)
1092 {
1093 log_info("Running withOffset case %d\n", (uint32_t)withOffset);
1094 cl_mem buffers1[MAX_BUFFERS][MAX_IMPORTS];
1095 cl_mem buffers2[MAX_BUFFERS][MAX_IMPORTS];
1096 if (withOffset)
1097 {
1098 pBufferSize = bufferSizeForOffset;
1099 }
1100 else
1101 {
1102 pBufferSize = bufferSize;
1103 }
1104 VulkanBufferList vkBufferList(numBuffers, vkDevice, pBufferSize,
1105 vkExternalMemoryHandleType);
1106 uint32_t interBufferOffset =
1107 (uint32_t)(vkBufferList[0].getSize());
1108
1109 for (size_t bIdx = 0; bIdx < numBuffers; bIdx++)
1110 {
1111 if (withOffset == 0)
1112 {
1113 vkBufferListDeviceMemory.push_back(
1114 new VulkanDeviceMemory(vkDevice, pBufferSize,
1115 memoryType,
1116 vkExternalMemoryHandleType));
1117 }
1118 if (withOffset == 1)
1119 {
1120 uint32_t totalSize =
1121 (uint32_t)(vkBufferList.size() * interBufferOffset);
1122 vkBufferListDeviceMemory.push_back(
1123 new VulkanDeviceMemory(vkDevice, totalSize,
1124 memoryType,
1125 vkExternalMemoryHandleType));
1126 }
1127 std::vector<clExternalMemory *> pExternalMemory1;
1128 std::vector<clExternalMemory *> pExternalMemory2;
1129 for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++)
1130 {
1131 pExternalMemory1.push_back(new clExternalMemory(
1132 vkBufferListDeviceMemory[bIdx],
1133 vkExternalMemoryHandleType,
1134 withOffset * bIdx * interBufferOffset, pBufferSize,
1135 context, deviceId));
1136 pExternalMemory2.push_back(new clExternalMemory(
1137 vkBufferListDeviceMemory[bIdx],
1138 vkExternalMemoryHandleType,
1139 withOffset * bIdx * interBufferOffset, pBufferSize,
1140 context2, deviceId));
1141 }
1142 externalMemory1.push_back(pExternalMemory1);
1143 externalMemory2.push_back(pExternalMemory2);
1144 }
1145
1146 clFinish(cmd_queue1);
1147 Params *params = (Params *)vkParamsDeviceMemory.map();
1148 params->numBuffers = numBuffers;
1149 params->bufferSize = pBufferSize;
1150 params->interBufferOffset = interBufferOffset * withOffset;
1151 vkParamsDeviceMemory.unmap();
1152 vkDescriptorSet.update(0, vkParamsBuffer);
1153 for (size_t bIdx = 0; bIdx < vkBufferList.size(); bIdx++)
1154 {
1155 size_t buffer_size = vkBufferList[bIdx].getSize();
1156 vkBufferListDeviceMemory[bIdx]->bindBuffer(
1157 vkBufferList[bIdx],
1158 bIdx * interBufferOffset * withOffset);
1159 for (size_t cl_bIdx = 0; cl_bIdx < numImports; cl_bIdx++)
1160 {
1161 buffers1[bIdx][cl_bIdx] =
1162 externalMemory1[bIdx][cl_bIdx]
1163 ->getExternalMemoryBuffer();
1164 buffers2[bIdx][cl_bIdx] =
1165 externalMemory2[bIdx][cl_bIdx]
1166 ->getExternalMemoryBuffer();
1167 }
1168 vkDescriptorSet.update((uint32_t)bIdx + 1,
1169 vkBufferList[bIdx]);
1170 }
1171
1172 vkCommandBuffer.begin();
1173 vkCommandBuffer.bindPipeline(vkComputePipeline);
1174 vkCommandBuffer.bindDescriptorSets(
1175 vkComputePipeline, vkPipelineLayout, vkDescriptorSet);
1176 vkCommandBuffer.dispatch(512, 1, 1);
1177 vkCommandBuffer.end();
1178
1179 for (int i = 0; i < numImports; i++)
1180 {
1181 update_buffer_kernel1[i] = (numBuffers == 1)
1182 ? kernel1[0]
1183 : ((numBuffers == 2) ? kernel1[1] : kernel1[2]);
1184 update_buffer_kernel2[i] = (numBuffers == 1)
1185 ? kernel2[0]
1186 : ((numBuffers == 2) ? kernel2[1] : kernel2[2]);
1187 }
1188
1189 // global work size should be less than or equal
1190 // to bufferSizeList[i]
1191 global_work_size[0] = pBufferSize;
1192
1193 for (uint32_t iter = 0; iter < maxIter; iter++)
1194 {
1195 if (iter == 0)
1196 {
1197 vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore);
1198 }
1199 else
1200 {
1201 vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer,
1202 vkVk2CLSemaphore);
1203 }
1204 clVk2CLExternalSemaphore->wait(cmd_queue1);
1205
1206 for (uint8_t launchIter = 0; launchIter < numImports;
1207 launchIter++)
1208 {
1209 err = clSetKernelArg(update_buffer_kernel1[launchIter],
1210 0, sizeof(uint32_t),
1211 (void *)&pBufferSize);
1212 for (int i = 0; i < numBuffers; i++)
1213 {
1214 err |= clSetKernelArg(
1215 update_buffer_kernel1[launchIter], i + 1,
1216 sizeof(cl_mem),
1217 (void *)&(buffers1[i][launchIter]));
1218 }
1219
1220 if (err != CL_SUCCESS)
1221 {
1222 print_error(err,
1223 "Error: Failed to set arg values for "
1224 "kernel\n ");
1225 goto CLEANUP;
1226 }
1227 err = clEnqueueNDRangeKernel(
1228 cmd_queue1, update_buffer_kernel1[launchIter], 1,
1229 NULL, global_work_size, NULL, 0, NULL, NULL);
1230 if (err != CL_SUCCESS)
1231 {
1232 print_error(err,
1233 "Error: Failed to launch "
1234 "update_buffer_kernel, error\n");
1235 goto CLEANUP;
1236 }
1237 }
1238 if (iter != (maxIter - 1))
1239 {
1240 clCl2VkExternalSemaphore->signal(cmd_queue1);
1241 }
1242 }
1243 clFinish(cmd_queue1);
1244 for (uint32_t iter = 0; iter < maxIter; iter++)
1245 {
1246 if (iter == 0)
1247 {
1248 vkQueue.submit(vkCommandBuffer, vkVk2CLSemaphore);
1249 }
1250 else
1251 {
1252 vkQueue.submit(vkCl2VkSemaphore, vkCommandBuffer,
1253 vkVk2CLSemaphore);
1254 }
1255 clVk2CLExternalSemaphore2->wait(cmd_queue2);
1256
1257 for (uint8_t launchIter = 0; launchIter < numImports;
1258 launchIter++)
1259 {
1260 err = clSetKernelArg(update_buffer_kernel2[launchIter],
1261 0, sizeof(uint32_t),
1262 (void *)&bufferSize);
1263 for (int i = 0; i < numBuffers; i++)
1264 {
1265 err |= clSetKernelArg(
1266 update_buffer_kernel2[launchIter], i + 1,
1267 sizeof(cl_mem),
1268 (void *)&(buffers2[i][launchIter]));
1269 }
1270
1271 if (err != CL_SUCCESS)
1272 {
1273 print_error(err,
1274 "Error: Failed to set arg values for "
1275 "kernel\n ");
1276 goto CLEANUP;
1277 }
1278 err = clEnqueueNDRangeKernel(
1279 cmd_queue2, update_buffer_kernel2[launchIter], 1,
1280 NULL, global_work_size, NULL, 0, NULL, NULL);
1281 if (err != CL_SUCCESS)
1282 {
1283 print_error(err,
1284 "Error: Failed to launch "
1285 "update_buffer_kernel, error\n ");
1286 goto CLEANUP;
1287 }
1288 }
1289 if (iter != (maxIter - 1))
1290 {
1291 clCl2VkExternalSemaphore2->signal(cmd_queue2);
1292 }
1293 }
1294 clFinish(cmd_queue2);
1295 error_3 = (uint8_t *)malloc(sizeof(uint8_t));
1296 if (NULL == error_3)
1297 {
1298 log_error("Not able to allocate memory\n");
1299 goto CLEANUP;
1300 }
1301
1302 error_1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
1303 sizeof(uint8_t), NULL, &err);
1304 if (CL_SUCCESS != err)
1305 {
1306 print_error(err, "Error: clCreateBuffer \n");
1307 goto CLEANUP;
1308 }
1309 error_2 = clCreateBuffer(context2, CL_MEM_WRITE_ONLY,
1310 sizeof(uint8_t), NULL, &err);
1311 if (CL_SUCCESS != err)
1312 {
1313 print_error(err, "Error: clCreateBuffer \n");
1314 goto CLEANUP;
1315 }
1316 uint8_t val = 0;
1317 err =
1318 clEnqueueWriteBuffer(cmd_queue1, error_1, CL_TRUE, 0,
1319 sizeof(uint8_t), &val, 0, NULL, NULL);
1320 if (err != CL_SUCCESS)
1321 {
1322 print_error(err, "Error: Failed read output, error \n");
1323 goto CLEANUP;
1324 }
1325
1326 err =
1327 clEnqueueWriteBuffer(cmd_queue2, error_2, CL_TRUE, 0,
1328 sizeof(uint8_t), &val, 0, NULL, NULL);
1329 if (err != CL_SUCCESS)
1330 {
1331 print_error(err, "Error: Failed read output, error \n");
1332 goto CLEANUP;
1333 }
1334
1335 calc_max_iter = maxIter * 2 * (numBuffers + 1);
1336 for (int i = 0; i < numBuffers; i++)
1337 {
1338 err = clSetKernelArg(verify_kernel, 0, sizeof(cl_mem),
1339 (void *)&(buffers1[i][0]));
1340 err |= clSetKernelArg(verify_kernel, 1, sizeof(int),
1341 &pBufferSize);
1342 err |= clSetKernelArg(verify_kernel, 2, sizeof(int),
1343 &calc_max_iter);
1344 err |= clSetKernelArg(verify_kernel, 3, sizeof(cl_mem),
1345 (void *)&error_1);
1346 if (err != CL_SUCCESS)
1347 {
1348 print_error(err,
1349 "Error: Failed to set arg values for "
1350 "verify_kernel \n");
1351 goto CLEANUP;
1352 }
1353 err = clEnqueueNDRangeKernel(cmd_queue1, verify_kernel, 1,
1354 NULL, global_work_size, NULL,
1355 0, NULL, NULL);
1356 if (err != CL_SUCCESS)
1357 {
1358 print_error(err,
1359 "Error: Failed to launch verify_kernel,"
1360 "error\n");
1361 goto CLEANUP;
1362 }
1363
1364 err = clEnqueueReadBuffer(cmd_queue1, error_1, CL_TRUE, 0,
1365 sizeof(uint8_t), error_3, 0, NULL,
1366 NULL);
1367 if (err != CL_SUCCESS)
1368 {
1369 print_error(err, "Error: Failed read output, error\n");
1370 goto CLEANUP;
1371 }
1372 if (*error_3 == 1)
1373 {
1374 log_error("&&&& vulkan_opencl_buffer test FAILED\n");
1375 goto CLEANUP;
1376 }
1377 }
1378 *error_3 = 0;
1379 for (int i = 0; i < vkBufferList.size(); i++)
1380 {
1381 err = clSetKernelArg(verify_kernel2, 0, sizeof(cl_mem),
1382 (void *)&(buffers2[i][0]));
1383 err |= clSetKernelArg(verify_kernel2, 1, sizeof(int),
1384 &pBufferSize);
1385 err |= clSetKernelArg(verify_kernel2, 2, sizeof(int),
1386 &calc_max_iter);
1387 err |= clSetKernelArg(verify_kernel2, 3, sizeof(cl_mem),
1388 (void *)&error_2);
1389 if (err != CL_SUCCESS)
1390 {
1391 print_error(err,
1392 "Error: Failed to set arg values for "
1393 "verify_kernel \n");
1394 goto CLEANUP;
1395 }
1396 err = clEnqueueNDRangeKernel(cmd_queue2, verify_kernel2, 1,
1397 NULL, global_work_size, NULL,
1398 0, NULL, NULL);
1399 if (err != CL_SUCCESS)
1400 {
1401 print_error(err,
1402 "Error: Failed to launch verify_kernel,"
1403 "error\n");
1404 goto CLEANUP;
1405 }
1406
1407 err = clEnqueueReadBuffer(cmd_queue2, error_2, CL_TRUE, 0,
1408 sizeof(uint8_t), error_3, 0, NULL,
1409 NULL);
1410 if (err != CL_SUCCESS)
1411 {
1412 print_error(err, "Error: Failed read output, error\n");
1413 goto CLEANUP;
1414 }
1415 if (*error_3 == 1)
1416 {
1417 log_error("&&&& vulkan_opencl_buffer test FAILED\n");
1418 goto CLEANUP;
1419 }
1420 }
1421 for (size_t i = 0; i < vkBufferList.size(); i++)
1422 {
1423 for (size_t j = 0; j < numImports; j++)
1424 {
1425 delete externalMemory1[i][j];
1426 delete externalMemory2[i][j];
1427 }
1428 }
1429 for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++)
1430 {
1431 delete vkBufferListDeviceMemory[i];
1432 }
1433 vkBufferListDeviceMemory.erase(vkBufferListDeviceMemory.begin(),
1434 vkBufferListDeviceMemory.end());
1435 for (size_t i = 0; i < externalMemory1.size(); i++)
1436 {
1437 externalMemory1[i].erase(externalMemory1[i].begin(),
1438 externalMemory1[i].begin()
1439 + numBuffers);
1440 externalMemory2[i].erase(externalMemory2[i].begin(),
1441 externalMemory2[i].begin()
1442 + numBuffers);
1443 }
1444 externalMemory1.clear();
1445 externalMemory2.clear();
1446 }
1447 }
1448 }
1449 CLEANUP:
1450 for (size_t i = 0; i < vkBufferListDeviceMemory.size(); i++)
1451 {
1452 if (vkBufferListDeviceMemory[i])
1453 {
1454 delete vkBufferListDeviceMemory[i];
1455 }
1456 }
1457 for (size_t i = 0; i < externalMemory1.size(); i++)
1458 {
1459 for (size_t j = 0; j < externalMemory1[i].size(); j++)
1460 {
1461 if (externalMemory1[i][j])
1462 {
1463 delete externalMemory1[i][j];
1464 }
1465 }
1466 }
1467 for (size_t i = 0; i < externalMemory2.size(); i++)
1468 {
1469 for (size_t j = 0; j < externalMemory2[i].size(); j++)
1470 {
1471 if (externalMemory2[i][j])
1472 {
1473 delete externalMemory2[i][j];
1474 }
1475 }
1476 }
1477 if (clVk2CLExternalSemaphore) delete clVk2CLExternalSemaphore;
1478 if (clCl2VkExternalSemaphore) delete clCl2VkExternalSemaphore;
1479 if (clVk2CLExternalSemaphore2) delete clVk2CLExternalSemaphore2;
1480 if (clCl2VkExternalSemaphore2) delete clCl2VkExternalSemaphore2;
1481 if (error_3) free(error_3);
1482 if (error_1) clReleaseMemObject(error_1);
1483 if (error_2) clReleaseMemObject(error_2);
1484 return err;
1485 }
1486
test_buffer_common(cl_device_id device_,cl_context context_,cl_command_queue queue_,int numElements_)1487 int test_buffer_common(cl_device_id device_, cl_context context_,
1488 cl_command_queue queue_, int numElements_)
1489 {
1490
1491 int current_device = 0;
1492 int device_count = 0;
1493 int devices_prohibited = 0;
1494 cl_int errNum = CL_SUCCESS;
1495 cl_platform_id platform = NULL;
1496 size_t extensionSize = 0;
1497 cl_uint num_devices = 0;
1498 cl_uint device_no = 0;
1499 const size_t bufsize = BUFFERSIZE;
1500 char buf[BUFFERSIZE];
1501 cl_device_id *devices;
1502 char *extensions = NULL;
1503 cl_kernel verify_kernel;
1504 cl_kernel verify_kernel2;
1505 cl_kernel kernel[3] = { NULL, NULL, NULL };
1506 cl_kernel kernel2[3] = { NULL, NULL, NULL };
1507 const char *program_source_const[3] = { kernel_text_numbuffer_1,
1508 kernel_text_numbuffer_2,
1509 kernel_text_numbuffer_4 };
1510 const char *program_source_const_verify;
1511 size_t program_source_length;
1512 cl_command_queue cmd_queue1 = NULL;
1513 cl_command_queue cmd_queue2 = NULL;
1514 cl_command_queue cmd_queue3 = NULL;
1515 cl_context context = NULL;
1516 cl_program program[3] = { NULL, NULL, NULL };
1517 cl_program program_verify, program_verify2;
1518 cl_context context2 = NULL;
1519
1520
1521 VulkanDevice vkDevice;
1522 uint32_t numBuffersList[] = { 1, 2, 4 };
1523 uint32_t bufferSizeList[] = { 4 * 1024, 64 * 1024, 2 * 1024 * 1024 };
1524 uint32_t bufferSizeListforOffset[] = { 256, 512, 1024 };
1525
1526 cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, 0, 0 };
1527 errNum = clGetPlatformIDs(1, &platform, NULL);
1528 if (errNum != CL_SUCCESS)
1529 {
1530 print_error(errNum, "Error: Failed to get platform\n");
1531 goto CLEANUP;
1532 }
1533
1534 errNum =
1535 clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices);
1536 if (CL_SUCCESS != errNum)
1537 {
1538 print_error(errNum, "clGetDeviceIDs failed in returning of devices\n");
1539 goto CLEANUP;
1540 }
1541 devices = (cl_device_id *)malloc(num_devices * sizeof(cl_device_id));
1542 if (NULL == devices)
1543 {
1544 errNum = CL_OUT_OF_HOST_MEMORY;
1545 print_error(errNum, "Unable to allocate memory for devices\n");
1546 goto CLEANUP;
1547 }
1548 errNum = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices,
1549 NULL);
1550 if (CL_SUCCESS != errNum)
1551 {
1552 print_error(errNum, "Failed to get deviceID.\n");
1553 goto CLEANUP;
1554 }
1555 contextProperties[1] = (cl_context_properties)platform;
1556 log_info("Assigned contextproperties for platform\n");
1557 for (device_no = 0; device_no < num_devices; device_no++)
1558 {
1559 errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS, 0,
1560 NULL, &extensionSize);
1561 if (CL_SUCCESS != errNum)
1562 {
1563 print_error(errNum,
1564 "Error in clGetDeviceInfo for getting device_extension "
1565 "size....\n");
1566 goto CLEANUP;
1567 }
1568 extensions = (char *)malloc(extensionSize);
1569 if (NULL == extensions)
1570 {
1571 print_error(errNum, "Unable to allocate memory for extensions\n");
1572 errNum = CL_OUT_OF_HOST_MEMORY;
1573 goto CLEANUP;
1574 }
1575 errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_EXTENSIONS,
1576 extensionSize, extensions, NULL);
1577 if (CL_SUCCESS != errNum)
1578 {
1579 print_error(errNum,
1580 "Error in clGetDeviceInfo for device_extension\n");
1581 goto CLEANUP;
1582 }
1583 errNum = clGetDeviceInfo(devices[device_no], CL_DEVICE_UUID_KHR,
1584 CL_UUID_SIZE_KHR, uuid, &extensionSize);
1585 if (CL_SUCCESS != errNum)
1586 {
1587 print_error(errNum, "clGetDeviceInfo failed\n");
1588 goto CLEANUP;
1589 }
1590 errNum =
1591 memcmp(uuid, vkDevice.getPhysicalDevice().getUUID(), VK_UUID_SIZE);
1592 if (errNum == 0)
1593 {
1594 break;
1595 }
1596 }
1597 if (device_no >= num_devices)
1598 {
1599 errNum = EXIT_FAILURE;
1600 print_error(errNum,
1601 "OpenCL error: "
1602 "No Vulkan-OpenCL Interop capable GPU found.\n");
1603 goto CLEANUP;
1604 }
1605 deviceId = devices[device_no];
1606 context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
1607 NULL, NULL, &errNum);
1608 if (CL_SUCCESS != errNum)
1609 {
1610 print_error(errNum, "error creating context\n");
1611 goto CLEANUP;
1612 }
1613 log_info("Successfully created context !!!\n");
1614
1615 cmd_queue1 = clCreateCommandQueue(context, devices[device_no], 0, &errNum);
1616 if (CL_SUCCESS != errNum)
1617 {
1618 errNum = CL_INVALID_COMMAND_QUEUE;
1619 print_error(errNum, "Error: Failed to create command queue!\n");
1620 goto CLEANUP;
1621 }
1622 cmd_queue2 = clCreateCommandQueue(context, devices[device_no], 0, &errNum);
1623 if (CL_SUCCESS != errNum)
1624 {
1625 errNum = CL_INVALID_COMMAND_QUEUE;
1626 print_error(errNum, "Error: Failed to create command queue!\n");
1627 goto CLEANUP;
1628 }
1629 log_info("clCreateCommandQueue successful\n");
1630 for (int i = 0; i < 3; i++)
1631 {
1632 program_source_length = strlen(program_source_const[i]);
1633 program[i] =
1634 clCreateProgramWithSource(context, 1, &program_source_const[i],
1635 &program_source_length, &errNum);
1636 errNum = clBuildProgram(program[i], 0, NULL, NULL, NULL, NULL);
1637 if (errNum != CL_SUCCESS)
1638 {
1639 print_error(errNum, "Error: Failed to build program \n");
1640 return errNum;
1641 }
1642 // create the kernel
1643 kernel[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum);
1644 if (errNum != CL_SUCCESS)
1645 {
1646 print_error(errNum, "clCreateKernel failed \n");
1647 return errNum;
1648 }
1649 }
1650
1651 program_source_const_verify = kernel_text_verify;
1652 program_source_length = strlen(program_source_const_verify);
1653 program_verify =
1654 clCreateProgramWithSource(context, 1, &program_source_const_verify,
1655 &program_source_length, &errNum);
1656 errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL);
1657 if (errNum != CL_SUCCESS)
1658 {
1659 log_error("Error: Failed to build program2\n");
1660 return errNum;
1661 }
1662 verify_kernel = clCreateKernel(program_verify, "checkKernel", &errNum);
1663 if (errNum != CL_SUCCESS)
1664 {
1665 print_error(errNum, "clCreateKernel failed \n");
1666 return errNum;
1667 }
1668
1669 if (multiCtx) // different context guard
1670 {
1671 context2 = clCreateContextFromType(
1672 contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum);
1673 if (CL_SUCCESS != errNum)
1674 {
1675 print_error(errNum, "error creating context\n");
1676 goto CLEANUP;
1677 }
1678 cmd_queue3 =
1679 clCreateCommandQueue(context2, devices[device_no], 0, &errNum);
1680 if (CL_SUCCESS != errNum)
1681 {
1682 errNum = CL_INVALID_COMMAND_QUEUE;
1683 print_error(errNum, "Error: Failed to create command queue!\n");
1684 goto CLEANUP;
1685 }
1686 for (int i = 0; i < 3; i++)
1687 {
1688 program_source_length = strlen(program_source_const[i]);
1689 program[i] =
1690 clCreateProgramWithSource(context2, 1, &program_source_const[i],
1691 &program_source_length, &errNum);
1692 errNum = clBuildProgram(program[i], 0, NULL, NULL, NULL, NULL);
1693 if (errNum != CL_SUCCESS)
1694 {
1695 print_error(errNum, "Error: Failed to build program \n");
1696 return errNum;
1697 }
1698 // create the kernel
1699 kernel2[i] = clCreateKernel(program[i], "clUpdateBuffer", &errNum);
1700 if (errNum != CL_SUCCESS)
1701 {
1702 print_error(errNum, "clCreateKernel failed \n");
1703 return errNum;
1704 }
1705 }
1706 program_source_length = strlen(program_source_const_verify);
1707 program_verify =
1708 clCreateProgramWithSource(context2, 1, &program_source_const_verify,
1709 &program_source_length, &errNum);
1710 errNum = clBuildProgram(program_verify, 0, NULL, NULL, NULL, NULL);
1711 if (errNum != CL_SUCCESS)
1712 {
1713 log_error("Error: Failed to build program2\n");
1714 return errNum;
1715 }
1716 verify_kernel2 = clCreateKernel(program_verify, "checkKernel", &errNum);
1717 if (errNum != CL_SUCCESS)
1718 {
1719 print_error(errNum, "clCreateKernel failed \n");
1720 return errNum;
1721 }
1722 }
1723
1724 for (size_t numBuffersIdx = 0; numBuffersIdx < ARRAY_SIZE(numBuffersList);
1725 numBuffersIdx++)
1726 {
1727 uint32_t numBuffers = numBuffersList[numBuffersIdx];
1728 log_info("Number of buffers: %d\n", numBuffers);
1729 for (size_t sizeIdx = 0; sizeIdx < ARRAY_SIZE(bufferSizeList);
1730 sizeIdx++)
1731 {
1732 uint32_t bufferSize = bufferSizeList[sizeIdx];
1733 uint32_t bufferSizeForOffset = bufferSizeListforOffset[sizeIdx];
1734 log_info("&&&& RUNNING vulkan_opencl_buffer test for Buffer size: "
1735 "%d\n",
1736 bufferSize);
1737 if (multiImport && !multiCtx)
1738 {
1739 errNum = run_test_with_multi_import_same_ctx(
1740 context, cmd_queue1, kernel, verify_kernel, vkDevice,
1741 numBuffers, bufferSize, bufferSizeForOffset);
1742 }
1743 else if (multiImport && multiCtx)
1744 {
1745 errNum = run_test_with_multi_import_diff_ctx(
1746 context, context2, cmd_queue1, cmd_queue3, kernel, kernel2,
1747 verify_kernel, verify_kernel2, vkDevice, numBuffers,
1748 bufferSize, bufferSizeForOffset);
1749 }
1750 else if (numCQ == 2)
1751 {
1752 errNum = run_test_with_two_queue(
1753 context, cmd_queue1, cmd_queue2, kernel, verify_kernel,
1754 vkDevice, numBuffers + 1, bufferSize);
1755 }
1756 else
1757 {
1758 errNum = run_test_with_one_queue(context, cmd_queue1, kernel,
1759 verify_kernel, vkDevice,
1760 numBuffers, bufferSize);
1761 }
1762 if (errNum != CL_SUCCESS)
1763 {
1764 print_error(errNum, "func_name failed \n");
1765 goto CLEANUP;
1766 }
1767 }
1768 }
1769
1770 CLEANUP:
1771 for (int i = 0; i < 3; i++)
1772 {
1773 if (program[i]) clReleaseProgram(program[i]);
1774 if (kernel[i]) clReleaseKernel(kernel[i]);
1775 }
1776 if (cmd_queue1) clReleaseCommandQueue(cmd_queue1);
1777 if (cmd_queue2) clReleaseCommandQueue(cmd_queue2);
1778 if (cmd_queue3) clReleaseCommandQueue(cmd_queue3);
1779 if (context) clReleaseContext(context);
1780 if (context2) clReleaseContext(context2);
1781
1782 if (devices) free(devices);
1783 if (extensions) free(extensions);
1784
1785 return errNum;
1786 }
1787