[CUDA] warning: linked binaries do not contain expected 'nvptx64-nvidia-cuda' target

[densamoilov]

When compiling a simple SYCL kernel for CUDA backend I get the following warning:

warning: linking module '/buildspace/dpcpp/llvm/build/install/lib/clang/16.0.0/../../clc/remangled-l64-signed_char.libspirv-nvptx64--nvidiacl.bc': Linking two modules of different target triples: '/buildspace/dpcpp/llvm/build/install/lib/clang/16.0.0/../../clc/remangled-l64-signed_char.libspirv-nvptx64--nvidiacl.bc' is 'nvptx64-unknown-nvidiacl' whereas 'main.cpp' is 'nvptx64-nvidia-cuda'
 [-Wlinker-warnings]

When I'm linking the object file I got I get the following warning:

clang-16: warning: linked binaries do not contain expected 'nvptx64-nvidia-cuda' target; found targets: 'nvptx64-nvidia-cuda-sm_50' [-Wsycl-target]

I know about the first, compile time warning but the second, link time warning is new. What does it mean and how can I fix it?

Reproducer:

#include <CL/sycl.hpp>

int main() {
    sycl::buffer<int> buf(1);
    auto dev = sycl::device(sycl::gpu_selector{});
    sycl::queue q(dev);

    q.submit([&](sycl::handler &cgh) {
        auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
        cgh.parallel_for(sycl::range<1>(1), [=](sycl::id<1> id) {
            acc[id] = 5;
        });
    });
    q.wait_and_throw();
    return 0;
}

I use DPC++ daily 2022-08-22.

[npmiller]

So this was introduced in #6591

And the issue is that for NVPTX and AMDGCN the bundles include the specific architecture, for example:

sycl-nvptx64-nvidia-cuda-sm_60

But that architecture is passed as an extra flag, for example:

-fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_60

But the code in #6591 is trying to match the triple in the bundles with the triple in -fsycl-targets, and so since that doesn't include the specific architecture it mismatches.

These extra flags are parsed further down during the BuildActions stages:

• https://github.com/intel/llvm/blob/sycl/clang/lib/Driver/Driver.cpp#L5477

Ideally we should check the specific architecture as well and also handle the default architecture for these targets in checkForOffloadMismatch, but I'm not too sure how best to approach this, whether we should move the parsing of the offload arch arguments closer to this check or duplicate it, or move the check further down.

@mdtoguchi any ideas on this?

[npmiller]

Both of these warnings should now be fixed:

• [SYCL][CUDA][HIP AMD] Correct cuda/hip spirv libclc .bc and module triple. #7245
• [SYCL][Driver] Move offload mismatch warning into SYCLActionBuilder #7434

[npmiller]

Closing this as the warnings have been fixed, please feel free to re-open if the issue persists