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