[SYCL] Target ordering breaks compilation

[Michoumichmich]

Describe the bug
When build for CUDA and CPU targets with -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice,nvptx64-nvidia-cuda-sycldevice everything works fine. But when targets are swapped there's an error.

To Reproduce

With SPIR64 first, it works

clang++ -fsycl -fsycl-unnamed-lambda  -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice,nvptx64-nvidia-cuda-sycldevice matrix_mult.dp.cpp
clang-13: warning: Unknown CUDA version. cuda.h: CUDA_VERSION=11030. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
Platform name: Intel(R) OpenCL
[...]
OpenCL program binary file was successfully created: /tmp/matrix_mult-242693.out

But with CUDA first it does not find the library.

clang++ -fsycl -fsycl-unnamed-lambda  -fsycl-targets=nvptx64-nvidia-cuda-sycldevice,spir64_x86_64-unknown-unknown-sycldevice matrix_mult.dp.cpp           
clang-13: warning: Unknown CUDA version. cuda.h: CUDA_VERSION=11030. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
clang-13: error: cannot find libdevice for . Provide path to different CUDA installation via --cuda-path, or pass -nocudalib to build without linking with libdevice.

I'm aware that I'm using an unsupported version of CUDA, but still think the behaviour is strange. Especially a week ago this same swap was causing the compiler to crash.

Environment:

