Altera_Forum
Honored Contributor
7 years agoAoc compilation error: Error: Channel has no point connection, connection missing,..
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 corporationAny 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;
}
}