Channel ID is ignored for io_pipes

[typohnebild]

Describe the bug

For ext::intel::kernel_readable_io_pipe/ext::intel::kernel_writeable_io_pipe, the ID in the channel_id struct is ignored. If only one read and one write channel are used, the pipes always get mapped to the 0-indexed read or write channel.

The minimal example below should write data to the second QSFP Port and should read these values from the first QSFP Port. With the help of our infrastructure, these two ports are linked to each other. So it is expected that read_chanel_id is mapped to the chan_id "kernel_input_ch1" and write_channel_id is mapped to the chan_id "kernel_output_ch0"

It worked for oneAPI version until 2023.1, but this does not work for newer ones anymore.

Since version 2023.2, it does not matter which ID is declared in the struct, but the read_chanel_id is mapped to kernel_input_ch0, and write_channel_id is mapped to kernel_output_ch0. In this case, we always get a loop back on the first port.

If there is more than one read / one write declared in the code, the mapping gets somehow random. 

To reproduce

1. Minimal Example:

#include <sycl/sycl.hpp>
#include <sycl/ext/intel/fpga_extensions.hpp>
#include <iostream>
#include <vector>

using namespace sycl;

// Vector size for this example
constexpr size_t kSize = 1024;

class VectorAdd;

struct read_channel_id {
  static constexpr unsigned id = 2;
};

struct write_channel_id {
  static constexpr unsigned id = 1;
};

struct channel_data {
  int data[8];
};

using read_channel = ext::intel::kernel_readable_io_pipe<read_channel_id, channel_data, 8>;
using write_channel = ext::intel::kernel_writeable_io_pipe<write_channel_id, channel_data, 8>;


int main() {

  std::vector<int> vec_a(kSize), vec_b(kSize), vec_r(kSize);
  for (int i = 0; i < kSize; i++) {
    vec_a[i] = rand();
    vec_b[i] = rand();
  }

#if defined(FPGA_EMULATOR)
  auto device_selector = ext::intel::fpga_emulator_selector_v;
#else
  auto device_selector = ext::intel::fpga_selector_v;
#endif

  try {
    queue q(device_selector);

    std::cout << "Running on device: "
              << q.get_device().get_info<info::device::name>() << "\n";

    {
      buffer buf_a(vec_a);
      buffer buf_b(vec_b);
      buffer buf_r(vec_r);

      q.submit([&](handler& h) {

        accessor a(buf_a, h, read_only);
        accessor b(buf_b, h, read_only);
        accessor r(buf_r, h, write_only, no_init);

        h.single_task<VectorAdd>([=]() [[intel::kernel_args_restrict]] {
          for (int i = 0; i < kSize/8; ++i) {
            channel_data write_data;
            for (int j=0; j < 8; j++) {
              write_data.data[j] = a[i*8 + j] + b[i*8 + j];
            }
            write_channel::write(write_data);
            channel_data read_data = read_channel::read();
            for (int j=0; j < 8; j++) {
              r[i*8 + j] = read_data.data[j];
            }
          }
        });
      });
    }
  }
  catch (sycl::exception const& e) {
    // Catches exceptions in the host code
    std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n";

    // Most likely the runtime couldn't find FPGA hardware!
    if (e.code().value() == CL_DEVICE_NOT_FOUND) {
      std::cerr << "If you are targeting an FPGA, please ensure that your "
                   "system has a correctly configured FPGA board.\n";
      std::cerr << "Run sys_check in the oneAPI root directory to verify.\n";
      std::cerr << "If you are targeting the FPGA emulator, compile with "
                   "-DFPGA_EMULATOR.\n";
    }
    std::terminate();
  }

  // Check the results.
  int correct = 0;
  for (int i = 0; i < kSize; i++) {
    if ( vec_r[i] == vec_a[i] + vec_b[i] ) {
      correct++;
    }
  }

  // Summarize and return.
  if (correct == kSize) {
    std::cout << "PASSED: results are correct\n";
  } else {
    std::cout << "FAILED: results are incorrect\n";
  }

  return !(correct == kSize);
}

2. Specify the command which should be used to compile the program

$ icpx  -O3 -fsycl -fintelfpga -Wall -Xsv -qactypes -DFPGA_HARDWARE -c minimal_example.cpp -o minimal_example.cpp.o
$ icpx -fsycl -fintelfpga -qactypes -Xshardware -Xstarget=p520_max_sg280l -o minimal_example.fpga minimal_example.cpp.o

4. Indicate what is wrong and what was expected
   For oneAPI 2023.1 the read_channel_id gets mapped to chan_id="kernel_input_1" according to the minimal_example.report.bc.xml

