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

Aoc compilation error: Error: Channel has no point connection, connection missing,..

Altera_Forum
Honored Contributor II
1,015 Views

In my OpenCL design, I need to split a stream into 4 streams. 

Then each stream gets processed independently. 

My design works as expected during simulation, but fails when compiling it to hardware. 

Here is the error message. It seems to be related with the split channel. 

I checked the log file, but it is not quite relevant to the error and I don't have a clue yet. 

 

The kernel code is also attached in the end.  

 

2 warnings generated. 

error: channel has no point connection: id= ch1 

avm_channel_id_ch1_read connection missing, or optimized away 

error: could not write system script to file. system generation aborted! 

error: system integrator failed. 

refer to graph_fpga/graph_fpga.log for details. 

 

 

By the way, I may working on the Harp system and the compiler's version is  

altera sdk for opencl, 64-bit offline compiler 

version 16.0.2 build 222 

copyright (c) 2016 altera corporation 

 

Any suggestions will be appreciated.channel int inCh __attribute__((depth(64))); channel int Ch0 __attribute__ ((depth(64))); channel int Ch1 __attribute__ ((depth(64))); channel int Ch2 __attribute__ ((depth(64))); channel int Ch3 __attribute__ ((depth(64))); channel int eofCh0 __attribute__ ((depth(2))); channel int eofCh1 __attribute__ ((depth(2))); channel int eofCh2 __attribute__ ((depth(2))); channel int eofCh3 __attribute__ ((depth(2))); __kernel void __attribute__((task)) readInStream( __global const int* restrict src, const int num ) { for(int i = 0; i < num; i++){ int data = src; write_channel_altera(inCh, data); } } // Split the input into multiple streams __kernel void __attribute__((task)) split( const int num ) { for(int i = 0; i < num; i++){ int idx = read_channel_altera(inCh); int chIdx = idx & 0xC0; bool success = false; do{ if(chIdx == 0){ success = write_channel_nb_altera(Ch0, idx); } else if(chIdx == 1){ success = write_channel_nb_altera(Ch1, idx); } else if(chIdx == 2){ success = write_channel_nb_altera(Ch2, idx); } else{ success = write_channel_nb_altera(Ch3, idx); } } while(!success); } write_channel_altera(eofCh0, 1); write_channel_altera(eofCh1, 1); write_channel_altera(eofCh2, 1); write_channel_altera(eofCh3, 1); } __kernel void __attribute__((task)) write0( __global int* restrict dst ) { int flag = 0; bool flagValid = false; bool dataValid = false; while(true){ int idx = read_channel_nb_altera(Ch0, &dataValid); if(dataValid){ dst = 1; } int tmp = read_channel_nb_altera(eofCh0, &flagValid); if(flagValid) flag = tmp; if(flag == 1 && dataValid == false) break; } } __kernel void __attribute__((task)) write1( __global int* restrict dst ) { int flag = 0; bool flagValid = false; bool dataValid = false; while(true){ int idx = read_channel_nb_altera(Ch1, &dataValid); if(dataValid){ dst = 1; } int tmp = read_channel_nb_altera(eofCh1, &flagValid); if(flagValid) flag = tmp; if(flag == 1 && dataValid == false) break; } } __kernel void __attribute__((task)) write2( __global int* restrict dst ) { int flag = 0; bool flagValid = false; bool dataValid = false; while(true){ int idx = read_channel_nb_altera(Ch2, &dataValid); if(dataValid){ dst = 1; } int tmp = read_channel_nb_altera(eofCh2, &flagValid); if(flagValid) flag = tmp; if(flag == 1 && dataValid == false) break; } } __kernel void __attribute__((task)) write3( __global int* restrict dst ) { int flag = 0; bool flagValid = false; bool dataValid = false; while(true){ int idx = read_channel_nb_altera(Ch3, &dataValid); if(dataValid){ dst = 1; } int tmp = read_channel_nb_altera(eofCh3, &flagValid); if(flagValid) flag = tmp; if(flag == 1 && dataValid == false) break; } }
0 Kudos
2 Replies
Altera_Forum
Honored Contributor II
246 Views

I am really surprised your code works correctly in the emulator since you have a clear mistake in your code: 

 

int chIdx = idx & 0xC0; 

 

chIdx will never be equal to numbers 1, 2 and 3 since you are zeroing out the first 6 (binary) digits of idx. Because of this, the write to channels Ch1 to Ch3 will be optimized out by the compiler since they will never be triggered. Needless to say, since each channel needs to have at least one read and one write point, compilation fails. If you replace 0xC0 with 0x03, which I think is the correct choice for your code, your kernel will compile correctly.
0 Kudos
Altera_Forum
Honored Contributor II
246 Views

Thank you very much. I see the problem. So two data paths are removed due to the coding mistakes 

and there are no reading for the two corresponding channels. 

 

The code is simplified based on my project. The original design split the data path into 4 paths and  

gets processed independently. As long as all the input data are processed, the final result is correct  

during the simulation. I didn't check the utilization of the data paths and missed the bug. 

 

Again thank you.
0 Kudos
Reply