Compilation Error with C++20 and CUDA Target (Windows)

[blinkfrog]

Compilation Error with C++20 and CUDA Target (Windows)

[fodinabor]

Just a note - you can usually edit the PR description&branch from the GUI, for rewriting the history: force push to your branch :)

Also: it's common to leave the issue open until it's fixed (and automatically closed) by merging the PR.

[nilsfriess]

It seems that MSVC expects std::fma(long double, long double, long double) to be implemented when compiling with -std=c++20 (even when it's not used anywhere) but this overload is not defined when compiling for CUDA using clang.
I searched through the LLVM issues to see if there are similar problems and actually found this issue by @fodinabor which seems to be basically the same problem.

Maybe it's sufficient to just declare this overload somewhere in your code

namespace std {
    __device__ long double fma(long double x, long double y, long double z);
}

At least Clang does something similar for other functions that need long double overloads that do not exists in CUDA.

[blinkfrog]

Thank you for the suggestion, @nilsfriess. Unfortunately, the solution doesn't seem to work in my case. Here's a brief summary of my attempts:

Initial Attempt:

I added the fma declaration you suggested to the beginning of my source file, before any includes.

namespace std {
    __device__ long double fma(long double x, long double y, long double z);
}

However, this approach did not resolve the issue.

I also tried to use syntax from that clang source file, but this didn't work too:

namespace std {
    __attribute__((device)) long double fma(long double, long double, long double);
}

Further Experimentation:

I discovered that adding the same lines directly to "C:\llvm\lib\clang\17\include\cuda_wrappers\cmath", specifically before the #include_next directive, resolves the error. This suggests the issue arises before my source files are processed.

Attempt with Implementation:

To confirm the order of processing, I added an implementation for fma both in my source and in the Clang header. The resulting redefinition error confirms that Clang's is processed before my code:

C:/Temp/sycltest/SyclCmakeProject5Fract/blur.cpp:11:37: error: redefinition of 'fma'
   11 | __attribute__((device)) long double fma(long double x, long double y, long double z) {return z};
      |                                     ^
C:\llvm\lib\clang\17\include\cuda_wrappers\cmath:29:37: note: previous definition is here
   29 | __attribute__((device)) long double fma(long double x, long double y, long double z) {return z;};
      |                                     ^
C:/Temp/sycltest/SyclCmakeProject5Fract/blur.cpp:11:95: error: expected ';' after return statement
   11 | __attribute__((device)) long double fma(long double x, long double y, long double z) {return z};

Given these observations, it seems the original problem occurs at a stage where my source file declarations are not yet considered. This is puzzling, as my understanding is that declarations in my source file should be processed prior to system headers.

Do you have any insights into why this might be happening, or any further suggestions to address the issue?

[fodinabor]

This ordering is a pretty delicate matter.. I think it might be important that we're including cmath beforehands or so..?

We have our own special header files for this which we include with -isystem $HIPSYCL_PATH/include/AdaptiveCpp/hipSYCL/std/hiplike this makes this file load early enough..

https://github.com/AdaptiveCpp/AdaptiveCpp/blob/develop/include/hipSYCL/std/hiplike/complex

So what happens if you adapt this file?
I believe clang CUDA force includes a few of those built-in files, so hooking into the correct header is important..
Maybe you actually need to add a file cmath to.that folder and add your workaround there..

[blinkfrog]

Thank you. Modifying complex didn't work (it looks like it isn't processed at all when compiling my project), however, creating cmath in hiplike directory worked!

This is the content of this file:

#if defined(__clang__) && (defined(__CUDACC__) || defined(__HIP__))
#if defined(_WIN32)
namespace std {
  __attribute__((device))
  long double fma(long double, long double, long double);
}
#endif
#endif

#include_next <cmath>

If needed, I can create a PR with this file added.

[fodinabor]

Perfect. I assume a PR for this should be great.

[nilsfriess]

But I guess this should also reported in LLVM? Because this is basically a Clang/MSVC compatibility issue and I guess every program compiled for CUDA with clang+msvc should have the same problem

[blinkfrog]

Created a PR (from a third attempt as I messed things up in the first two attempts 🤦 )

[blinkfrog]

Thank you, will remember this next time.