<INTERFACE port="iord_bl_io_acl_c_ZN4sycl3_V13ext5intel23kernel_readable_io_pipeI15read_channel_id12channel_dataLm8EE9m_StorageE_pipe_channel" cosim_name="" type="host" width="256" true_lsu_width="4294967295" burstwidth="1" pipelined="true" addr_space="__channel" aspace_id="32" optype="read" criticality="1024" buffer_location="" group_id="0" connected_to_agent="0" depth="8" user_depth="8" ready_latency="0" chan_id="kernel_input_ch1" is_fifo="1" init="" almost_full_value="8" bypass_channel="0" inter_kernel_pipelining="1" write_success_ignored="0" sfks_id="0" is_sfks_entry="0" is_sfks_token="0" nb_sfks_entry_write="0" read_ready_ignored="0" bits_per_symbol ="0" uses_packets ="0" stall_free ="0" uses_empty ="0" first_symbol_in_high_order_bits ="0" uses_stall_latency ="0" channel_type ="host_pipe" pipe_mapping_id="0" impl_in_csr ="0" csr_address ="0" func_wrapper_interface_name ="" cra_interface_name ="" />
<INTERFACE port="iowr_bl_io_acl_c_ZN4sycl3_V13ext5intel24kernel_writeable_io_pipeI16write_channel_id12channel_dataLm8EE9m_StorageE_pipe_channel" cosim_name="" type="host" width="256" true_lsu_width="4294967295" burstwidth="1" pipelined="true" addr_space="__channel" aspace_id="32" optype="write" criticality="1024" buffer_location="" group_id="0" connected_to_agent="0" depth="8" user_depth="8" ready_latency="0" chan_id="kernel_output_ch0" is_fifo="1" init="" almost_full_value="-1" bypass_channel="0" inter_kernel_pipelining="1" write_success_ignored="0" sfks_id="0" is_sfks_entry="0" is_sfks_token="0" nb_sfks_entry_write="0" read_ready_ignored="0" bits_per_symbol ="0" uses_packets ="0" stall_free ="0" uses_empty ="0" first_symbol_in_high_order_bits ="0" uses_stall_latency ="0" channel_type ="host_pipe" pipe_mapping_id="0" impl_in_csr ="0" csr_address ="0" func_wrapper_interface_name ="" cra_interface_name ="" />

While for oneAPI 2024.0 (and all later versions) it gets mapped to chan_id="kernel_input_0"

<INTERFACE port="iord_bl_io_acl_c_read_channel_id_pipe_channel" cosim_name="" type="host" width="256" true_lsu_width="4294967295" burstwidth="1" pipelined="true" addr_space="__channel" aspace_id="32" optype="read" criticality="1024" buffer_location="" group_id="0" connected_to_agent="0" depth="8" user_depth="8" ready_latency="0" chan_id="kernel_input_ch0" is_fifo="1" init="" almost_full_value="8" bypass_channel="0" inter_kernel_pipelining="1" write_success_ignored="0" sfks_id="0" is_sfks_entry="0" is_sfks_token="0" nb_sfks_entry_write="0" read_ready_ignored="0" bits_per_symbol ="0" uses_packets ="0" stall_free ="0" uses_empty ="0" first_symbol_in_high_order_bits ="0" uses_stall_latency ="0" channel_type ="host_pipe" pipe_mapping_id="0" hostpipe_name ="" impl_in_csr ="0" />
<INTERFACE port="iowr_bl_io_acl_c_write_channel_id_pipe_channel" cosim_name="" type="host" width="256" true_lsu_width="4294967295" burstwidth="1" pipelined="true" addr_space="__channel" aspace_id="32" optype="write" criticality="1024" buffer_location="" group_id="0" connected_to_agent="0" depth="8" user_depth="8" ready_latency="0" chan_id="kernel_output_ch0" is_fifo="1" init="" almost_full_value="-1" bypass_channel="0" inter_kernel_pipelining="1" write_success_ignored="0" sfks_id="0" is_sfks_entry="0" is_sfks_token="0" nb_sfks_entry_write="0" read_ready_ignored="0" bits_per_symbol ="0" uses_packets ="0" stall_free ="0" uses_empty ="0" first_symbol_in_high_order_bits ="0" uses_stall_latency ="0" channel_type ="host_pipe" pipe_mapping_id="0" hostpipe_name ="" impl_in_csr ="0" />

Environment

OS: Red Hat Enterprise Linux 8.8 / Ubuntu
Target device and vendor: Intel Stratix FPGA 2800 GX on Bittware 520n (1SG280LU3F50E1VGS1)
Quartus Version: 20.04
BSP: p520_max_sg280l

The channel section of the board_spec.xml looks like this

  <channels>
    <interface name="board" port="io_to_dev_ch0" type="streamsource" width="256" chan_id="kernel_input_ch0"/>
    <interface name="board" port="dev_to_io_ch0" type="streamsink" width="256" chan_id="kernel_output_ch0"/>
    <interface name="board" port="io_to_dev_ch1" type="streamsource" width="256" chan_id="kernel_input_ch1"/>
    <interface name="board" port="dev_to_io_ch1" type="streamsink" width="256" chan_id="kernel_output_ch1"/>
    <interface name="board" port="io_to_dev_ch2" type="streamsource" width="256" chan_id="kernel_input_ch2"/>
    <interface name="board" port="dev_to_io_ch2" type="streamsink" width="256" chan_id="kernel_output_ch2"/>
    <interface name="board" port="io_to_dev_ch3" type="streamsource" width="256" chan_id="kernel_input_ch3"/>
    <interface name="board" port="dev_to_io_ch3" type="streamsink" width="256" chan_id="kernel_output_ch3"/>
  </channels>

Additional context

No response