Intel® Quartus® Prime Software
Intel® Quartus® Prime Design Software, Design Entry, Synthesis, Simulation, Verification, Timing Analysis, System Design (Platform Designer, formerly Qsys)
16556 Discussions

how to use block channel and how to set channel depth

Altera_Forum
Honored Contributor II
1,351 Views

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
Altera_Forum
Honored Contributor II
376 Views

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
Altera_Forum
Honored Contributor II
376 Views

 

--- 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
Altera_Forum
Honored Contributor II
376 Views

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
Reply