[clang] [Clang][CUDA][Driver] Diagnose unsupported -emit-llvm with fatbinary (PR #177577)

[via cfe-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-driver

@llvm/pr-subscribers-clang

Author: cseslowpoke (cseslowpoke)

<details>
<summary>Changes</summary>

Fix #<!-- -->172426 

---
Full diff: https://github.com/llvm/llvm-project/pull/177577.diff


3 Files Affected:

- (modified) clang/include/clang/Basic/DiagnosticDriverKinds.td (+4) 
- (modified) clang/lib/Driver/ToolChains/Cuda.cpp (+4) 
- (added) clang/test/Driver/cuda-emit-llvm-fatbinary-error.cu (+7) 


``````````diff
diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td
index db0f521b73544..8a7e70ddf118e 100644
--- a/clang/include/clang/Basic/DiagnosticDriverKinds.td
+++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td
@@ -66,6 +66,10 @@ def err_drv_no_cuda_libdevice : Error<
   "via '--cuda-path', or pass '-nocudalib' to build without linking with "
   "libdevice">;
 
+def err_drv_cuda_fatbinary_requires_elf_or_ptx : Error<
+  "cannot create CUDA fatbinary from LLVM IR/bitcode; use '--cuda-device-only' "
+  "or '--cuda-host-only', or compile device code to PTX/SASS (e.g. drop '-emit-llvm')">;
+
 def err_drv_no_rocm_device_lib : Error<
   "cannot find ROCm device library%select{| for %1| for ABI version %1"
   "%select{|, which requires ROCm %3 or higher}2}0; provide its path via "
diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp
index 6cc73ff5fc1f6..6b6e38def2ed7 100644
--- a/clang/lib/Driver/ToolChains/Cuda.cpp
+++ b/clang/lib/Driver/ToolChains/Cuda.cpp
@@ -546,6 +546,10 @@ void NVPTX::FatBinary::ConstructJob(Compilation &C, const JobAction &JA,
     CmdArgs.push_back("-g");
 
   for (const auto &II : Inputs) {
+    if (II.getType() == types::TY_LLVM_BC || II.getType() == types::TY_LLVM_IR) {
+      C.getDriver().Diag(diag::err_drv_cuda_fatbinary_requires_elf_or_ptx);
+      return;
+    }
     auto *A = II.getAction();
     assert(A->getInputs().size() == 1 &&
            "Device offload action is expected to have a single input");
diff --git a/clang/test/Driver/cuda-emit-llvm-fatbinary-error.cu b/clang/test/Driver/cuda-emit-llvm-fatbinary-error.cu
new file mode 100644
index 0000000000000..ee007e461f6e9
--- /dev/null
+++ b/clang/test/Driver/cuda-emit-llvm-fatbinary-error.cu
@@ -0,0 +1,7 @@
+// RUN: not %clangxx -### --target=x86_64-linux-gnu -S -emit-llvm --cuda-gpu-arch=sm_52 \
+// RUN:   --cuda-path=%S/Inputs/CUDA/usr/local/cuda %s 2>&1 | FileCheck %s
+//
+// CHECK: error: cannot create CUDA fatbinary from LLVM IR/bitcode
+// CHECK-NOT: "--create"
+
+__global__ void k() {}

``````````

</details>


https://github.com/llvm/llvm-project/pull/177577