https://github.com/aaronj0 updated https://github.com/llvm/llvm-project/pull/229788
>From d18cec4ecfcd8521de19f5a85bebb1b436895293 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Mon, 5 Oct 2026 14:55:40 +0200 Subject: [PATCH 1/2] [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..23e438c2f5ae8 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 (LangOpts.IncrementalExtensions && LinkModules.empty()) { + 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 >From d5502897c1cc49aa5a22232b6b687379cf41b4e2 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Thu, 8 Oct 2026 12:14:38 +0200 Subject: [PATCH 2/2] [clang-repl] Address review comments Add an OnlyInternalized flag to loadLinkModules and pass it from LinkInModules, so plain bitcode files are not reopened for every PTU. Link a plain bitcode file in the unittest and delete it after the interpreter is created. --- clang/include/clang/CodeGen/ModuleLinker.h | 9 ++- clang/lib/CodeGen/CodeGenAction.cpp | 11 ++-- clang/lib/CodeGen/ModuleLinker.cpp | 5 +- .../Interpreter/DeviceOffloadTest.cpp | 61 ++++++++++++------- 4 files changed, 54 insertions(+), 32 deletions(-) diff --git a/clang/include/clang/CodeGen/ModuleLinker.h b/clang/include/clang/CodeGen/ModuleLinker.h index 4f94e6057f1e8..fcc413083dcdb 100644 --- a/clang/include/clang/CodeGen/ModuleLinker.h +++ b/clang/include/clang/CodeGen/ModuleLinker.h @@ -34,10 +34,13 @@ struct LinkModule { }; /// Load every bitcode file listed in CodeGenOpts.LinkBitcodeFiles into -/// \p LinkModules. Returns true on error (diagnostic already reported). -/// Appends to \p LinkModules; does not clear it. +/// \p LinkModules. With \p OnlyInternalized, skip the entries that are not +/// internalized, the plain -mlink-bitcode-file inputs. Returns true on error +/// (diagnostic already reported). Appends to \p LinkModules; does not clear +/// it. bool loadLinkModules(CompilerInstance &CI, llvm::LLVMContext &Ctx, - llvm::SmallVectorImpl<LinkModule> &LinkModules); + llvm::SmallVectorImpl<LinkModule> &LinkModules, + bool OnlyInternalized = false); namespace CodeGen { /// Adds attributes to \p F according to our \p CodeGenOpts and \p LangOpts, as diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 23e438c2f5ae8..785aee4ecb837 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -28,7 +28,6 @@ #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" @@ -177,12 +176,10 @@ 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 (LangOpts.IncrementalExtensions && LinkModules.empty()) { - if (loadLinkModules(CI, M->getContext(), LinkModules)) - return true; - llvm::erase_if(LinkModules, - [](const LinkModule &LM) { return !LM.Internalize; }); - } + if (LangOpts.IncrementalExtensions && LinkModules.empty() && + loadLinkModules(CI, M->getContext(), LinkModules, + /*OnlyInternalized=*/true)) + return true; for (auto &LM : LinkModules) { assert(LM.Module && "LinkModule does not actually have a module"); diff --git a/clang/lib/CodeGen/ModuleLinker.cpp b/clang/lib/CodeGen/ModuleLinker.cpp index 6c876276e344b..1eab6bd2e7afb 100644 --- a/clang/lib/CodeGen/ModuleLinker.cpp +++ b/clang/lib/CodeGen/ModuleLinker.cpp @@ -16,12 +16,15 @@ using namespace clang; bool clang::loadLinkModules(CompilerInstance &CI, llvm::LLVMContext &Ctx, - llvm::SmallVectorImpl<LinkModule> &LinkModules) { + llvm::SmallVectorImpl<LinkModule> &LinkModules, + bool OnlyInternalized) { if (!LinkModules.empty()) return false; for (const CodeGenOptions::BitcodeFileToLink &F : CI.getCodeGenOpts().LinkBitcodeFiles) { + if (OnlyInternalized && !F.Internalize) + continue; auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename); if (!BCBuf) { CI.getDiagnostics().Report(diag::err_cannot_open_file) diff --git a/clang/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp index 2a12ae5065e4f..6df5504029e5e 100644 --- a/clang/unittests/Interpreter/DeviceOffloadTest.cpp +++ b/clang/unittests/Interpreter/DeviceOffloadTest.cpp @@ -110,40 +110,59 @@ 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()); +// A bitcode library with one float(float) function that doubles its argument, +// or an empty one when Name is empty. +static void writeLibrary(int FD, llvm::StringRef Name) { + llvm::LLVMContext Ctx; + llvm::Module Lib("lib", Ctx); + Lib.setTargetTriple(llvm::Triple("nvptx64-nvidia-cuda")); + Lib.setDataLayout(Lib.getTargetTriple().computeDataLayout()); + if (!Name.empty()) { llvm::Type *FloatTy = llvm::Type::getFloatTy(Ctx); - llvm::Function *Twice = llvm::Function::Create( + llvm::Function *Fn = 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); + llvm::GlobalValue::ExternalLinkage, Name, Lib); + llvm::IRBuilder<> Builder(llvm::BasicBlock::Create(Ctx, "", Fn)); + Builder.CreateRet(Builder.CreateFAdd(Fn->getArg(0), Fn->getArg(0))); } + llvm::raw_fd_ostream OS(FD, /*shouldClose=*/true); + llvm::WriteBitcodeToFile(Lib, OS); +} + +TEST_F(DeviceOffloadTest, BuiltinBitcodeLinkedIntoEveryModule) { + // Two libraries stand in for libdevice and for a plain bitcode file. The + // plain one is linked whole into the initial module, which stays empty, so + // it carries nothing. + llvm::SmallString<128> BuiltinPath, PlainPath; + int BuiltinFD, PlainFD; + ASSERT_FALSE(llvm::sys::fs::createTemporaryFile("builtin", "bc", BuiltinFD, + BuiltinPath)); + llvm::FileRemover BuiltinRemover(BuiltinPath); + writeLibrary(BuiltinFD, "twice"); + ASSERT_FALSE( + llvm::sys::fs::createTemporaryFile("plain", "bc", PlainFD, PlainPath)); + llvm::FileRemover PlainRemover(PlainPath); + writeLibrary(PlainFD, /*Name=*/""); 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}); + // The entries the driver adds for -mlink-builtin-bitcode and for + // -mlink-bitcode-file. + auto &Files = (*DeviceCI)->getCodeGenOpts().LinkBitcodeFiles; + Files.push_back({std::string(BuiltinPath), /*PropagateAttrs=*/true, + /*Internalize=*/true, llvm::Linker::Flags::LinkOnlyNeeded}); + Files.push_back({std::string(PlainPath), /*PropagateAttrs=*/false, + /*Internalize=*/false, llvm::Linker::Flags::None}); 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()); + // The plain file is linked once, into the initial module, and never + // reopened. + ASSERT_FALSE(llvm::sys::fs::remove(PlainPath)); // Every device module that calls twice() must define it. for (const char *Code : _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
