[SYCL][CUDA] `__CUDA_ARCH__` defined when compiling for CUDA backend since `sycl-nightly/20221129`

[AuroraPerego]

While trying to compile my application with different nightlies at some point the compilation started to fail.
I found out that the reason is that __CUDA_ARCH__ is defined, while it wasn't before and I was wondering whether this is wanted or not.

How to reproduce:

clang++ -fsycl -std=c++17 dummy.cpp -dM -E -fsycl-targets=nvptx64-nvidia-cuda -fno-bundle-offload-arch --cuda-path=/usr/local/cuda -Wno-unknown-cuda-version -Wno-linker-warnings | grep CUDA_ARCH

dummy.cpp is just:

#include <CL/sycl.hpp>

int main(){
    return 0;
}

• until the sycl-nightly/20221128 (clang version 16.0.0 (https://github.com/intel/llvm 85f2b97)) this command gives no output
• since the sycl-nightly/20221129 (clang version 16.0.0 (https://github.com/intel/llvm 142d9f5)) the output is:
  #define __CUDA_ARCH__ 500

The commit that led to this change is this.

• OS: Linux
• CUDA version: 11.7

[hdelan]

Can you share the error that occurred when compiling this dummy.cpp please?

[AuroraPerego]

This file compiles, it was just a quick way to see which macros are defined. What is not compiling is the project I'm working on right now that uses Eigen. Since __CUDA_ARCH__ is defined it falls in a #ifdef (here)were it should not go and raises this error:

In file included from /afs/cern.ch/user/a/aperego/public/pixeltrack-standalone/external/eigen/Eigen/Core:170:
/afs/cern.ch/user/a/aperego/public/pixeltrack-standalone/external/eigen/Eigen/src/Core/MathFunctions.h:988:15: fatal error: no member named 'isfinite' in the global namespace
    return (::isfinite)(x);
            ~~^

If you want I can share the steps to reproduce this error, but it's directly related to the fact the __CUDA_ARCH__ is defined and wasn't happening before sycl-nightly/20221129. If this is expected I'll pass to Eigen a flag that ignores __CUDA_ARCH__ to be able to compile anyway.

[AuroraPerego]

The definition of __CUDA_ARCH__ has been likely introduced by #6524 via https://github.com/intel/llvm/pull/6524/files#diff-84691cf70860f37e55bcd6c5275edb7d31c60e43d9c9f310a2356be3a78e8418

[hdelan]

__CUDA_ARCH__ plays a role in SYCL to allow for compile time branching of device code for CUDA backend. A lot of code relies on this. So it seems that Eigen/src/Core/util/Macros.h needs to be updated to allow for __CUDA_ARCH__ to be defined in SYCL as well. ie it would be best if instead of __CUDA_ARCH__ suggesting that the CUDA API is being used, it might be worthwhile to check that __SYCL_DEVICE_ONLY__ is also not defined.

[fwyzard]

In terms of "don't break existing code", wouldn't it be a better approach to not define __CUDA_ARCH__ during SYCL device compilation, and instead use a different preprocessor macro to identify the SYCL target architecture ?

[abagusetty]

This might likely be an alternative: https://github.com/intel/llvm/blob/e5de913f884acd9566449c7e8c72827cb6009a2b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_architecture.asciidoc

[fwyzard]

@abagusetty thanks for the link.

I do like the idea of having a dedicated macro for the SYCL backends, instead of reusing __CUDA_ARCH__.

However, if the possible values are defined by an enum, I don't see how they could be used in preprocessor conditional compilation ?

Of course they can be used with if constexpr, traits, etc. - but sometimes you really want to #if something in or out.

Maybe the extensions could be made more useful by also defining some corresponding macro names ?
Like SYCL_EXT_ONEAPI_DEVICE_ARCHITECTURE_NVIDIA_GPU_SM_50 with the same value as sycl::ext::oneapi::experimental::architecture::nvidia_gpu_sm_50, etc..

And some aliases to identify the first and last entry of each architecture, so one could do

if constexpr(SYCL_EXT_ONEAPI_DEVICE_ARCHITECTURE >= architecture::nvidia_gpu_first && SYCL_EXT_ONEAPI_DEVICE_ARCHITECTURE <= architecture::nvidia_gpu_last) {
  ...

etc.

[GeorgeWeb]

Hi @AuroraPerego

We expect this issue to soon be resolved with the changes from PR #8257.

TLDR: We define a unique __SYCL_CUDA_ARCH__ macro and use it instead for SYCL device targets and leave __CUDA_ARCH__ as before for CUDA targets.