Forum Discussion

Altera_Forum's avatar
Altera_Forum
Icon for Honored Contributor rankHonored Contributor
10 years ago

Channel has no point Connection

I just got a weird error when trying to build my OpenCL code with aocl:


Error Channel has no point connection: ID= in_channel_a
  avm_channel_id_in_channel_a_write connection misin, or optimized away

I simplified the code down to the following, which still produces the error:

# pragma OPENCL EXTENSION cl_altera_channels : enable 
channel uint2 in_channel_a __attribute__((depth(1)));
channel uint2 in_channel_b __attribute__((depth(1)));
channel uint2 out_channel __attribute__((depth(1)));
__kernel void in_streamer_a(__global const uint2* in, uint n) {
    for(uint i = 0; i < n; ++i) {
        write_channel_altera(in_channel_a, in);
    }
}
__kernel void in_streamer_b(__global const uint2* in, uint n) {
    for(uint i = 0; i < n; ++i) {
        write_channel_altera(in_channel_b, in);
    }
}
__kernel void out_streamer_a(__global uint2* restrict out, uint n) {
    for(uint i = 0; i < n; ++i) {
        out = read_channel_altera(out_channel);
    }
}
__kernel void simplified_worker_kernel(uint n_steps) {
    for(int i = -1; i < n_steps; ++i) {
        uint2 write_data;
        uint2 current_a;
        uint2 current_b;
        current_a = read_channel_altera(in_channel_a);
        current_b = read_channel_altera(in_channel_b);
        write_data = current_a + current_b;
        write_channel_altera(out_channel, write_data);
    }
}

As far as I can tell, each channel is read at exactly one place and written to at exactly one place. The connections should be obvious. What am I missing?

3 Replies

  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor

    I seem to have found the issue. When I change the loop from

    for(int i = -1; i < n_steps; ++i)

    to

    for(int i = -1; i < (int)n_steps; ++i)

    The error disappears.

    This means the compiler evaluates "(int)-1 < (uint)0" as false. As a result the entire Loop is optimized away, so in_channel_a for example is never read. I still think this is a compiler error though, does anyone here know what the OpenCL spec says about signed-unsigned comparison?
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor

    what is your n_steps declared as? let say if you place i = 0 , it this work?

    for(int i = 0; i < n_steps; ++i)
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor

    Yes, that works. Casting to int works too. For my personal needs the problem is solved, but I'm still wondering why the error occurred in the first place. I'm sure I won't be the last to encounter it.