l.0-parallel reduction codegen mishandles simul_reduce

As reported by @kaushikcfd on Riot:

import loopy as lp

knl = lp.make_kernel(
        "{[i]: 0<=i<4}",
        """
        a = simul_reduce(sum, i, 7*i)
        b = simul_reduce(sum, i, 10*i)
        """)

knl = lp.tag_inames(knl, "i:l.0")
knl = lp.realize_reduction(knl)

print(lp.generate_code_v2(knl).device_code())

generates two sets of inames.

cc @mattwala

changed the description

There are 2 parts to this problem --

1. map_reduction_local should generate same sets of intermediate inames for the simul_reduce And so after realize_reduction the kernel in the description should become --

---------------------------------------------------------------------------
INSTRUCTIONS:
    for red_init_i
↱     acc_i[red_init_i] = 0  {id=insn_i_init}
│↱    neutral_i = 0  {id=insn_i_init_neutral}
││  end red_init_i
││  for i
└└↱   acc_i[i] = neutral_i + 7*i  {id=insn_i_transfer}
  │ end i
  │ for red_i_s0_0
↱ └   acc_i[red_i_s0_0] = acc_i[red_i_s0_0] + acc_i[red_i_s0_0 + 2]  {id=red_i_stage_0}
│   end red_i_s0_0
│   for red_i_s1_0
└↱    acc_i[red_i_s1_0] = acc_i[red_i_s1_0] + acc_i[red_i_s1_0 + 1]  {id=red_i_stage_1}
 │  end red_i_s1_0
 │  for red_i
 └    a = acc_i[0]  {id=insn_1}
    end red_i
    for red_init_i
↱     acc_i_0[red_init_i] = 0  {id=insn_0_i_init}
│↱    neutral_i_0 = 0  {id=insn_0_i_init_neutral}
││  end red_init_i
││  for i
└└↱   acc_i_0[i] = neutral_i_0 + 10*i  {id=insn_0_i_transfer}
  │ end i
  │ for red_i_s0_0
↱ └   acc_i_0[red_i_s0_0] = acc_i_0[red_i_s0_0] + acc_i_0[red_i_s0_0 + 2]  {id=red_i_stage_0_0}
│   end red_i_s0_0
│   for red_i_s1_0
└↱    acc_i_0[red_i_s1_0] = acc_i_0[red_i_s1_0] + acc_i_0[red_i_s1_0 + 1]  {id=red_i_stage_1_0}
 │  end red_i_s1_0
 │  for red_i
 └    b = acc_i_0[0]  {id=insn_0_0}
    end red_i
---------------------------------------------------------------------------

2. Currently, the above improvised kernel obtained after step 1 generates a sub-optimal code(notice the unnecessary local barriers):

__kernel void __attribute__ ((reqd_work_group_size(4, 1, 1))) loopy_kernel(__global int *__restrict__ a, __global int *__restrict__ b)
{
  __local int acc_i[4];
  __local int acc_i_0[4];
  int neutral_i;
  int neutral_i_0;

  neutral_i = 0;
  acc_i[lid(0)] = 0;
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i (insn_i_transfer depends on insn_i_init) */;
  acc_i[lid(0)] = neutral_i + 7 * lid(0);
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i (red_i_stage_0 depends on insn_i_transfer) */;
  if (1 + -1 * lid(0) >= 0)
    acc_i[lid(0)] = acc_i[lid(0)] + acc_i[2 + lid(0)];
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i (red_i_stage_1 depends on red_i_stage_0) */;
  if (lid(0) == 0)
    acc_i[0] = acc_i[0] + acc_i[1];
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i (insn_1 depends on red_i_stage_1) */;
  a[0] = acc_i[0];
  neutral_i_0 = 0;
  acc_i_0[lid(0)] = 0;
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i_0 (insn_0_i_transfer depends on insn_0_i_init) */;
  acc_i_0[lid(0)] = neutral_i_0 + 10 * lid(0);
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i_0 (red_i_stage_0_0 depends on insn_0_i_transfer) */;
  if (1 + -1 * lid(0) >= 0)
    acc_i_0[lid(0)] = acc_i_0[lid(0)] + acc_i_0[2 + lid(0)];
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i_0 (red_i_stage_1_0 depends on red_i_stage_0_0) */;
  if (lid(0) == 0)
    acc_i_0[0] = acc_i_0[0] + acc_i_0[1];
  barrier(CLK_LOCAL_MEM_FENCE) /* for acc_i_0 (insn_0_0 depends on red_i_stage_1_0) */;
  b[0] = acc_i_0[0];
}

I agree with Kaushik. I don't think there is any special handling of simul_reduce as opposed to reduce right now.

mentioned in merge request !316