[HIP] DPC++ does not use correctly rounded sqrt, even when using -fno-fast-math

[illuhad]

Describe the bug

DPC++ does not use use correctly rounded sqrt() on AMD GPUs, even if -fno-fast-math is explicitly passed.

This is contrary to the behavior of both hipcc and AdaptiveCpp which by default correctly round sqrt, and can lead to misleading benchmark results, or at least make them difficult to interpret.

To reproduce

1. Print compiler invocation e.g. using icpx -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx906 /dev/null -fno-fast-math -###
2. Observe in the output that it links with the bitcode library oclc_correctly_rounded_sqrt_off.bc

Environment

This was observed on Linux with oneAPI 2024.0.2.

Additional context

I suspect that oclc_correctly_rounded_sqrt_off is simply the default due to all ROCm bitcode library configuration knobs being initialized with false (daz = off, finite_only = off, unsafe_math = off, correctly_rounded_sqrt = off).

correctly_rounded_sqrt however may have to be treated differently, because it is the only one where a setting of off does not correspond to precision.

[pranav-sivaraman]

Investigating into this a bit more I found that you need to pass in the following flag -fsycl-fp32-prec-sqrt to get the correctly round sqrt. It seems this is also needed for the CUDA backend as well according to the Getting Started guide.

[illuhad]

@pranav-sivaraman Thanks for looking into this! I still think it is pretty misleading and unintuitive that this is happening in spite of -fno-fast-math, which most users will interpret to communicate the intent of precise arithmetic.

[npmiller]

Is this for sqrt or sycl::sqrt?

SYCL built-ins follow the precision requirements of OpenCL, so they don't require sycl::sqrt to be correctly rounded, and given that -fno-fast-math is the default for clang and that using the correctly rounded sqrt gives a significant performance hit, I think the current behavior is correct for sycl::sqrt.

For the cmath/std:: sqrt maybe that should be different, although we'd have to investigate if we can have both with different precision.

We had some discussions around the sqrt precision when the SYCL flag was added:

• [CUDA] sycl::sqrt leads to IEEE754 incompatible results on NVidia cards #4041
• [SYCL][CUDA] Add -fcuda-prec-sqrt flag #5141
• [SYCL] Add -fsycl-fp32-prec-sqrt flag #5309

[illuhad]

My observation is for sycl::sqrt. It may be that -fno-fast-math is the default for clang, but it is not the default for icpx. So when a user explicitly passes -fno-fast-math to icpx, they do communicate the intent to do something different.

I don't want to reopen the discussion about the rationale behind this default if you've already discussed this at length, but my two cents are:

• The behavior is surprising because other heterogeneous compilers that DPC++ may be compared against do not behave this way. This includes native compilers like hipcc and other SYCL compilers like AdaptiveCpp.
• It is therefore almost guaranteed that performance comparisons are going to end up with misleading results.
• People are less likely to compare to OpenCL, which as a separate source model, fits a different niche. So the OpenCL behavior does not matter much.
• Just because the SYCL specification technically allows something, it does not mean that it is the right thing to do.