Math tests are susceptible to constant folding

[npmiller]

Problem

A lot of the math tests use kernels that look like this (using FMA as an example but it also applies to other built-ins):

      [=] {
        float inputData_0(0.7879067377f);
        float inputData_1(0.8543576171f);
        float inputData_2(0.3422439025f);

        return sycl::fma(inputData_0, inputData_1, inputData_2);
      },

This is taken directly from the generated build/tests/math_builtin_api/math_builtin_float_base_1.cpp, just slightly reformatted and static_assert removed for brevity.

This pattern is problematic because it makes it very straightforward for the compiler to do constant folding and simply compute the result during compilation.

Sure enough looking at the LLVM IR generated by DPC++ for the FMA operation in this test, we get:

  store float 0x3FF03F1200000000, ptr %arrayidx.ascast.i.i, align 4, !tbaa !16

The FMA operation was computed by the compiler and the kernel ends up simply storing the result to memory.

What this means is that it is currently possible to pass a lot of the math tests without ever actually doing any math on the device.

Now it can be worthwhile to check that the compiler's constant folding gives correct results, but I'm not sure if that's what these tests are intending to do.

Possible workaround

At least with current DPC++ a possible workaround is to move the constants out of the kernel lambda and simply have them get captured, that way we generate the following IR for the FMA operation:

  %1 = tail call noundef float @llvm.fma.f32(float %__SYCLKernel.sroa.10.32.copyload, float %__SYCLKernel.sroa.12.32.copyload, float %__SYCLKernel.sroa.13.32.copyload)

Which means the kernel will actually perform a FMA operation on the device.

This might be straightforward to do in the current tests, however I'm not 100% sure if it's enough as it could theoretically be possible for the compiler to still find these constants and do constant folding, even though it's not able to do it right now.

[bader]

Use something like (x86) rdtsc or user input to get the seed. Use the seed to generate random input for tested kernels. If test fails, print the seed to reproduce the input.