Add an optional diagnostic for the use of double-precision ops in kernels

[al42and]

Is your feature request related to a problem? Please describe

In some (many?) cases, the kernels are intended to be run on a GPU, most of which have reduced FP64 performance compared to FP32. However, it's very easy in C++ to accidentally introduce double-precision arithmetic.

Such issues can be detected with a profiler, but it might be helpful to have an opt-in, compile-time diagnostic to warn about double-precision ops in the device kernels. No warning is to be emitted unless the user explicitly requested it.

Toy example (based on real cases). Excerpt:

  sycl::event ev = queue.submit([&](sycl::handler &cgh) {
    auto buffer_dev = buffer.get_access<sycl::access_mode::write>(cgh);
    cgh.parallel_for<class DummyKernel>(
        sycl::range<1>{numThreads}, [=](sycl::id<1> threadId) {
          int x = threadId.get(0);
          float x2 = x * 2.0;      // Forgot F suffix
          float y = fma(x, 2, x2); // Used fma instead of sycl::fma<float>
          buffer_dev[threadId] = y;
        });
  });

Part of the SPIR-V output (clang++ -fsycl -ffast-math -Wall simple.cpp -O3 -o simple && SYCL_DUMP_IMAGES=1 ./simple):

   %conv_i_i = OpUConvert %uint %33
  %conv2_i_i = OpConvertSToF %double %conv_i_i
    %mul_i_i = OpFMul %double %conv2_i_i %double_2
  %conv3_i_i = OpFConvert %float %mul_i_i
  %conv5_i_i = OpFConvert %double %conv3_i_i
         %47 = OpExtInst %double %1 fma %conv2_i_i %double_2 %conv5_i_i
  %conv6_i_i = OpFConvert %float %47

So, we're getting unnecessary conversions between FP64 and FP32, and FP64 ops when FP32 were probably intended.

Describe the solution you would like

Have an optional -Wsycl-fp64-on-device warning switch.

Describe alternatives you have considered

• A kernel attribute enabling the diagnostic for a single kernel?
• Things like float x2 = x * 2.0 probably can get be optimized in the ffast-math mode?

Additional context

The example above is compiled correctly, so it's not a bug. But it is likely not what the user intended, so a diagnostic would be a quick way to catch such problems without even having to run a profiler.

[Michoumichmich]

About the emulated ops, the fact that SYCL returns size_t (thus 64 bit) from most of the space queries causes unintentional performance penalties on CUDA (a few percent, but it's enough to get the same performance as "native" CUDA sometimes). So maybe a more general warning about code that leads to emulated types ?

As for the specific case of float x2 = x * 2.0; it looks more like a compiler bug. It should use fp32 as the results is exactly the same. Even GCC without fast-math uses fp32: https://godbolt.org/z/GeeTGEvdb

Isn't -Wc++11-narrowing enough? You could also pass -Xcuda-ptxas --warn-on-double-precision-use to clang++

[bader]

FYI: https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_API.html#math-intrinsics-options

OpenCL proposes alternative solution for this problem:

-cl-single-precision-constant - This option forces implicit conversions of double-precision floating-point literals to single precision. This option is ignored for programs created with IL.

[al42and]

So maybe a more general warning about code that leads to emulated types ?

That would also be useful, although probably prone to false positives

You could also pass -Xcuda-ptxas --warn-on-double-precision-use to clang++

@Michoumichmich, thanks, that's a good suggestion (which I'll add to our CUDA code), but it only applies when targeting the CUDA backend. Overall, I feel that nvcc is much more aggressive than clang in doing the double-to-float narrowing.

Isn't -Wc++11-narrowing enough?

With 962417d, the example from the issue description does not trigger any warnings with -Wc++11-narrowing -Wall -Wextra.

-cl-single-precision-constant - This option forces implicit conversions of double-precision floating-point literals to single precision. This option is ignored for programs created with IL.

@bader, thanks. A similar flag could indeed be helpful, although I think a general warning, similar to CUDA's --warn-on-double-precision-use, would be even more useful (e.g., the fma part of the example). Such a warning seems to be easy to implement and widely applicable.

[srividya-sundaram]

@bader @AlexeySachkov I have a WIP PR which emits a warning when double literals are used in device code.
Do we also need to address the follwing cases?

• Spelling the type out (double, but also, long double?)
• Calling a builtin which returns a double (these don’t have function signatures that would have the type spelled out), like __builtin_fabs
• Default argument promotion in a variadic function call. e.g., void func(int i, ….); func(1, 1.0f); the 1.0f will be promoted to double when passed to func()

I did learn that in OpenCL, double precision is supported by the device, for OpenCL C 3.0 or newer the {opencl_c_fp64} feature macro is present. But, It is not clear from the SYCL spec or Clanf code base if device code supports variadic calls or builtins.

Can either of you or SYCL language design folks please clarify 2 and 3 above.

Thank you!

[bader]

@srividya-sundaram, I don't think that requested compiler feature should be designed by the language design folks. I think it's different from diagnosing the use of unsupported optional core feature.

Based on this description:

In some (many?) cases, the kernels are intended to be run on a GPU, most of which have reduced FP64 performance compared to FP32. However, it's very easy in C++ to accidentally introduce double-precision arithmetic.

Such issues can be detected with a profiler, but it might be helpful to have an opt-in, compile-time diagnostic to warn about double-precision ops in the device kernels. No warning is to be emitted unless the user explicitly requested it.

I conclude that if emitted code has operations with double LLVM type, there must be a warning.

It's not clear if each spelling of double or long double leads to double precision FP operations (e.g. passing double * around), so reporting a warning for each spelling might produce false alarms and should be accurate in this case.

It's clear that cases 2 and 3 produces either direct use of doubles or conversion from/to doubles, which are exactly the cases compiler is expected to detect as shown in the example from the description.

WRT OpenCL diagnostics, according to my understanding it aims to diagnose the case when the code with doubles are compiled for the target, which doesn't support doubles. The option requested in this issue should provide diagnostics for the targets supporting double precision FP operations.

[srividya-sundaram]

@bader Thank you for the clarification. Could you please take a look at PR6323 ?

[al42and]

#6323 allows detecting the problematic code I had in mind, so this can be closed.

[al42and]

It's not just me making this mistake :) #7277