KhronosGroup / KhronosGroup/SYCL-CTS

Math tests are susceptible to constant folding

Open
#1,070 1 comment 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
75
Forks
96
Avg merge
5d 18h
Merged PRs (30d)
4

Description

### 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):

```cpp
[=] {
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:

```llvm
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:

```llvm
%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.

Contributor guide

Open the contributing guide

Research direction

Start with the generated tests in build/tests/math_builtin_api/math_builtin_float_base_1.cpp and inspect the LLVM IR produced for the FMA and other math kernels. Compare the current constant-in-kernel pattern with captured inputs, then verify that the resulting kernels retain device math operations rather than folded constants. Done means the affected math tests cannot pass without executing the built-ins on the device.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
testing-qa
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.