How does the producer, consumer, and manager strategy in the programming guide work?

Thanks for the suggestion.

I implemented a design filtering a stream of data as you suggested using non-blocking channel.

(Basically the design scans the input data and selects data that is larger than 10,

then the selected data is multiplied by a constant. Finally, the result is written back to memory. )

It works as expected. I have the kernel code attached in case people in the forum may be interested.

Also the code seems to be a bit complex, so it will be appreciated if you have more efficient implementation.

Some minor questions:

When I use the flag_valid as the end of the condition in k1 and k2,

Some of the data in the tail is not handled. Basically it indicates that

when the flag_valid is captured in k1, the data in ch01 has not been processed.

Even though I tried to add memory fence in the code, the result remains

wrong. I don't actually understand this part.

channel int ch01 __attribute__ ((depth(16)));
channel int ch12 __attribute__ ((depth(16)));
channel int k0_end __attribute__ ((depth(4)));
channel int k1_end __attribute__ ((depth(4)));
// Read data from DDR 
__kernel void __attribute__((task)) k0(
        __global const int* restrict  src,
        const int num)
{
    int i;
    for(i = 0; i < num; i++){
        bool success = false;
        do{
            success = write_channel_nb_altera(ch01, src);
        }
        while(!success);
    }
    //printf("%d data is sent to ch01\n", i);
    //mem_fence(CLK_GLOBAL_MEM_FENCE | CLK_CHANNEL_MEM_FENCE);
    write_channel_altera(k0_end, 1);
}
// Filtering
__kernel void k1(
        const int threshold,
        const int cons)
{
    bool flag_valid = false;
    bool data_valid = false;
    int flag = 0;
    int data;
    int count = 0;
    int total = 0;
    while(true){
        data = read_channel_nb_altera(ch01, &data_valid);
        if(data_valid){
            total++;
            if(data > threshold){
                data = data * cons;
                count++;
                // write to channel12
                bool success = false;
                do{
                    success = write_channel_nb_altera(ch12, data);
                }
                while(!success);
            }
        }
        int tmp = read_channel_nb_altera(k0_end, &flag_valid);
        if(flag_valid) flag = tmp;
        if(flag == 1 && data_valid == false){
            //mem_fence(CLK_CHANNEL_MEM_FENCE);
            write_channel_altera(k1_end, 1);
            break;
        }
    }
    //printf("%d data go through the filter\n", total);
    //printf("%d data is changed.\n", count);
}
// Write data to DDR
__kernel void __attribute__((task)) k2(
        __global int* restrict dst
        )
{
    bool flag_valid = false;
    bool data_valid = false;
    int flag = 0;
    int data;
    int idx = 0;
    while(true){
        data = read_channel_nb_altera(ch12, &data_valid);
        // write data to DDR
        if(data_valid){
            dst = data;
            idx++;
        }
        int tmp = read_channel_nb_altera(k1_end, &flag_valid);
        if(flag_valid) flag = tmp;
        if(flag == 1 && data_valid == false){
            //mem_fence(CLK_GLOBAL_MEM_FENCE | CLK_CHANNEL_MEM_FENCE);
            break;
        }
    }
    //printf("%d data is written to memory.\n", idx);
}

Cheng Liu