[BUG] Multiple conditions evaluation results in wrong code gen

[jjfumero]

Describe the bug

The generated code omits a condition when having a complex if condition:

The result code for OpenCL is as follows:

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable  
__kernel void testIfInt6(__global uchar *_heap_base, ulong _frame_base, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics)
{
  int i_9, i_8, i_7, i_11, i_1, i_2; 
  bool z_10; 
  ulong ul_6, ul_0; 
  long l_3, l_5, l_4; 

  __global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];


  // BLOCK 0
  ul_0  =  (ulong) _frame[3];
  i_1  =  get_global_id(0);
  // BLOCK 1 MERGES [0 5 ]
  i_2  =  i_1;
  for(;i_2 < 256;)  {
    // BLOCK 2
    l_3  =  (long) i_2;
    l_4  =  l_3 << 2;
    l_5  =  l_4 + 24L;
    ul_6  =  ul_0 + l_5;
    i_7  =  *((__global int *) ul_6);
    i_8  =  get_global_size(0);
    i_9  =  i_8 + i_2;
    z_10  =  i_7 < 2;
    if(z_10)
    {
      // BLOCK 3
      *((__global int *) ul_6)  =  100;
      i_7  =  100;
    }  // B3
    else
    {
      // BLOCK 4
      *((__global int *) ul_6)  =  200;
      i_7  =  200;
    }  // B4
    // BLOCK 5 MERGES [3 4 ]
    i_11  =  i_9;
    i_2  =  i_11;
  }  // B5
  // BLOCK 6
  return;
}  //  kernel

How To Reproduce

This test case reproduces the error.

    public static void testIfInt6(int[] a) {
        for (@Parallel int i = 0; i < a.length; i++) {
            if (a[i] >= 0 && a[i] <= 1) {
                a[i] = 100;
            } else {
                a[i] = 200;
            }
        }
    }


 @Test
    public void test06() {
        final int numElements = 256;
        int[] a = new int[numElements];
        int[] expectedResult = new int[numElements];

        Arrays.fill(a, -1);
        Arrays.fill(expectedResult, 200);

        new TaskSchedule("s0") //
                .task("t0", TestKernels::testIfInt6, a) //
                .streamOut(a) //
                .execute(); //

        assertArrayEquals(expectedResult, a);
    }

The problem is that during one of the low-tier compiler phases the multiple conditions is removed. I suspect this is during one of the canonicalization phases.

[gsaxena888]

Has this issue been resolved? (We have various math statements, eg., "max", "min" that presumably use if conditions underneath and we may have similar or more complex if conditions as described above.)

[stratika]

hi @gsaxena888, this bug is still present. We have a unit-test that breaks:

tornado -ea  --printKernel --jvm "-Xmx6g -Dtornado.recover.bailout=False -Dtornado.unittests.verbose=True "  -m  tornado.unittests/uk.ac.manchester.tornado.unittests.tools.TornadoTestRunner  --params "uk.ac.manchester.tornado.unittests.foundation.TestIf#test06"

I am wondering if with the latest compiler updates, we still observe similar behaviour as described above by @jjfumero. We will have a look, and report back.

[jjfumero]

Yes, this bug is still present for the OpenCL backend. but it is fixed for the SPIR-V and PTX backends. The reason the OpenCL is more complicated is due to the uncontrolled flow graph to control the flow graph during the code generation of the OpenCL C code.

We want to fully transition to SPIR-V code to be dispatched through the OpenCL runtime, so our OpenCL C code will be deprecated, and only used on old OpenCL driver implementations (FPGAs, and Apple M1/M2/M3).