[clang] f698bf3 - [clang-repl] Fix PTX emission for CUDA inputs with no device code (#226977)

[via cfe-commits]

Author: Aaron Jomy
Date: 2026-10-02T12:14:28Z
New Revision: f698bf310346a643113604245c6f1e87fae139f4

URL: https://github.com/llvm/llvm-project/commit/f698bf310346a643113604245c6f1e87fae139f4
DIFF: https://github.com/llvm/llvm-project/commit/f698bf310346a643113604245c6f1e87fae139f4.diff

LOG: [clang-repl] Fix PTX emission for CUDA inputs with no device code (#226977)

Previously, GeneratePTX() returned an error when PassManager::run()
returned false. That value reports whether any pass changed the module,
not whether emission succeeded. IIUC, addPassesToEmitFile() should be
the only failure point, which is already checked.

The check was harmless until e8b75c172810 ("[NVPTX] Add NewPM
boilerplate to NVPTXAssignValidGlobalNames"). Before it, that pass
returned true unconditionally (if the pass succeeded), so run() always
reported a change. Now a device module without a function definition,
produced in the case of host-only inputs, reports no change and fails
the check, falsely erroring out with `Failed to emit PTX code.` Due to
this, every host-only input to `clang-repl --cuda` is now rejected on
main. The host-supports-cuda lit probe consists mostly of such inputs,
so clang/test/Interpreter/CUDA reports UNSUPPORTED instead of failing.

I've also added a test with host-only inputs so a regression like this
could be caught.

Added: 
    clang/test/Interpreter/CUDA/empty-device-module.cu

Modified: 
    clang/lib/Interpreter/DeviceOffload.cpp
    clang/unittests/Interpreter/DeviceOffloadTest.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Interpreter/DeviceOffload.cpp b/clang/lib/Interpreter/DeviceOffload.cpp
index da571eae349e49..f9ef76c067c78a 100644
--- a/clang/lib/Interpreter/DeviceOffload.cpp
+++ b/clang/lib/Interpreter/DeviceOffload.cpp
@@ -68,9 +68,7 @@ llvm::Expected<llvm::StringRef> IncrementalCUDADeviceParser::GeneratePTX() {
         llvm::inconvertibleErrorCode());
   }
 
-  if (!PM.run(*PTU.TheModule))
-    return llvm::make_error<llvm::StringError>("Failed to emit PTX code.",
-                                               llvm::inconvertibleErrorCode());
+  PM.run(*PTU.TheModule);
 
   PTXCode += '\0';
   while (PTXCode.size() % 8)

diff  --git a/clang/test/Interpreter/CUDA/empty-device-module.cu b/clang/test/Interpreter/CUDA/empty-device-module.cu
new file mode 100644
index 00000000000000..fd9ba760e5a690
--- /dev/null
+++ b/clang/test/Interpreter/CUDA/empty-device-module.cu
@@ -0,0 +1,18 @@
+// Tests host-only inputs. They produce an empty device module, and emitting
+// PTX for it must not be reported as a failure just because no pass changed
+// the module.
+// RUN: cat %s | clang-repl --cuda | FileCheck %s
+
+extern "C" int printf(const char*, ...);
+
+int host_only = 42;
+printf("host_only: %d\n", host_only);
+// CHECK: host_only: 42
+
+__global__ void kernel() {}
+
+kernel<<<1,1>>>();
+printf("CUDA Error: %d\n", cudaGetLastError());
+// CHECK-NEXT: CUDA Error: 0
+
+%quit

diff  --git a/clang/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp
index 964ee6c1dc572c..8a2385e05a7ff4 100644
--- a/clang/unittests/Interpreter/DeviceOffloadTest.cpp
+++ b/clang/unittests/Interpreter/DeviceOffloadTest.cpp
@@ -20,7 +20,9 @@
 #include "llvm/MC/TargetRegistry.h"
 #include "llvm/Support/Error.h"
 #include "llvm/Support/TargetSelect.h"
+#include "llvm/Support/VirtualFileSystem.h"
 #include "llvm/TargetParser/Triple.h"
+#include "llvm/Testing/Support/Error.h"
 
 #include "gtest/gtest.h"
 
@@ -73,4 +75,30 @@ TEST_F(DeviceOffloadTest, FirstDeviceModuleVerifies) {
 #endif
 }
 
+TEST_F(DeviceOffloadTest, EmptyDeviceModule) {
+  // Without the runtime headers and libdevice no CUDA toolkit is needed.
+  IncrementalCompilerBuilder CB;
+  CB.SetCompilerArgs({"-nocudainc", "-nocudalib"});
+  auto DeviceCI = CB.CreateDevice(OffloadType::CUDA);
+  ASSERT_THAT_EXPECTED(DeviceCI, llvm::Succeeded());
+  auto HostCI = CB.CreateHost(OffloadType::CUDA);
+  ASSERT_THAT_EXPECTED(HostCI, llvm::Succeeded());
+  auto Interp = Interpreter::createWithDevice(
+      OffloadType::CUDA, std::move(*HostCI), std::move(*DeviceCI));
+  ASSERT_THAT_EXPECTED(Interp, llvm::Succeeded());
+
+  // A host-only input leaves the device module without a function. Its PTX
+  // must still be emitted and handed to the host side.
+  auto PTU = (*Interp)->Parse("int i = 0;");
+  ASSERT_THAT_EXPECTED(PTU, llvm::Succeeded());
+
+  const CompilerInstance *CI = (*Interp)->getCompilerInstance();
+  llvm::StringRef Fatbin = CI->getCodeGenOpts().OffloadBinaryToEmbedFile;
+  ASSERT_FALSE(Fatbin.empty());
+  auto Buf = CI->getVirtualFileSystem().getBufferForFile(
+      Fatbin, /*FileSize=*/-1, /*RequiresNullTerminator=*/false);
+  ASSERT_TRUE(static_cast<bool>(Buf)) << Buf.getError().message();
+  EXPECT_TRUE((*Buf)->getBuffer().contains(".target"));
+}
+
 } // end anonymous namespace