How to use Pipe in OpenCL2.0

[lucas_w_2]

Hi guys!
 i am tring to use Pipe in OpenCL2.0. But I am confused about the built-in Pipe read and write functions,such as write_pipe, reserver_read_pipe.

Code :

reserve_id_t reserve_write_pipe (
                                        pipe gentype p,
                                       uint num_packets)

Code :

int write_pipe (pipe gentype p,
                                reserve_id_t reserve_id,
                                uint index,
                               const gentype *p

1. For the reserve_write_pipe() function, Description is that 

Code :

Reserve num_packets entries for writing to pipe p. Returns a valid reservation ID if the reservation is successful.

I don't know what's the meaning of the num_packets argument.
2. For write_pipe() function,Description is that 

Code :

Write packet specified by ptr to the reserved area of the pipe referred to by reserve_id and index.

What's the meaning of "reserved area of the pipe"? and What's the meaning of "reserve_id " and "index"?
3. If i want to write to Pipe by order(such as work-item 0 writes 0 to the pipe,work-items 1 write 1 to the pipe,.. work-item n writes n to the pipe,so the data in the pipe is "012345……n" ,and 0 is the first one in the pipe),what should I do ?
Can someone help me please? Thanks.
hi_buddy

[Robert_I_Intel]

I would highly discourage you to use pipes, but if you insist:

1. Let's say you have an int pipe and num_packets is 5: you reserve 5 int slots in the reservation id that is returned (let's say 42), so then you can write to these slots by providing an index (0 thru 4) and your reservation id (42).

2. What you can do is the first work item in a work group makes a reservation the size of the work group. Then every work item writes to the slot by providing its local id as an index.

Just curious: why do you want to use pipes? Do you actually have an algorithm that you think will map to pipes or just want to press all the buttons in OpenCL 2.0 spec?

[lucas_w_2]

Hi Robert. I just want to have a try , because Pipe is the new feature in OpenCL2.0

2. What you can do is the first work item in a work group makes a reservation the size of the work group. Then every work item writes to the slot by providing its local id as an index.

^{I don't understand what have you said above. Can you explain that ? thanks a lot}

[Robert_I_Intel]

local reserve_id_t rid;
if (get_local_id(0) == 0)
    rid = reserve_write_pipe (p, get_local_size(0));

barrier(CLK_LOCAL_MEM_FENCE);
int foo = get_local_id(0);
write_pipe (p, rid, get_local_id(0), &foo);

work_group_commit_write_pipe(p, rid);

[Robert_I_Intel]

You can even make it simpler:

reserve_id_t rid = work_group_reserve_write_pipe(p, get_local_size(0));
int foo = get_local_id(0); 
write_pipe (p, rid, get_local_id(0), &foo); 
work_group_commit_write_pipe(p, rid); 

[lucas_w_2]

Hi Robert,

1. I have modified the code under your guidance, now  I can get the sequential data across a workgroup,not across the NDRange. it seems that there is no way to control the order across the NDRange.

2.I'm curious that why you highly discourage me to use pipes.

Thanks.

[Robert_I_Intel]

Hi Lucas,

1. Correct: work groups execute independently, so you won't be able to control the order across the NDRange.

2. Pipes are a weird beast in OpenCL: usually, when you talk pipes, concurrently executing processes come to mind. Typically, people expect kernels to run concurrently and communicate via pipes: that's not what is going on here. Pipes are just glorified buffers - there is no special magic and anything you can accomplish with pipes you can do with buffers. There is no concurrency either: one kernel completes execution and then the other starts, so you need to size the pipe so it could fit ALL the output of the kernel - it is just a memory buffer :(