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

[Aaron Jomy via cfe-commits]

https://github.com/aaronj0 updated https://github.com/llvm/llvm-project/pull/226977

>From a28f2c126c57ff528b5f498c429dcaf7e453b3c2 Mon Sep 17 00:00:00 2001
From: Aaron Jomy <aaronjomyjoseph at gmail.com>
Date: Thu, 24 Sep 2026 16:02:54 +0200
Subject: [PATCH 1/2] [clang-repl] Fix PTX emission for CUDA inputs with no
 device code

GeneratePTX() treated a false return from legacy::PassManager::run() as
a failure. That value reports whether any pass changed the module, not
whether emission succeeded. Emission failures surface from
addPassesToEmitFile(), which is already checked.

The check was harmless while NVPTXAssignValidGlobalNames returned true
unconditionally. Since e8b75c172810 the pass returns Changed, so the
empty device module a host-only input produces reports no change and
every such input to `clang-repl --cuda` fails with
`error: Failed to emit PTX code`.

The host-supports-cuda lit probe is mostly host-only inputs, so
clang/test/Interpreter/CUDA reports UNSUPPORTED instead of failing.

After, host-only inputs succeed and the probe passes. The lit test
feeds host-only inputs and then launches a kernel.
---
 clang/lib/Interpreter/DeviceOffload.cpp        |  4 +---
 .../Interpreter/CUDA/empty-device-module.cu    | 18 ++++++++++++++++++
 2 files changed, 19 insertions(+), 3 deletions(-)
 create mode 100644 clang/test/Interpreter/CUDA/empty-device-module.cu

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

>From 4185559089d99e5f43350709aa60861488ca3ccf Mon Sep 17 00:00:00 2001
From: Aaron Jomy <aaronjomyjoseph at gmail.com>
Date: Thu, 1 Oct 2026 17:51:42 +0200
Subject: [PATCH 2/2] [clang-repl] Test the empty CUDA device module without a
 GPU

Add a unittest that builds a CUDA interpreter without a toolkit, parses
one host-only input and checks that the fatbin the host side embeds
holds the PTX of the empty device module. The test only requires the
NVPTX backend and does not need a CUDA toolkit or a GPU present, unlike
the lit test, which launches a kernel.

Adds EmptyDeviceModule to DeviceOffloadTest.cpp.
---
 .../Interpreter/DeviceOffloadTest.cpp         | 28 +++++++++++++++++++
 1 file changed, 28 insertions(+)

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