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 "harness/compat.h"
17
18 #include <assert.h>
19 #include <iomanip>
20 #include <iostream>
21 #include <sstream>
22 #include <stdio.h>
23 #include <string.h>
24 #include <string>
25 #include <sys/stat.h>
26 #include <sys/types.h>
27
28 #include "procs.h"
29 #include "harness/errorHelpers.h"
30
31 #define STRING_LENGTH 1024
32
createKernelSourceCode(std::stringstream & stream,int num_pipes)33 void createKernelSourceCode(std::stringstream &stream, int num_pipes)
34 {
35 int i;
36
37 stream << "__kernel void test_multiple_pipe_write(__global int *src, ";
38 for (i = 0; i < num_pipes; i++)
39 {
40 stream << "__write_only pipe int pipe" << i << ", ";
41 }
42 stream << R"(int num_pipes )
43 {
44 int gid = get_global_id(0);
45 reserve_id_t res_id;
46
47
48 if(gid < (get_global_size(0))/num_pipes)
49 {
50 res_id = reserve_write_pipe(pipe0, 1);
51 if(is_valid_reserve_id(res_id))
52 {
53 write_pipe(pipe0, res_id, 0, &src[gid]);
54 commit_write_pipe(pipe0, res_id);
55 }
56 })";
57
58 for (i = 1; i < num_pipes; i++)
59 {
60 // clang-format off
61 stream << R"(
62 else if(gid < ()" << (i + 1) << R"(*get_global_size(0))/num_pipes)
63 {
64 res_id = reserve_write_pipe(pipe)" << i << R"(, 1);
65 if(is_valid_reserve_id(res_id))
66 {
67 write_pipe(pipe)" << i << R"(, res_id, 0, &src[gid]);
68 commit_write_pipe(pipe)" << i << R"(, res_id);
69 }
70 }
71 )";
72 // clang-format om
73 }
74 stream << R"(
75 }
76
77 __kernel void test_multiple_pipe_read(__global int *dst, )";
78
79 for (i = 0; i < num_pipes; i++)
80 {
81 stream << "__read_only pipe int pipe" << i << ", ";
82 }
83 stream << R"(int num_pipes )
84 {
85 int gid = get_global_id(0);
86 reserve_id_t res_id;
87
88
89 if(gid < (get_global_size(0))/num_pipes)
90 {
91 res_id = reserve_read_pipe(pipe0, 1);
92 if(is_valid_reserve_id(res_id))
93 {
94 read_pipe(pipe0, res_id, 0, &dst[gid]);
95 commit_read_pipe(pipe0, res_id);
96 }
97 })";
98
99 for (i = 1; i < num_pipes; i++)
100 {
101 // clang-format off
102 stream << R"(
103 else if(gid < ()" << (i + 1) << R"(*get_global_size(0))/num_pipes)
104 {
105 res_id = reserve_read_pipe(pipe)" << i << R"(, 1);
106 if(is_valid_reserve_id(res_id))
107 {
108 read_pipe(pipe)" << i << R"(, res_id, 0, &dst[gid]);
109 commit_read_pipe(pipe)" << i << R"(, res_id);
110 }
111 })";
112 // clang-format on
113 }
114 stream << "}";
115 }
116
verify_result(void * ptr1,void * ptr2,int n)117 static int verify_result(void *ptr1, void *ptr2, int n)
118 {
119 int i;
120 int sum_input = 0, sum_output = 0;
121 cl_char *inptr = (cl_char *)ptr1;
122 cl_char *outptr = (cl_char *)ptr2;
123
124 for(i = 0; i < n; i++)
125 {
126 sum_input += inptr[i];
127 sum_output += outptr[i];
128 }
129 if(sum_input != sum_output){
130 return -1;
131 }
132 return 0;
133 }
134
verify_result_int(void * ptr1,void * ptr2,int n)135 static int verify_result_int(void *ptr1, void *ptr2, int n)
136 {
137 int i;
138 int sum_input = 0, sum_output = 0;
139 cl_int *inptr = (cl_int *)ptr1;
140 cl_int *outptr = (cl_int *)ptr2;
141
142 for(i = 0; i < n; i++)
143 {
144 sum_input += inptr[i];
145 sum_output += outptr[i];
146 }
147 if(sum_input != sum_output){
148 return -1;
149 }
150 return 0;
151 }
152
test_pipe_max_args(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)153 int test_pipe_max_args(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
154 {
155
156 clMemWrapper pipes[1024];
157 clMemWrapper buffers[2];
158 void *outptr;
159 cl_int *inptr;
160 clProgramWrapper program;
161 clKernelWrapper kernel[2];
162 size_t global_work_size[3];
163 cl_int err;
164 cl_int size;
165 int num_pipe_elements = 1024;
166 int i, j;
167 int max_pipe_args;
168 std::stringstream source;
169 clEventWrapper producer_sync_event = NULL;
170 clEventWrapper consumer_sync_event = NULL;
171 BufferOwningPtr<cl_int> BufferInPtr;
172 BufferOwningPtr<cl_int> BufferOutPtr;
173
174 MTdataHolder d(gRandomSeed);
175 const char *kernelName[] = { "test_multiple_pipe_write",
176 "test_multiple_pipe_read" };
177
178 size_t min_alignment = get_min_alignment(context);
179
180 err = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS,
181 sizeof(max_pipe_args), (void *)&max_pipe_args, NULL);
182 if (err)
183 {
184 print_error(err, " clGetDeviceInfo failed\n");
185 return -1;
186 }
187 if(max_pipe_args < 16){
188 log_error("The device should support minimum 16 pipe objects that could be passed as arguments to the kernel");
189 return -1;
190 }
191
192 global_work_size[0] = (cl_uint)num_pipe_elements * max_pipe_args;
193 size = sizeof(int) * num_pipe_elements * max_pipe_args;
194
195 inptr = (cl_int *)align_malloc(size, min_alignment);
196
197 for(i = 0; i < num_pipe_elements * max_pipe_args; i++){
198 inptr[i] = (int)genrand_int32(d);
199 }
200 BufferInPtr.reset(inptr, nullptr, 0, size, true);
201
202 buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
203 test_error_ret(err, " clCreateBuffer failed", -1);
204
205 outptr = align_malloc(size, min_alignment);
206 BufferOutPtr.reset(outptr, nullptr, 0, size, true);
207 buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err);
208 test_error_ret(err, " clCreateBuffer failed", -1);
209
210 for(i = 0; i < max_pipe_args; i++){
211 pipes[i] = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_pipe_elements, NULL, &err);
212 test_error_ret(err, " clCreatePipe failed", -1);
213 }
214
215 createKernelSourceCode(source, max_pipe_args);
216
217 std::string kernel_source = source.str();
218 const char *sources[] = { kernel_source.c_str() };
219
220 // Create producer kernel
221 err = create_single_kernel_helper(context, &program, &kernel[0], 1, sources,
222 kernelName[0]);
223 test_error_ret(err, " Error creating program", -1);
224
225 //Create consumer kernel
226 kernel[1] = clCreateKernel(program, kernelName[1], &err);
227 test_error_ret(err, " Error creating kernel", -1);
228
229 err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
230 for( i = 0; i < max_pipe_args; i++){
231 err |= clSetKernelArg(kernel[0], i+1, sizeof(cl_mem), (void*)&pipes[i]);
232 }
233 err |= clSetKernelArg(kernel[0], max_pipe_args + 1, sizeof(int), (void*)&max_pipe_args);
234 err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&buffers[1]);
235 for( i = 0; i < max_pipe_args; i++){
236 err |= clSetKernelArg(kernel[1], i+1, sizeof(cl_mem), (void*)&pipes[i]);
237 }
238 err |= clSetKernelArg(kernel[1], max_pipe_args + 1, sizeof(int), (void*)&max_pipe_args);
239 test_error_ret(err, " clSetKernelArg failed", -1);
240
241 // Launch Producer kernel
242 err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
243 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
244
245 // Launch Consumer kernel
246 err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
247 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
248
249 err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL);
250 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
251
252 err = clWaitForEvents(1, &consumer_sync_event);
253 test_error_ret(err, " clWaitForEvents failed", -1);
254
255 if( verify_result( inptr, outptr, num_pipe_elements*sizeof(cl_int))){
256 log_error("test_pipe_max_args failed\n");
257 }
258 else {
259 log_info("test_pipe_max_args passed\n");
260 }
261
262 return 0;
263 }
264
265
test_pipe_max_packet_size(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)266 int test_pipe_max_packet_size(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
267 {
268 clMemWrapper pipe;
269 clMemWrapper buffers[2];
270 void *outptr;
271 cl_char *inptr;
272 clProgramWrapper program;
273 clKernelWrapper kernel[2];
274 size_t global_work_size[3];
275 cl_int err;
276 size_t size;
277 int num_pipe_elements = 1024;
278 int i;
279 cl_uint max_pipe_packet_size;
280 clEventWrapper producer_sync_event = NULL;
281 clEventWrapper consumer_sync_event = NULL;
282 BufferOwningPtr<cl_int> BufferInPtr;
283 BufferOwningPtr<cl_int> BufferOutPtr;
284 MTdataHolder d(gRandomSeed);
285 const char *kernelName[] = { "test_pipe_max_packet_size_write",
286 "test_pipe_max_packet_size_read" };
287
288 size_t min_alignment = get_min_alignment(context);
289
290 global_work_size[0] = (cl_uint)num_pipe_elements;
291
292 std::stringstream source;
293
294 err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_PACKET_SIZE,
295 sizeof(max_pipe_packet_size),
296 (void *)&max_pipe_packet_size, NULL);
297 test_error_ret(err, " clCreatePipe failed", -1);
298
299 if (max_pipe_packet_size < 1024)
300 {
301 log_error(
302 "The device should support minimum packet size of 1024 bytes");
303 return -1;
304 }
305
306 if(max_pipe_packet_size > (32*1024*1024/num_pipe_elements))
307 {
308 max_pipe_packet_size = 32*1024*1024/num_pipe_elements;
309 }
310
311 size = max_pipe_packet_size * num_pipe_elements;
312
313 inptr = (cl_char *)align_malloc(size, min_alignment);
314
315 for(i = 0; i < size; i++){
316 inptr[i] = (char)genrand_int32(d);
317 }
318 BufferInPtr.reset(inptr, nullptr, 0, size, true);
319
320 buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
321 test_error_ret(err, " clCreateBuffer failed", -1);
322
323 outptr = align_malloc(size, min_alignment);
324 BufferOutPtr.reset(outptr, nullptr, 0, size, true);
325
326 buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err);
327 test_error_ret(err, " clCreateBuffer failed", -1);
328
329 pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, max_pipe_packet_size, num_pipe_elements, NULL, &err);
330 test_error_ret(err, " clCreatePipe failed", -1);
331
332 // clang-format off
333 source << R"(
334 typedef struct{
335 char a[)" << max_pipe_packet_size << R"(];
336 }TestStruct;
337
338 __kernel void test_pipe_max_packet_size_write(__global TestStruct *src, __write_only pipe TestStruct out_pipe)
339 {
340 int gid = get_global_id(0);
341 reserve_id_t res_id;
342
343 res_id = reserve_write_pipe(out_pipe, 1);
344 if(is_valid_reserve_id(res_id))
345 {
346 write_pipe(out_pipe, res_id, 0, &src[gid]);
347 commit_write_pipe(out_pipe, res_id);
348 }
349 }
350
351 __kernel void test_pipe_max_packet_size_read(__read_only pipe TestStruct in_pipe, __global TestStruct *dst)
352 {
353 int gid = get_global_id(0);
354 reserve_id_t res_id;
355
356 res_id = reserve_read_pipe(in_pipe, 1);
357 if(is_valid_reserve_id(res_id))
358 {
359 read_pipe(in_pipe, res_id, 0, &dst[gid]);
360 commit_read_pipe(in_pipe, res_id);
361 }
362 }
363 )";
364 // clang-format on
365
366 std::string kernel_source = source.str();
367 const char *sources[] = { kernel_source.c_str() };
368
369 // Create producer kernel
370 err = create_single_kernel_helper(context, &program, &kernel[0], 1, sources,
371 kernelName[0]);
372 test_error_ret(err, " Error creating program", -1);
373
374 //Create consumer kernel
375 kernel[1] = clCreateKernel(program, kernelName[1], &err);
376 test_error_ret(err, " Error creating kernel", -1);
377
378 err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
379 err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
380 err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
381 err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
382 test_error_ret(err, " clSetKernelArg failed", -1);
383
384 // Launch Producer kernel
385 err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
386 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
387
388 // Launch Consumer kernel
389 err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
390 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
391
392 err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &consumer_sync_event, NULL);
393 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
394
395 if( verify_result( inptr, outptr, size)){
396 log_error("test_pipe_max_packet_size failed\n");
397 }
398 else {
399 log_info("test_pipe_max_packet_size passed\n");
400 }
401
402 return 0;
403 }
404
test_pipe_max_active_reservations(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)405 int test_pipe_max_active_reservations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
406 {
407 clMemWrapper pipe;
408 clMemWrapper buffers[2];
409 clMemWrapper buf_reservations;
410 clMemWrapper buf_status;
411 clMemWrapper buf_reserve_id_t_size;
412 clMemWrapper buf_reserve_id_t_size_aligned;
413 cl_int *inptr;
414 void *outptr;
415 int size, i;
416 clProgramWrapper program;
417 clKernelWrapper kernel[3];
418 size_t global_work_size[3];
419 cl_int err;
420 int status = 0;
421 cl_uint max_active_reservations = 0;
422 cl_ulong max_global_size = 0;
423 int reserve_id_t_size;
424 int temp;
425 clEventWrapper sync_event = NULL;
426 clEventWrapper read_event = NULL;
427 BufferOwningPtr<cl_int> BufferInPtr;
428 BufferOwningPtr<cl_int> BufferOutPtr;
429 MTdataHolder d(gRandomSeed);
430 const char *kernelName[3] = { "test_pipe_max_active_reservations_write",
431 "test_pipe_max_active_reservations_read",
432 "pipe_get_reserve_id_t_size" };
433
434 size_t min_alignment = get_min_alignment(context);
435
436 std::stringstream source;
437
438 global_work_size[0] = 1;
439
440 err = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS,
441 sizeof(max_active_reservations),
442 (void *)&max_active_reservations, NULL);
443 test_error_ret(err, " clGetDeviceInfo failed", -1);
444
445 err = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE,
446 sizeof(max_global_size), (void *)&max_global_size,
447 NULL);
448 test_error_ret(err, " clGetDeviceInfo failed", -1);
449
450 max_active_reservations = (max_active_reservations > max_global_size)
451 ? 1 << 16
452 : max_active_reservations;
453
454 if (max_active_reservations < 1)
455 {
456 log_error("The device should support minimum active reservations of 1");
457 return -1;
458 }
459
460 // To get reserve_id_t size
461 buf_reserve_id_t_size = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, sizeof(reserve_id_t_size), NULL, &err);
462 test_error_ret(err, " clCreateBuffer failed", -1);
463
464 // clang-format off
465 source << R"(
466 __kernel void test_pipe_max_active_reservations_write(__global int *src, __write_only pipe int out_pipe, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)
467 {
468 __global reserve_id_t *res_id_ptr;
469 int reserve_idx;
470 int commit_idx;
471
472 for(reserve_idx = 0; reserve_idx < )" << max_active_reservations << R"(; reserve_idx++)
473 {
474 res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);
475 *res_id_ptr = reserve_write_pipe(out_pipe, 1);
476 if(is_valid_reserve_id(res_id_ptr[0]))
477 {
478 write_pipe(out_pipe, res_id_ptr[0], 0, &src[reserve_idx]);
479 }
480 else
481 {
482 *status = -1;
483 return;
484 }
485 }
486
487 for(commit_idx = 0; commit_idx < )" << max_active_reservations << R"(; commit_idx++)
488 {
489 res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);
490 commit_write_pipe(out_pipe, res_id_ptr[0]);
491 }
492 }
493
494 __kernel void test_pipe_max_active_reservations_read(__read_only pipe int in_pipe, __global int *dst, __global char *reserve_id, __global int *reserve_id_t_size_aligned, __global int *status)
495 {
496 __global reserve_id_t *res_id_ptr;
497 int reserve_idx;
498 int commit_idx;
499
500 for(reserve_idx = 0; reserve_idx < )" << max_active_reservations << R"(; reserve_idx++)
501 {
502 res_id_ptr = (__global reserve_id_t*)(reserve_id + reserve_idx*reserve_id_t_size_aligned[0]);
503 *res_id_ptr = reserve_read_pipe(in_pipe, 1);
504
505 if(is_valid_reserve_id(res_id_ptr[0]))
506 {
507 read_pipe(in_pipe, res_id_ptr[0], 0, &dst[reserve_idx]);
508 }
509 else
510 {
511 *status = -1;
512 return;
513 }
514 }
515
516 for(commit_idx = 0; commit_idx < )" << max_active_reservations << R"(; commit_idx++)
517 {
518 res_id_ptr = (__global reserve_id_t*)(reserve_id + commit_idx*reserve_id_t_size_aligned[0]);
519 commit_read_pipe(in_pipe, res_id_ptr[0]);
520 }
521 }
522
523 __kernel void pipe_get_reserve_id_t_size(__global int *reserve_id_t_size)
524 {
525 *reserve_id_t_size = sizeof(reserve_id_t);
526 }
527 )";
528 // clang-format on
529
530 std::string kernel_source = source.str();
531 const char *sources[] = { kernel_source.c_str() };
532
533 // Create producer kernel
534 err = create_single_kernel_helper(context, &program, &kernel[0], 1, sources,
535 kernelName[0]);
536 test_error_ret(err, " Error creating program", -1);
537
538 // Create consumer kernel
539 kernel[1] = clCreateKernel(program, kernelName[1], &err);
540 test_error_ret(err, " Error creating kernel", -1);
541
542 // Create size query kernel for reserve_id_t
543 kernel[2] = clCreateKernel(program, kernelName[2], &err);
544 test_error_ret(err, " Error creating kernel", -1);
545
546 err = clSetKernelArg(kernel[2], 0, sizeof(cl_mem), (void*)&buf_reserve_id_t_size);
547 test_error_ret(err, " clSetKernelArg failed", -1);
548
549 //Launch size query kernel for reserve_id_t
550 err = clEnqueueNDRangeKernel( queue, kernel[2], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event );
551 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
552
553 err = clEnqueueReadBuffer(queue, buf_reserve_id_t_size, true, 0, sizeof(reserve_id_t_size), &reserve_id_t_size, 1, &sync_event, &read_event);
554 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
555
556 err = clWaitForEvents(1, &read_event);
557 test_error_ret(err, " clWaitForEvents failed", -1);
558
559 // Round reserve_id_t_size to the nearest power of 2
560 temp = 1;
561 while(temp < reserve_id_t_size)
562 temp *= 2;
563 reserve_id_t_size = temp;
564
565 size = sizeof(cl_int) * max_active_reservations;
566 inptr = (cl_int *)align_malloc(size, min_alignment);
567
568 for(i = 0; i < max_active_reservations; i++){
569 inptr[i] = (int)genrand_int32(d);
570 }
571 BufferInPtr.reset(inptr, nullptr, 0, size, true);
572
573 buffers[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
574 test_error_ret(err, " clCreateBuffer failed", -1);
575
576 outptr = align_malloc(size, min_alignment);
577 BufferOutPtr.reset(outptr, nullptr, 0, size, true);
578
579 buffers[1] = clCreateBuffer(context, CL_MEM_HOST_READ_ONLY, size, NULL, &err);
580 test_error_ret(err, " clCreateBuffer failed", -1);
581
582 buf_reserve_id_t_size_aligned = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(reserve_id_t_size), &reserve_id_t_size, &err);
583 test_error_ret(err, " clCreateBuffer failed", -1);
584
585 //For error status
586 buf_status = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(int), &status, &err);
587 test_error_ret(err, " clCreateBuffer failed", -1);
588
589 pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), max_active_reservations, NULL, &err);
590 test_error_ret(err, " clCreatePipe failed", -1);
591
592 // Global buffer to hold all active reservation ids
593 buf_reservations = clCreateBuffer(context, CL_MEM_HOST_NO_ACCESS, reserve_id_t_size*max_active_reservations, NULL, &err);
594 test_error_ret(err, " clCreateBuffer failed", -1);
595
596 err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
597 err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
598 err |= clSetKernelArg(kernel[0], 2, sizeof(cl_mem), (void*)&buf_reservations);
599 err |= clSetKernelArg(kernel[0], 3, sizeof(cl_mem), (void*)&buf_reserve_id_t_size_aligned);
600 err |= clSetKernelArg(kernel[0], 4, sizeof(cl_mem), (void*)&buf_status);
601 test_error_ret(err, " clSetKernelArg failed", -1);
602
603 err = clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
604 err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
605 err |= clSetKernelArg(kernel[1], 2, sizeof(cl_mem), (void*)&buf_reservations);
606 err |= clSetKernelArg(kernel[1], 3, sizeof(cl_mem), (void*)&buf_reserve_id_t_size_aligned);
607 err |= clSetKernelArg(kernel[1], 4, sizeof(cl_mem), (void*)&buf_status);
608 test_error_ret(err, " clSetKernelArg failed", -1);
609
610 clReleaseEvent(sync_event);
611
612 // Launch Producer kernel
613 err = clEnqueueNDRangeKernel(queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event);
614 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
615
616 err = clEnqueueReadBuffer(queue, buf_status, true, 0, sizeof(int), &status, 1, &sync_event, NULL);
617 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
618
619 if(status != 0)
620 {
621 log_error("test_pipe_max_active_reservations failed\n");
622 return -1;
623 }
624
625 clReleaseEvent(sync_event);
626 // Launch Consumer kernel
627 err = clEnqueueNDRangeKernel(queue, kernel[1], 1, NULL, global_work_size, NULL, 0, NULL, &sync_event);
628 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
629
630 err = clEnqueueReadBuffer(queue, buf_status, true, 0, sizeof(int), &status, 1, &sync_event, NULL);
631 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
632
633 if(status != 0)
634 {
635 log_error("test_pipe_max_active_reservations failed\n");
636 return -1;
637 }
638
639 err = clEnqueueReadBuffer(queue, buffers[1], true, 0, size, outptr, 1, &sync_event, NULL);
640 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
641
642 if( verify_result_int( inptr, outptr, max_active_reservations)){
643 log_error("test_pipe_max_active_reservations failed\n");
644 return -1;
645 }
646 else {
647 log_info("test_pipe_max_active_reservations passed\n");
648 }
649
650 return 0;
651 }