• OS: RHEL 8-3
• Target device and vendor: Intel CPU, Nvidia GTX 1660 Ti (sm_75)
• DPC++ version: clang version 13.0.0 (https://github.com/intel/llvm b52f203)
• Dependencies version:
  [opencl:0] CPU : Intel(R) OpenCL 2.1 [2021.11.3.0.17_160000]
  [cuda:0] GPU : NVIDIA CUDA BACKEND 0.0 [CUDA 11.3]
  [host:0] HOST: SYCL host platform 1.2 [1.2]

[kbobrovs]

@mdtoguchi, FYI. AFAIK you are about to do some refactoring in the driver anyway, so maybe you could take a look at this one too.

[AidanBeltonS]

The error appears to be introduced due to a bad set of action inputs
While the error within the program is introduced due to the collapse of an offloading preprocessor action in BuildJobsForActionNoCache
the root issue is that during the compilation phase the bounded archs and triples become swapped.
This does not cause an error with spir64 as it does not use the bounded arch, however with nvptx an error is thrown as it does
need to know the gpu's arch.

NVPTX,SPIRV action's

                     +- 0: input, "subbuffer.cpp", c++, (host-sycl)
                  +- 1: append-footer, {0}, c++, (host-sycl)
               +- 2: preprocessor, {1}, c++-cpp-output, (host-sycl)
               |     +- 3: input, "subbuffer.cpp", c++, (device-sycl, sm_50)
               |  +- 4: preprocessor, {3}, c++-cpp-output, (device-sycl, sm_50)
               |- 5: compiler, {4}, ir, (device-sycl, sm_50)
            +- 6: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (spir64-unknown-unknown:sm_50)" {5}, c++-cpp-output
         +- 7: compiler, {6}, ir, (host-sycl)
      +- 8: backend, {7}, assembler, (host-sycl)

SPIRV,NVPTX action's

                     +- 0: input, "subbuffer.cpp", c++, (host-sycl)
                  +- 1: append-footer, {0}, c++, (host-sycl)
               +- 2: preprocessor, {1}, c++-cpp-output, (host-sycl)
               |     +- 3: input, "subbuffer.cpp", c++, (device-sycl)
               |  +- 4: preprocessor, {3}, c++-cpp-output, (device-sycl)
               |- 5: compiler, {4}, ir, (device-sycl)
            +- 6: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (nvptx64-nvidia-cuda)" {5}, c++-cpp-output            ---------Here, need bound arch
         +- 7: compiler, {6}, ir, (host-sycl)
      +- 8: backend, {7}, assembler, (host-sycl)

Full phase output NVPTX, SPIRV

NVPTX,SPIRV

clang version 14.0.0 (https://github.com/AidanBeltonS/llvm d4e18d9858a620e498cb87d1c7503ff6d3bc4ee2)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/aidan/llvm/build/bin
                     +- 0: input, "subbuffer.cpp", c++, (host-sycl)
                  +- 1: append-footer, {0}, c++, (host-sycl)
               +- 2: preprocessor, {1}, c++-cpp-output, (host-sycl)
               |     +- 3: input, "subbuffer.cpp", c++, (device-sycl, sm_50)
               |  +- 4: preprocessor, {3}, c++-cpp-output, (device-sycl, sm_50)
               |- 5: compiler, {4}, ir, (device-sycl, sm_50)
            +- 6: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (spir64-unknown-unknown:sm_50)" {5}, c++-cpp-output
         +- 7: compiler, {6}, ir, (host-sycl)
      +- 8: backend, {7}, assembler, (host-sycl)
   +- 9: assembler, {8}, object, (host-sycl)
+- 10: linker, {9}, image, (host-sycl)
|                    +- 11: input, "subbuffer.cpp", c++, (device-sycl)
|                 +- 12: preprocessor, {11}, c++-cpp-output, (device-sycl)
|              +- 13: compiler, {12}, ir, (device-sycl)
|           +- 14: linker, {13}, ir, (device-sycl)
|           |  +- 15: input, "/home/aidan/llvm/build/bin/../lib/libsycl-crt.o", object
|           |- 16: clang-offload-unbundler, {15}, object
|           |  +- 17: input, "/home/aidan/llvm/build/bin/../lib/libsycl-complex.o", object
|           |- 18: clang-offload-unbundler, {17}, object
|           |  +- 19: input, "/home/aidan/llvm/build/bin/../lib/libsycl-complex-fp64.o", object
|           |- 20: clang-offload-unbundler, {19}, object
|           |  +- 21: input, "/home/aidan/llvm/build/bin/../lib/libsycl-cmath.o", object
|           |- 22: clang-offload-unbundler, {21}, object
|           |  +- 23: input, "/home/aidan/llvm/build/bin/../lib/libsycl-cmath-fp64.o", object
|           |- 24: clang-offload-unbundler, {23}, object
|           |  +- 25: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cassert.o", object
|           |- 26: clang-offload-unbundler, {25}, object
|           |  +- 27: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cstring.o", object
|           |- 28: clang-offload-unbundler, {27}, object
|           |  +- 29: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-complex.o", object
|           |- 30: clang-offload-unbundler, {29}, object
|           |  +- 31: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-complex-fp64.o", object
|           |- 32: clang-offload-unbundler, {31}, object
|           |  +- 33: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cmath.o", object
|           |- 34: clang-offload-unbundler, {33}, object
|           |  +- 35: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cmath-fp64.o", object
|           |- 36: clang-offload-unbundler, {35}, object
|        +- 37: linker, {14, 16, 18, 20, 22, 24, 26, 28, 30, 32, 34, 36}, ir, (device-sycl)
|     +- 38: sycl-post-link, {37}, tempfiletable, (device-sycl)
|     |  +- 39: file-table-tform, {38}, tempfilelist, (device-sycl)
|     |- 40: llvm-spirv, {39}, tempfilelist, (device-sycl)
|  +- 41: file-table-tform, {38, 40}, tempfiletable, (device-sycl)
|- 42: clang-offload-wrapper, {41}, object, (device-sycl)
|        +- 43: linker, {5}, ir, (device-sycl, sm_50)
|     +- 44: sycl-post-link, {43}, ir, (device-sycl, sm_50)
|     |  +- 45: file-table-tform, {44}, ir, (device-sycl, sm_50)
|     |  |  +- 46: backend, {45}, assembler, (device-sycl, sm_50)
|     |  |  |- 47: assembler, {46}, object, (device-sycl, sm_50)
|     |  |- 48: linker, {46, 47}, cuda-fatbin, (device-sycl, sm_50)
|     |- 49: foreach, {45, 48}, cuda-fatbin, (device-sycl, sm_50)
|  +- 50: file-table-tform, {44, 49}, tempfiletable, (device-sycl, sm_50)
|- 51: clang-offload-wrapper, {50}, object, (device-sycl, sm_50)
52: offload, "host-sycl (x86_64-unknown-linux-gnu)" {10}, "device-sycl (spir64-unknown-unknown)" {42}, "device-sycl (nvptx64-nvidia-cuda:sm_50)" {51}, image

Full phase output SPIRV, NVPTX

SPIRV,NVPTX

                     +- 0: input, "subbuffer.cpp", c++, (host-sycl)
                  +- 1: append-footer, {0}, c++, (host-sycl)
               +- 2: preprocessor, {1}, c++-cpp-output, (host-sycl)
               |     +- 3: input, "subbuffer.cpp", c++, (device-sycl)
               |  +- 4: preprocessor, {3}, c++-cpp-output, (device-sycl)
               |- 5: compiler, {4}, ir, (device-sycl)
            +- 6: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (nvptx64-nvidia-cuda)" {5}, c++-cpp-output            ---------Here, need bound arch
         +- 7: compiler, {6}, ir, (host-sycl)
      +- 8: backend, {7}, assembler, (host-sycl)
   +- 9: assembler, {8}, object, (host-sycl)
+- 10: linker, {9}, image, (host-sycl)
|                 +- 11: input, "subbuffer.cpp", c++, (device-sycl, sm_50)
|              +- 12: preprocessor, {11}, c++-cpp-output, (device-sycl, sm_50)
|           +- 13: compiler, {12}, ir, (device-sycl, sm_50)
|        +- 14: linker, {13}, ir, (device-sycl, sm_50)
|     +- 15: sycl-post-link, {14}, ir, (device-sycl, sm_50)
|     |  +- 16: file-table-tform, {15}, ir, (device-sycl, sm_50)
|     |  |  +- 17: backend, {16}, assembler, (device-sycl, sm_50)
|     |  |  |- 18: assembler, {17}, object, (device-sycl, sm_50)
|     |  |- 19: linker, {17, 18}, cuda-fatbin, (device-sycl, sm_50)
|     |- 20: foreach, {16, 19}, cuda-fatbin, (device-sycl, sm_50)
|  +- 21: file-table-tform, {15, 20}, tempfiletable, (device-sycl, sm_50)
|- 22: clang-offload-wrapper, {21}, object, (device-sycl, sm_50)
|           +- 23: linker, {5}, ir, (device-sycl)
|           |  +- 24: input, "/home/aidan/llvm/build/bin/../lib/libsycl-crt.o", object
|           |- 25: clang-offload-unbundler, {24}, object
|           |  +- 26: input, "/home/aidan/llvm/build/bin/../lib/libsycl-complex.o", object
|           |- 27: clang-offload-unbundler, {26}, object
|           |  +- 28: input, "/home/aidan/llvm/build/bin/../lib/libsycl-complex-fp64.o", object
|           |- 29: clang-offload-unbundler, {28}, object
|           |  +- 30: input, "/home/aidan/llvm/build/bin/../lib/libsycl-cmath.o", object
|           |- 31: clang-offload-unbundler, {30}, object
|           |  +- 32: input, "/home/aidan/llvm/build/bin/../lib/libsycl-cmath-fp64.o", object
|           |- 33: clang-offload-unbundler, {32}, object
|           |  +- 34: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cassert.o", object
|           |- 35: clang-offload-unbundler, {34}, object
|           |  +- 36: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cstring.o", object
|           |- 37: clang-offload-unbundler, {36}, object
|           |  +- 38: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-complex.o", object
|           |- 39: clang-offload-unbundler, {38}, object
|           |  +- 40: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-complex-fp64.o", object
|           |- 41: clang-offload-unbundler, {40}, object
|           |  +- 42: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cmath.o", object
|           |- 43: clang-offload-unbundler, {42}, object
|           |  +- 44: input, "/home/aidan/llvm/build/bin/../lib/libsycl-fallback-cmath-fp64.o", object
|           |- 45: clang-offload-unbundler, {44}, object
|        +- 46: linker, {23, 25, 27, 29, 31, 33, 35, 37, 39, 41, 43, 45}, ir, (device-sycl)
|     +- 47: sycl-post-link, {46}, tempfiletable, (device-sycl)
|     |     +- 48: file-table-tform, {47}, tempfilelist, (device-sycl)
|     |  +- 49: llvm-spirv, {48}, tempfilelist, (device-sycl)
|     |- 50: backend-compiler, {49}, image, (device-sycl)
|  +- 51: file-table-tform, {47, 50}, tempfiletable, (device-sycl)
|- 52: clang-offload-wrapper, {51}, object, (device-sycl)
53: offload, "host-sycl (x86_64-unknown-linux-gnu)" {10}, "device-sycl (nvptx64-nvidia-cuda:sm_50)" {22}, "device-sycl (spir64_x86_64-unknown-unknown)" {52}, image

Section of code which causes swap

// Device compilation generates LLVM BC.
if (CurPhase == phases::Compile) {
  for (Action *&A : SYCLDeviceActions) {
    types::ID OutputType = types::TY_LLVM_BC;
    if ((SYCLDeviceOnly || Args.hasArg(options::OPT_emit_llvm)) &&
        Args.hasArg(options::OPT_S))
      OutputType = types::TY_LLVM_IR;
    if (SYCLDeviceOnly) {
      if (Args.hasFlag(options::OPT_fno_sycl_use_bitcode,
                        options::OPT_fsycl_use_bitcode, false)) {
        auto *CompileAction =
            C.MakeAction<CompileJobAction>(A, types::TY_LLVM_BC);
        A = C.MakeAction<SPIRVTranslatorJobAction>(CompileAction,
                                                    types::TY_SPIRV);
        continue;
      }
    }
    A = C.MakeAction<CompileJobAction>(A, OutputType);
    DeviceCompilerInput = A;
  }
  const auto *TC = ToolChains.front();
  const char *BoundArch = nullptr;
  if (TC->getTriple().isNVPTX() || TC->getTriple().isAMDGCN())
    BoundArch = GpuArchList.front().second.c_str();
  DA.add(*DeviceCompilerInput, *TC, BoundArch, Action::OFK_SYCL);
  // Clear the input file, it is already a dependence to a host
  // action.
  DeviceCompilerInput = nullptr;
  return SYCLDeviceOnly ? ABRT_Ignore_Host : ABRT_Success;
}

The SYCLDeviceActions are ordered based upon the toolchain.
So A is likely to be the last action, however we retrieve the first toolchain.
In effect the triple is wrong and the bound arch is correct.

A temporary fix is to pull from the back of the toolchain.

diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 490baf333772..242e0881f966 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -4082,7 +4082,7 @@ class OffloadingActionBuilder final {
           A = C.MakeAction<CompileJobAction>(A, OutputType);
           DeviceCompilerInput = A;
         }
-        const auto *TC = ToolChains.front();
+        const auto *TC = ToolChains.back();
         const char *BoundArch = nullptr;
         if (TC->getTriple().isNVPTX() || TC->getTriple().isAMDGCN())
           BoundArch = GpuArchList.front().second.c_str();

Proposed fix:

This issue with the patch will likely turn up again in a different form as GPUArchList, ToolChains, and SYCLTripleList could easily not be in the same order.
I think introducing a fix with a patch similar to #2087 would be reasonable as it would make multi-targetting much easier to handle and allow for multiple archs per triple.

[AidanBeltonS]

@Michoumichmich could you confirm if the above patch resolves the multi-targetting order error?
So I can confirm that this is the problem?

[Michoumichmich]

Hello,
Sorry for reopening that issue, but I failed on me today.
When I build with nvptx64-nvidia-cuda,spir64 or nvptx64-nvidia-cuda,spir64_x86_64 I get the runtime error:
llvm/sycl/source/detail/program_manager/program_manager.cpp:1084: void sycl::detail::ProgramManager::addImages(pi_device_binaries): Assertion 'Result.second && "Kernel sets are not disjoint"' failed. when I launch a kernel on the GPU or the CPU.
If I switch the targets, it works again.
I'm running on d970bb2a9d5032726e26dfced8c2450693536d5c

[Michoumichmich]

@bader should I open another issue ?