Community
cancel
Showing results for 
Search instead for 
Did you mean: 
Highlighted
Valued Contributor III
983 Views

how to use block channel and how to set channel depth

Hi!  

When I use 'block channel' to pass data from one kernel to another, my program got stuck, the code post as follow: 

//---------------------------------------------------------------------------------------------------------------# define FLIT_MAX_PAYLOAD_LEN 61 

 

 

typedef unsigned char uint8_t, u8; 

 

 

typedef struct { 

int cnt; 

uint8_t payload[FLIT_MAX_PAYLOAD_LEN]; 

} msg_flit_t;//msg_channel 

 

# pragma OPENCL EXTENSION cl_altera_channels : enable 

channel msg_flit_t cnt_out_0 __attribute__((depth(2))); 

global int count; 

 

 

__kernel void producer(){ 

int i; 

int j; 

msg_flit_t x; 

 

 

for(i = 0; i < FLIT_MAX_PAYLOAD_LEN; i++){ 

x.payload[i] = i; 

count = 1; 

for(j = 0; j < 50; j++){ 

x.payload[0] += 1; 

write_channel_altera(cnt_out_0, x); 

 

 

__kernel void cosumer(){ 

msg_flit_t x; 

 

 

for(int i = 0; i < 50; i++){ 

x = read_channel_altera(cnt_out_0); 

printf("%d\n", x.payload[0]); 

//--------------------------------------------------------------------------------------------------------------- 

 

1> The loop iteration is 50-times, which is > 2(depth of channel), the program got stuck and never come back again, Why?  

2> How to set the depth of channel if the producer and cosumer kernel is as follow: 

 

__kernel void producer() 

....... 

while(1) { 

write_altera_channel(); 

....... 

 

__kernel void consumer() 

....... 

while(1) { 

read_altera_channel(); 

....... 

}
0 Kudos
3 Replies
Highlighted
Valued Contributor III
8 Views

Re: how to use block channel and how to set channel depth

Hi hackeren, 

 

I have seen this issue in my own code but of course I can't remember exactly how I fixed it 🙂 It was either 

 

1. My host code was incorrect and the second kernel wasn't correctly enqueued or 

2. The emulation environment doesn't entirely correctly emulate the hardware and the consumer kernel was not launched before the producer kernel had filled up the channel, increasing the channel depth for emulation only may fix it. 

 

Wish I had noted what I did a bit better.
0 Kudos
Valued Contributor III
8 Views

Re: how to use block channel and how to set channel depth

 

--- Quote Start ---  

Hi hackeren, 

 

I have seen this issue in my own code but of course I can't remember exactly how I fixed it 🙂 It was either 

 

1. My host code was incorrect and the second kernel wasn't correctly enqueued or 

2. The emulation environment doesn't entirely correctly emulate the hardware and the consumer kernel was not launched before the producer kernel had filled up the channel, increasing the channel depth for emulation only may fix it. 

 

Wish I had noted what I did a bit better. 

--- Quote End ---  

 

 

TKS for your advice ! I create two queues, one for producer kernel and another for consumer kernel, instead of only one command queue for all the kernels.
0 Kudos
Highlighted
Valued Contributor III
8 Views

Re: how to use block channel and how to set channel depth

Yes, that is correct. I have the same in my implementation. One error I had at one point was using different contexts for each kernel, as I recall this was an issue. Did you try increasing the channel depth?

0 Kudos