KhronosGroup / KhronosGroup/SYCL-CTS
Math tests are susceptible to constant folding
- 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
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