https://github.com/aaronj0 created https://github.com/llvm/llvm-project/pull/229788
BackendConsumer::LinkInModules links the -mlink-builtin-bitcode modules into the first module it finalizes and clears the list. The interpreter runs the consumer once per PTU, so only the initial PTU gets libdevice. A kernel using device math in a later PTU references undefined __nv_* symbols and the PTX JIT rejects it with cudaErrorInvalidPtx. Reload the builtin bitcode for every PTU in incremental mode. The modules get internalized into whatever links them, so each PTU needs its own copy. Plain -mlink-bitcode-file inputs are still linked once. Add a lit test that calls sinf and expf from a kernel (needs a GPU) and a unittest that links a one-function bitcode library as builtin bitcode and checks that two device PTUs define the function rather than declare it. Assisted by Fable 5.1, Human in the loop >From ef87b6b34e333e4816611c7a68faa6940a5065d2 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Mon, 5 Oct 2026 14:55:40 +0200 Subject: [PATCH] [clang-repl] Link builtin bitcode into every PTU BackendConsumer::LinkInModules links the -mlink-builtin-bitcode modules into the first module it finalizes and clears the list. The interpreter runs the consumer once per PTU, so only the initial PTU gets libdevice. A kernel using device math in a later PTU references undefined __nv_* symbols and the PTX JIT rejects it with cudaErrorInvalidPtx. Reload the builtin bitcode for every PTU in incremental mode. The modules get internalized into whatever links them, so each PTU needs its own copy. Plain -mlink-bitcode-file inputs are still linked once. Add a lit test that calls sinf and expf from a kernel (needs a GPU) and a unittest that links a one-function bitcode library as builtin bitcode and checks that two device PTUs define the function rather than declare it. --- clang/lib/CodeGen/CodeGenAction.cpp | 11 ++++ clang/test/Interpreter/CUDA/libdevice-math.cu | 23 +++++++ clang/unittests/Interpreter/CMakeLists.txt | 1 + .../Interpreter/DeviceOffloadTest.cpp | 62 +++++++++++++++++++ 4 files changed, 97 insertions(+) create mode 100644 clang/test/Interpreter/CUDA/libdevice-math.cu diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 7b2d39069702a..e8667d63aa658 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -28,6 +28,7 @@ #include "clang/Frontend/MultiplexConsumer.h" #include "clang/Lex/Preprocessor.h" #include "clang/Serialization/ASTWriter.h" +#include "llvm/ADT/STLExtras.h" #include "llvm/ADT/ScopeExit.h" #include "llvm/Bitcode/BitcodeReader.h" #include "llvm/IR/DebugInfo.h" @@ -173,6 +174,16 @@ void BackendConsumer::HandleInterestingDecl(DeclGroupRef D) { // Links each entry in LinkModules into our module. Returns true on error. bool BackendConsumer::LinkInModules(llvm::Module *M) { + // Builtin bitcode is internalized into the module that links it, so each + // PTU of an incremental session needs its own copy. Plain bitcode files are + // still linked once. + if (LinkModules.empty() && LangOpts.IncrementalExtensions) { + if (loadLinkModules(CI, M->getContext(), LinkModules)) + return true; + llvm::erase_if(LinkModules, + [](const LinkModule &LM) { return !LM.Internalize; }); + } + for (auto &LM : LinkModules) { assert(LM.Module && "LinkModule does not actually have a module"); diff --git a/clang/test/Interpreter/CUDA/libdevice-math.cu b/clang/test/Interpreter/CUDA/libdevice-math.cu new file mode 100644 index 0000000000000..f47daf49b8229 --- /dev/null +++ b/clang/test/Interpreter/CUDA/libdevice-math.cu @@ -0,0 +1,23 @@ +// Tests device math from libdevice +// RUN: cat %s | clang-repl --cuda | FileCheck %s + +extern "C" int printf(const char*, ...); + +__global__ void test_kernel(float* value) { *value = sinf(0.0f) + expf(0.0f); } + +float var; +float* devptr = nullptr; +printf("cudaMalloc: %d\n", cudaMalloc((void **) &devptr, sizeof(float))); +// CHECK: cudaMalloc: 0 + +test_kernel<<<1,1>>>(devptr); +printf("CUDA Error: %d\n", cudaGetLastError()); +// CHECK-NEXT: CUDA Error: 0 + +printf("cudaMemcpy: %d\n", cudaMemcpy(&var, devptr, sizeof(float), cudaMemcpyDeviceToHost)); +// CHECK-NEXT: cudaMemcpy: 0 + +printf("Value: %.1f\n", var); +// CHECK-NEXT: Value: 1.0 + +%quit diff --git a/clang/unittests/Interpreter/CMakeLists.txt b/clang/unittests/Interpreter/CMakeLists.txt index 6d999104054c3..4f04e5f057645 100644 --- a/clang/unittests/Interpreter/CMakeLists.txt +++ b/clang/unittests/Interpreter/CMakeLists.txt @@ -11,6 +11,7 @@ set(CLANG_LIBS_TO_LINK else() set(LLVM_COMPONENTS_TO_LINK ${LLVM_TARGETS_TO_BUILD} + BitWriter Core MC OrcJIT diff --git a/clang/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp index 039c919fe1e1e..2a12ae5065e4f 100644 --- a/clang/unittests/Interpreter/DeviceOffloadTest.cpp +++ b/clang/unittests/Interpreter/DeviceOffloadTest.cpp @@ -14,13 +14,22 @@ #include "InterpreterTestFixture.h" +#include "clang/Basic/CodeGenOptions.h" #include "clang/Frontend/CompilerInstance.h" #include "clang/Interpreter/Interpreter.h" +#include "llvm/Bitcode/BitcodeWriter.h" +#include "llvm/IR/IRBuilder.h" +#include "llvm/IR/LLVMContext.h" +#include "llvm/IR/Module.h" +#include "llvm/Linker/Linker.h" #include "llvm/MC/TargetRegistry.h" #include "llvm/Support/Error.h" +#include "llvm/Support/FileSystem.h" +#include "llvm/Support/FileUtilities.h" #include "llvm/Support/TargetSelect.h" #include "llvm/Support/VirtualFileSystem.h" +#include "llvm/Support/raw_ostream.h" #include "llvm/TargetParser/Triple.h" #include "llvm/Testing/Support/Error.h" @@ -101,4 +110,57 @@ TEST_F(DeviceOffloadTest, EmptyDeviceModule) { EXPECT_TRUE((*Buf)->getBuffer().contains(".target")); } +TEST_F(DeviceOffloadTest, BuiltinBitcodeLinkedIntoEveryModule) { + // A bitcode library with one function stands in for libdevice. + llvm::SmallString<128> LibPath; + int LibFD; + ASSERT_FALSE(llvm::sys::fs::createTemporaryFile("lib", "bc", LibFD, LibPath)); + llvm::FileRemover LibRemover(LibPath); + { + llvm::LLVMContext Ctx; + llvm::Module Lib("lib", Ctx); + Lib.setTargetTriple(llvm::Triple("nvptx64-nvidia-cuda")); + Lib.setDataLayout(Lib.getTargetTriple().computeDataLayout()); + llvm::Type *FloatTy = llvm::Type::getFloatTy(Ctx); + llvm::Function *Twice = llvm::Function::Create( + llvm::FunctionType::get(FloatTy, {FloatTy}, /*isVarArg=*/false), + llvm::GlobalValue::ExternalLinkage, "twice", Lib); + llvm::IRBuilder<> Builder(llvm::BasicBlock::Create(Ctx, "", Twice)); + Builder.CreateRet(Builder.CreateFAdd(Twice->getArg(0), Twice->getArg(0))); + llvm::raw_fd_ostream OS(LibFD, /*shouldClose=*/true); + llvm::WriteBitcodeToFile(Lib, OS); + } + + IncrementalCompilerBuilder CB; + CB.SetCompilerArgs({"-nocudainc", "-nocudalib"}); + auto DeviceCI = CB.CreateCudaDevice(); + ASSERT_THAT_EXPECTED(DeviceCI, llvm::Succeeded()); + // The entry the driver adds for libdevice with -mlink-builtin-bitcode. + (*DeviceCI)->getCodeGenOpts().LinkBitcodeFiles.push_back( + {std::string(LibPath), /*PropagateAttrs=*/true, /*Internalize=*/true, + llvm::Linker::Flags::LinkOnlyNeeded}); + auto HostCI = CB.CreateCudaHost(); + ASSERT_THAT_EXPECTED(HostCI, llvm::Succeeded()); + auto Interp = + Interpreter::createWithCUDA(std::move(*HostCI), std::move(*DeviceCI)); + ASSERT_THAT_EXPECTED(Interp, llvm::Succeeded()); + + // Every device module that calls twice() must define it. + for (const char *Code : + {"extern \"C\" __attribute__((device)) float twice(float);" + "__attribute__((device)) float f(float x) { return twice(x); }", + "__attribute__((device)) float g(float x) { return twice(x); }"}) { + auto PTU = (*Interp)->Parse(Code); + ASSERT_THAT_EXPECTED(PTU, llvm::Succeeded()); + const CompilerInstance *CI = (*Interp)->getCompilerInstance(); + auto Buf = CI->getVirtualFileSystem().getBufferForFile( + CI->getCodeGenOpts().OffloadBinaryToEmbedFile, /*FileSize=*/-1, + /*RequiresNullTerminator=*/false); + ASSERT_TRUE(static_cast<bool>(Buf)) << Buf.getError().message(); + llvm::StringRef PTX = (*Buf)->getBuffer(); + EXPECT_FALSE(PTX.contains(".extern")) << PTX; + EXPECT_TRUE(PTX.contains("twice")) << PTX; + } +} + } // end anonymous namespace _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
