- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
Background
I am trying to implement a DSP chain with Variable Size FFT Intel FPGA IP. This IP takes the size of the current FFT as a log2 value, supplied as the 4 most significant bits of the AVST data bus. So for an FFT operating on complex floats, the AVST input/output interface has 68 bits on the data bus.
For each set of data to be processed by the DSP chain, the FFT size is known ahead of time. The DSP chain is configured using a Nios, which writes to CSRs of the IP cores to configure their parameters, start the cores, and monitor them. To provide the FFT size to the VFFT IP through its AVST data channels, I am writing the size to a 16 bit register in the SYCL IP block preceding the VFFT. This IP block should concatenate this value to the MSBs of its output pipe along with its intended computations.
This seemed to work with the following rough outline of the SCYL kernel, showing just the logic, no boilerplate:
```
// Defining the properties of the output pipe of SYCL kernel
using VarFFTPipeProps = decltype(sycl::ext::oneapi::experimental::properties(
sycl::ext::intel::experimental::first_symbol_in_high_order_bits<true>,
sycl::ext::intel::experimental::protocol_avalon_streaming_uses_ready,
sycl::ext::intel::experimental::bits_per_symbol<68>));
...
/* The pipe is defined with a 68 bit int. We tried to define a struct with ac_complex<float> and ac_int<4, false> members, but we couldn't get the compiler to see that struct as 68 bits.
sycl::ext::intel::experimental::pipe<KernelOutput, ac_int<68, false>, 0, VarFFTPipeProps>;*/
...
/* The input pipe is of data type ac_complex<float>. The logic is to maintain the bit structure of the complex and create an int array to bit_fill the output value. Assume a uint16_t parameter "log2_fft_size" is passed into the kernel via CSR*/
ac_complex<float> data_in;
data_in = PipeIn::read();
...
// other kernel logic, assume data still resides in data_in and is not changed
...
int real, imag;
/* not sure about the correct way to maintain the bit structure, whether I could use reinterpret_cast or if the API has some way to reinterpret a set of bits as another type.
Compiler seems to do okay with memcpy here despite the target hardware being FPGA RTL */
memcpy(&real, &data_in.real(), sizeof(data_in.real()));
memcpy(&imag, &data_in.imag(), sizeof(data_in.imag()));
int out_bit_vec[] = {log2_fft_size, imag, real};
ac_int<68, false> out_data;
out_data.bit_fill(out_bit_vec);
PipeOut::write(out_data);
```
This has appeared to work for us. However, this VFFT is followed by another SYCL kernel, followed by a VFFT block for the inverse. So this middle core between the VFFTs needs basically to pass through the four most significant bits representing the log2 fft size.
This middle kernel has input and output pipes of the same type as the output of the previous kernel, with an `ac_int<68, false>` data type. In an attempt to extract the uppermost bits from the input pipe in this kernel:
```
ac_int<68, false> data_in;
data_in = PipeIn::read();
// extract upper bits
ac_int<4,false> log2_fft_size;
log2_fft_size = data_in.slc<4>(64);
```
The rest of the kernel code does some other DSP logic, and then performs the same bit_fill with int array process to form the output value. This results in an optimizer failure. I have uploaded the output of `make report` as a text file. But if I just directly initialize the 4 bit integer:
```
ac_int<4,false> log2_fft_size = 9;
```
the kernel compiles fine. It seems that this operation should be supported. I have tried bit shifting the input data, I have tried using a temporary 72 bit integer which is initialized with the input data hoping that slicing a whole byte would alleviate, but to no avail. I cannot figure out how to extract those four upper bits from the input data pipe.
What is the solution to this problem? This project is targeting an Agilex device, using Quartus/OneAPI version 24.1. Please let me know if there is anything I can try!
Link Copied

- Subscribe to RSS Feed
- Mark Topic as New
- Mark Topic as Read
- Float this Topic for Current User
- Bookmark
- Subscribe
- Printer Friendly Page