[CUDA] Do not embed a fatbinary when using the new driver

[jhuber6]

Summary

Previously, when using the new driver we created a fatbinary with the
PTX and Cubin output. This was mainly done in an attempt to create some
backwards compatibility with the existing CUDA support that embeds the
fatbinary in each TU. This will most likely be more work than necessary
to actually implement. The linker wrapper cannot do anything with these
embedded PTX files because we do not know how to link them, and if we
did want to include multiple files it should go through the
clang-offload-packager instead. Also this didn't repsect the setting
that disables embedding PTX (although it wasn't used anyway).

[jhuber6]

Remove comment that is no longer true now that getInputFilename always returns a .cubin variant for object types.

[tra]

The linker wrapper cannot do anything with these embedded PTX files because we do not know how to link them,

Neither, apparently does nvlink. It does have --emip-ptx <file> option, but only if LTO is enabled, which matches the new driver behavior.

[jhuber6]

In D128441#3605809, @tra wrote:

The linker wrapper cannot do anything with these embedded PTX files because we do not know how to link them,

Neither, apparently does nvlink. It does have --emip-ptx <file> option, but only if LTO is enabled, which matches the new driver behavior.

Thanks for the review. I'm not sure exactly how CUDA does it, but for their RDC support they do somehow link PTX from multiple TU's at runtime for JIT. I'm guessing they just compile each file upon initialization and link them with nvlink. I think using LTO for JIT support is the saner option in that case.

[tra]

This change breaks clang++ --cuda-device-only compilation. Clang does not create any output in this case. Reverting the change fixes the problem.

Reproducible with:

echo '__global__ void k(){}' | bin/clang++  --offload-arch=sm_70 -x cuda -  --cuda-device-only -v  -c -o foo123.o

Compilation succeeds, but there's no foo123.o to be found.

[jhuber6]

In D128441#3702800, @tra wrote:

This change breaks clang++ --cuda-device-only compilation. Clang does not create any output in this case. Reverting the change fixes the problem.

Reproducible with:

echo '__global__ void k(){}' | bin/clang++  --offload-arch=sm_70 -x cuda -  --cuda-device-only -v  -c -o foo123.o

Compilation succeeds, but there's no foo123.o to be found.

Is it spitting it out as foo123.cubin instead?

[tra]

In D128441#3702816, @jhuber6 wrote:

Is it spitting it out as foo123.cubin instead?

That's the output name it passes to ptxas, but it's treated as a temporary file and is removed at the end, so the user gets nothing.

clang/lib/Driver/Driver.cpp

for (phases::ID Phase : PL) {
           // collapsed with any other actions so we can use it in the device.
           HostAction->setCannotBeCollapsedWithNextDependentAction();
           OffloadAction::HostDependence HDep(
               *HostAction, *C.getSingleOffloadToolChain<Action::OFK_Host>(),
               TCAndArch->second.data(), Kind);
           OffloadAction::DeviceDependences DDep;
           DDep.add(*A, *TCAndArch->first, TCAndArch->second.data(), Kind);
           A = C.MakeAction<OffloadAction>(HDep, DDep);
-        } else if (isa<AssembleJobAction>(A) && Kind == Action::OFK_Cuda) {
-          // The Cuda toolchain uses fatbinary as the linker phase to bundle the
-          // PTX and Cubin output.
-          ActionList FatbinActions;
-          for (Action *A : {A, A->getInputs()[0]}) {
-            OffloadAction::DeviceDependences DDep;
-            DDep.add(*A, *TCAndArch->first, TCAndArch->second.data(), Kind);
-            FatbinActions.emplace_back(
-                C.MakeAction<OffloadAction>(DDep, A->getType()));
-          }
-          A = C.MakeAction<LinkJobAction>(FatbinActions, types::TY_CUDA_FATBIN);
         }
         ++TCAndArch;
       }
     }
 
     auto TCAndArch = TCAndArchs.begin();
     for (Action *A : DeviceActions) {
       DDeps.add(*A, *TCAndArch->first, TCAndArch->second.data(), Kind);

clang/lib/Driver/ToolChains/Cuda.cpp

for (const auto& II : Inputs) {
     if (II.getType() == types::TY_PP_Asm &&
         !shouldIncludePTX(Args, gpu_arch_str))
       continue;
     // We need to pass an Arch of the form "sm_XX" for cubin files and
     // "compute_XX" for ptx.
     const char *Arch = (II.getType() == types::TY_PP_Asm)
                            ? CudaArchToVirtualArchString(gpu_arch)
                            : gpu_arch_str;
-    CmdArgs.push_back(Args.MakeArgString(llvm::Twine("--image=profile=") +
-                                         Arch + ",file=" + II.getFilename()));
+    CmdArgs.push_back(
+        Args.MakeArgString(llvm::Twine("--image=profile=") + Arch +
+                           ",file=" + getToolChain().getInputFilename(II)));
   }
 
   for (const auto& A : Args.getAllArgValues(options::OPT_Xcuda_fatbinary))
     CmdArgs.push_back(Args.MakeArgString(A));
 
   const char *Exec = Args.MakeArgString(TC.GetProgramPath("fatbinary"));
   C.addCommand(std::make_unique<Command>(
       JA, *this,
CudaToolChain::CudaToolChain(const Driver &D, const llvm::Triple &Triple,
   }
   // Lookup binaries into the driver directory, this is used to
   // discover the clang-offload-bundler executable.
   getProgramPaths().push_back(getDriver().Dir);
 }
 
 std::string CudaToolChain::getInputFilename(const InputInfo &Input) const {
   // Only object files are changed, for example assembly files keep their .s
-  // extensions. CUDA also continues to use .o as they don't use nvlink but
-  // fatbinary.
-  if (!(OK == Action::OFK_OpenMP && Input.getType() == types::TY_Object))
+  // extensions. 
+  if (Input.getType() != types::TY_Object)
     return ToolChain::getInputFilename(Input);
 
   // Replace extension for object files with cubin because nvlink relies on
   // these particular file names.
   SmallString<256> Filename(ToolChain::getInputFilename(Input));
   llvm::sys::path::replace_extension(Filename, "cubin");
   return std::string(Filename.str());
 }

clang/test/Driver/cuda-openmp-driver.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // RUN: %clang -### -target x86_64-linux-gnu -nocudalib -ccc-print-bindings -fgpu-rdc \
 // RUN:        --offload-new-driver --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
 // RUN: | FileCheck -check-prefix BINDINGS %s
 
-// BINDINGS: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX_SM_35:.+]]"
+//      BINDINGS: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX_SM_35:.+]]"
 // BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Linker", inputs: ["[[CUBIN_SM_35]]", "[[PTX_SM_35]]"], output: "[[FATBIN_SM_35:.+]]"
 // BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]"], output: "[[PTX_SM_70:.+]]"
 // BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Linker", inputs: ["[[CUBIN_SM_70]]", "[[PTX_SM_70:.+]]"], output: "[[FATBIN_SM_70:.+]]"
-// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Packager", inputs: ["[[FATBIN_SM_35]]", "[[FATBIN_SM_70]]"], output: "[[BINARY:.+]]"
+// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Packager", inputs: ["[[CUBIN_SM_35]]", "[[CUBIN_SM_70]]"], output: "[[BINARY:.+]]"
 // BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
 // BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[HOST_OBJ]]"], output: "a.out"
 
 // RUN: %clang -### -nocudalib --offload-new-driver %s 2>&1 | FileCheck -check-prefix RDC %s
 // RDC: error: Using '--offload-new-driver' requires '-fgpu-rdc'
 
 // RUN: %clang -### -target x86_64-linux-gnu -nocudalib -ccc-print-bindings -fgpu-rdc \
 // RUN:        --offload-new-driver --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
 // RUN: | FileCheck -check-prefix BINDINGS-HOST %s
 
 // BINDINGS-HOST: # "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[OUTPUT:.+]]"
 // BINDINGS-HOST: # "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[OUTPUT]]"], output: "a.out"
 
 // RUN: %clang -### -target x86_64-linux-gnu -nocudalib -ccc-print-bindings -fgpu-rdc \
 // RUN:        --offload-new-driver --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
 // RUN: | FileCheck -check-prefix BINDINGS-DEVICE %s
 
 // BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
 // BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "NVPTX::Linker", inputs: ["[[CUBIN]]", "[[PTX]]"], output: "{{.*}}.fatbin"
 
 // RUN: %clang -### -target x86_64-linux-gnu -nocudalib --cuda-feature=+ptx61 --offload-arch=sm_70 %s 2>&1 | FileCheck -check-prefix MANUAL-FEATURE %s
 // MANUAL-FEATURE: -cc1{{.*}}-target-feature{{.*}}+ptx61
 
 // RUN: %clang -### -target x86_64-linux-gnu -nocudalib -ccc-print-bindings --offload-link %s 2>&1 \
 // RUN: | FileCheck -check-prefix DEVICE-LINK %s
 
 // DEVICE-LINK: "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[INPUT:.+]]"], output: "a.out"

clang/test/Driver/cuda-phases.cu

 // NEW_DRIVER: 1: preprocessor, {0}, cuda-cpp-output
 // NEW_DRIVER: 2: compiler, {1}, ir
 // NEW_DRIVER: 3: input, "[[INPUT]]", cuda, (device-cuda, sm_52)
 // NEW_DRIVER: 4: preprocessor, {3}, cuda-cpp-output, (device-cuda, sm_52)
 // NEW_DRIVER: 5: compiler, {4}, ir, (device-cuda, sm_52)
 // NEW_DRIVER: 6: backend, {5}, assembler, (device-cuda, sm_52)
 // NEW_DRIVER: 7: assembler, {6}, object, (device-cuda, sm_52)
 // NEW_DRIVER: 8: offload, "device-cuda (nvptx64-nvidia-cuda:sm_52)" {7}, object
-// NEW_DRIVER: 9: offload, "device-cuda (nvptx64-nvidia-cuda:sm_52)" {6}, assembler
-// NEW_DRIVER: 10: linker, {8, 9}, cuda-fatbin, (device-cuda, sm_52)
-// NEW_DRIVER: 11: offload, "device-cuda (nvptx64-nvidia-cuda:sm_52)" {10}, cuda-fatbin
-// NEW_DRIVER: 12: input, "[[INPUT]]", cuda, (device-cuda, sm_70)
-// NEW_DRIVER: 13: preprocessor, {12}, cuda-cpp-output, (device-cuda, sm_70)
-// NEW_DRIVER: 14: compiler, {13}, ir, (device-cuda, sm_70)
-// NEW_DRIVER: 15: backend, {14}, assembler, (device-cuda, sm_70)
-// NEW_DRIVER: 16: assembler, {15}, object, (device-cuda, sm_70)
-// NEW_DRIVER: 17: offload, "device-cuda (nvptx64-nvidia-cuda:sm_70)" {16}, object
-// NEW_DRIVER: 18: offload, "device-cuda (nvptx64-nvidia-cuda:sm_70)" {15}, assembler
-// NEW_DRIVER: 19: linker, {17, 18}, cuda-fatbin, (device-cuda, sm_70)
-// NEW_DRIVER: 20: offload, "device-cuda (nvptx64-nvidia-cuda:sm_70)" {19}, cuda-fatbin
-// NEW_DRIVER: 21: clang-offload-packager, {11, 20}, image
-// NEW_DRIVER: 22: offload, " (powerpc64le-ibm-linux-gnu)" {2}, " (powerpc64le-ibm-linux-gnu)" {21}, ir
-// NEW_DRIVER: 23: backend, {22}, assembler, (host-cuda)
-// NEW_DRIVER: 24: assembler, {23}, object, (host-cuda)
-// NEW_DRIVER: 25: clang-linker-wrapper, {24}, image, (host-cuda)
+// NEW_DRIVER: 9: input, "[[INPUT]]", cuda, (device-cuda, sm_70)
+// NEW_DRIVER: 10: preprocessor, {9}, cuda-cpp-output, (device-cuda, sm_70)
+// NEW_DRIVER: 11: compiler, {10}, ir, (device-cuda, sm_70)
+// NEW_DRIVER: 12: backend, {11}, assembler, (device-cuda, sm_70)
+// NEW_DRIVER: 13: assembler, {12}, object, (device-cuda, sm_70)
+// NEW_DRIVER: 14: offload, "device-cuda (nvptx64-nvidia-cuda:sm_70)" {13}, object
+// NEW_DRIVER: 15: clang-offload-packager, {8, 14}, image
+// NEW_DRIVER: 16: offload, " (powerpc64le-ibm-linux-gnu)" {2}, " (powerpc64le-ibm-linux-gnu)" {15}, ir
+// NEW_DRIVER: 17: backend, {16}, assembler, (host-cuda)
+// NEW_DRIVER: 18: assembler, {17}, object, (host-cuda)
+// NEW_DRIVER: 19: clang-linker-wrapper, {18}, image, (host-cuda)