[bug] FPGA kernels generate an extra bracket

[stratika]

Describe the bug

I tried yesterday to run the DFTMT example on our AMD Alveo 280 FPGA card. I noticed that there is an error in OpenCL related to the generated kernel and one extra curly bracket that is generated.

How To Reproduce

make BACKEND=opencl
tornado --jvm "-Ds0.t0.device=0:1 -Dtornado.fpga.conf.file=/home/thanos/repositories/TornadoVM/bin/sdk/etc/xilinx-fpga.conf -Xmx20g -Xms20g" --printKernel --threadInfo -m tornado.examples/uk.ac.manchester.tornado.examples.dynamic.DFTMT --params="256 default 1"

__kernel void computeDFT(__global long *_kernel_context, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics, __global uchar *inreal, __global uchar *inimag, __global uchar *outreal, __global uchar *outimag, __global uchar *inputSize)
{
  ulong ul_37, ul_36, ul_3, ul_14, ul_2, ul_1, ul_0, ul_16; 
  int i_32, i_33, i_4, i_6, i_38, i_5, i_10, i_11; 
  float f_8, f_7, f_9, f_15, f_18, f_17, f_20, f_19, f_22, f_21, f_24, f_23, f_26, f_25, f_28, f_27, f_30, f_29, f_31; 
  long l_13, l_12, l_35, l_34; 

  // BLOCK 0
  ul_0  =  (ulong) inreal;
  ul_1  =  (ulong) inimag;
  ul_2  =  (ulong) outreal;
  ul_3  =  (ulong) outimag;
  i_4  =  get_global_size(0);
  i_5  =  get_global_id(0);
  // BLOCK 1 MERGES [0 5 ]
  i_6  =  i_5;
  // BLOCK 2
  f_7  =  (float) i_6;
  // BLOCK 3 MERGES [2 4 ]
  f_8  =  0.0F;
  f_9  =  0.0F;
  i_10  =  0;
  __attribute__((xcl_pipeline_loop(1)))
  for(;i_10 < 256;)
  {
    // BLOCK 4
    i_11  =  i_10 + 6;
    l_12  =  (long) i_11;
    l_13  =  l_12 << 2;
    ul_14  =  ul_0 + l_13;
    f_15  =  *((__global float *) ul_14);
    ul_16  =  ul_1 + l_13;
    f_17  =  *((__global float *) ul_16);
    f_18  =  *((__global float *) ul_14);
    f_19  =  *((__global float *) ul_16);
    f_20  =  (float) i_10;
    f_21  =  f_20 * 6.2831855F;
    f_22  =  f_21 * f_7;
    f_23  =  f_22 / 256.0F;
    f_24  =  native_sin(f_23);
    f_25  =  native_cos(f_23);
    f_26  =  f_25 * f_19;
    f_27  =  fma(f_24, f_18, f_26);
    f_28  =  f_9 - f_27;
    f_29  =  f_17 * f_24;
    f_30  =  fma(f_25, f_15, f_29);
    f_31  =  f_8 + f_30;
    i_32  =  i_10 + 1;
    f_8  =  f_31;
    f_9  =  f_28;
    i_10  =  i_32;
  }  // B4
  // BLOCK 5
  i_33  =  i_6 + 6;
  l_34  =  (long) i_33;
  l_35  =  l_34 << 2;
  ul_36  =  ul_2 + l_35;
  *((__global float *) ul_36)  =  f_8;
  ul_37  =  ul_3 + l_35;
  *((__global float *) ul_37)  =  f_9;
  i_38  =  i_4 + i_6;
  i_6  =  i_38;
  // BLOCK 6
  return;
}  // B6
}  //  kernel.      <---- this is additional curly bracket

Expected behavior

The expected behavior is to generate kernels that do not have this syntax error.

Computing system setup (please complete the following information):

• OS: Ubuntu 22.04.5
• OpenCL and Driver versions: OpenCL 1.0, Alveo U280 (Vitis 2023.2)
• TornadoVM commit id: 0d44252

---

[jjfumero]

@mairooni can help with this issue

[stratika]

sure, I think I have found the commit that it started breaking. I will sync with @mairooni.

[stratika]

We synced with @mairooni. The changes in OCLBlockVisitor.java (here) seem to have caused the issue. I will have a look next week.