Kernel hangs when using OpenCL pipes

Question

I am trying to write an OpenCL kernel that uses OpenCL pipes. The kernel code is given below.

    uint tid = get_global_id(0);
    uint numWorkItems = get_global_size(0);
    int i;
    int rid;
    int temp = 0, temp1 = 0;
    int val;

    int    szgr = get_local_size(0);
    int    lid  = get_local_id(0);

    for(i = tid + start_index; i < rLen; i = i + numWorkItems){
            temp = 0;
            val = input[i];
            temp = hashTable[val - 1];
            if(temp){
                    temp1 = projection[val - 1];
            }

            reserve_id_t rid1 = work_group_reserve_write_pipe(c0, szgr);

            while(is_valid_reserve_id(rid1) == false){
            rid1 = work_group_reserve_write_pipe(c0, szgr);
            }

            if(is_valid_reserve_id(rid1))
            {

                    write_pipe(c0,rid1,lid, &temp);
                    work_group_commit_write_pipe(c0, rid1);
            }

            reserve_id_t rid2 = work_group_reserve_write_pipe(c1, szgr);

            while(is_valid_reserve_id(rid2) == false){
            rid2 = work_group_reserve_write_pipe(c1, szgr);
            }

            if(is_valid_reserve_id(rid2))
            {

                    write_pipe(c1,rid2,lid, &temp1);
                    work_group_commit_write_pipe(c1, rid2);
            }
    }

But the work_group_reserve_write_pipe function always fails and because of this the kernels hangs at the while loop. If I remove this while loop then the code doesnt hang but writing to the pipe doesnt happen. Can someone tell me why this is happening?

The pipe is declared as a _write_only pipe.

score 1 · Accepted Answer · answered Feb 21 '17 at 08:58

About work_group_reserve_write_pipe:

This built-in function must be encountered by all work-items in a work-group executing the kernel with the same argument values; otherwise the behavior is undefined.

the loop starts from tid + start_index so after some loop iterations, some work items doesn't hit this instruction. Also a while loop is doing same undefined behaviour.

Kernel hangs when using OpenCL pipes

1 Answers1