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 <[email protected]> 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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
