BugReport: CInstructions in inlining and vectorisation

[sv2518]

Hello!

It seems like CInstructions are not treated correctly in the inlining and vectorisation process.

I have written an iterative solver in loopy. This loopy kernel is used in other kernels in the Firedrake codebase. Because loopy does not support while loops I am using a for-loop and a stopping criterion. Because loopy does not support a break statement, I am using a CInstruction to generate the stop criterion. I ran into trouble in two cases.

• In the inlining because the variable names in the criterion do not change (but all other variables in the kernel are)
• In the vectorisation because we need to linearise the stop criterion variable (I was doing this by hand previously but I think it should be loopy's job)

I spent all day writing a minimal failing example, but I am stuck with a scheduling error. I don't get the scheduling error in the Firedrake codebase, so I am doing something wrong in the construction of the MFE. The MFE is here https://gist.github.com/sv2518/3e1cbc3bd894320bc4dcabfeb3c86457

[sv2518]

It would also be great if you could support Cinstructions in the lp.t_unit_to_python function but that's a separate issue

[kaushikcfd]

Thanks for the report + MFE!

In the inlining because the variable names in the criterion do not change (but all other variables in the kernel are)

The inline was struggling because the predicates were expressed as pure-strings over which performing manipulation isn't supported in loopy. Instead, initializing the predicate expressions as shown in the regression of #635 should support inline-renaming.

In the vectorization because we need to linearise the stop criterion variable (I was doing this by hand previously but I think it should be loopy's job)

I think if written using CInstruction.predicates the vectorization logic should be happy as well.

It would also be great if you could support Cinstructions in the lp.t_unit_to_python function but that's a separate issue

This is a bit involved in the way the logic there is implemented over there. I don't have a quick-fix for that yet.

[sv2518]

Can confirm that the inlining logic has been fixed, but we need to do some more work on the vectorisation
We currently generate something like the following where crit is a vectorised type

      if (crit < 1e-08)
      {
        #pragma omp simd
        for (int32_t n_shift_batch = 0; n_shift_batch <= 3; ++n_shift_batch)
        {
          break;
        }
      }

[sv2518]

Fixed as part of #557