https://github.com/aaronj0 updated https://github.com/llvm/llvm-project/pull/229788
>From ea0660c8ab1b2e3e1d32f240cf483d1ee16e91a7 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Wed, 21 Oct 2026 12:49:02 +0200 Subject: [PATCH 1/2] [clang][CodeGen] Add flag to skip non-internalized bitcode files in loadLinkModules --- clang/include/clang/CodeGen/ModuleLinker.h | 9 ++++++--- clang/lib/CIR/FrontendAction/CIRGenAction.cpp | 3 ++- clang/lib/CodeGen/CodeGenAction.cpp | 6 ++++-- clang/lib/CodeGen/ModuleLinker.cpp | 5 ++++- 4 files changed, 16 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/CodeGen/ModuleLinker.h b/clang/include/clang/CodeGen/ModuleLinker.h index f24b40f7cba236..6d973a46ce3da3 100644 --- a/clang/include/clang/CodeGen/ModuleLinker.h +++ b/clang/include/clang/CodeGen/ModuleLinker.h @@ -35,10 +35,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); /// Link every module in \p LinkModules into \p M, propagating default /// function attributes and internalizing as each entry requests. Linker diff --git a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp index 483d02cb782ce7..3cdf8f22cea017 100644 --- a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp +++ b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp @@ -379,7 +379,8 @@ CIRGenAction::~CIRGenAction() { MLIRMod.release(); } bool CIRGenAction::BeginSourceFileAction(CompilerInstance &CI) { Ctx->getLibraryOptions().assign(CI.getLibraryOptions()); Ctx->setVirtualFileSystem(CI.getVirtualFileSystemPtr()); - if (clang::loadLinkModules(CI, *Ctx, LinkModules)) + if (clang::loadLinkModules(CI, *Ctx, LinkModules, + /*OnlyInternalized=*/false)) return false; return ASTFrontendAction::BeginSourceFileAction(CI); } diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 61ca1904d52685..603913b2eea327 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -372,7 +372,8 @@ CodeGenAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) { return nullptr; // Load bitcode modules to link with, if we need to. - if (clang::loadLinkModules(CI, *VMContext, LinkModules)) + if (clang::loadLinkModules(CI, *VMContext, LinkModules, + /*OnlyInternalized=*/false)) return nullptr; CoverageSourceInfo *CoverageInfo = nullptr; @@ -450,7 +451,8 @@ CodeGenAction::loadModule(MemoryBufferRef MBRef) { } // Load bitcode modules to link with, if we need to. - if (clang::loadLinkModules(CI, *VMContext, LinkModules)) + if (clang::loadLinkModules(CI, *VMContext, LinkModules, + /*OnlyInternalized=*/false)) return nullptr; // Handle textual IR and bitcode file with one single module. diff --git a/clang/lib/CodeGen/ModuleLinker.cpp b/clang/lib/CodeGen/ModuleLinker.cpp index b50b412622f439..c8fdc5ad326051 100644 --- a/clang/lib/CodeGen/ModuleLinker.cpp +++ b/clang/lib/CodeGen/ModuleLinker.cpp @@ -20,12 +20,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) >From 277a60769bed23ed210fb7905601223586c8f899 Mon Sep 17 00:00:00 2001 From: Aaron Jomy <[email protected]> Date: Wed, 21 Oct 2026 12:49:02 +0200 Subject: [PATCH 2/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 | 8 ++ clang/test/Interpreter/CUDA/libdevice-math.cu | 23 ++++++ clang/unittests/Interpreter/CMakeLists.txt | 1 + .../Interpreter/DeviceOffloadTest.cpp | 79 +++++++++++++++++++ 4 files changed, 111 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 603913b2eea327..c54cc093218c91 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -172,6 +172,14 @@ 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() && + loadLinkModules(CI, M->getContext(), LinkModules, + /*OnlyInternalized=*/true)) + return true; + return linkInModules(*M, LinkModules, CodeGenOpts, LangOpts, TargetOpts, DiagConsumer); } diff --git a/clang/test/Interpreter/CUDA/libdevice-math.cu b/clang/test/Interpreter/CUDA/libdevice-math.cu new file mode 100644 index 00000000000000..f47daf49b82293 --- /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 ad5836d8a1b7fc..f4909d06ce73d2 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 f3b7ddf837cde7..bac7fdd53408f7 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" @@ -114,4 +123,74 @@ TEST(DeviceOffloadBuilderTest, MultipleOffloadArchsAreRejected) { "supported. Pass a single --offload-arch")); } +// 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 *Fn = llvm::Function::Create( + llvm::FunctionType::get(FloatTy, {FloatTy}, /*isVarArg=*/false), + 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) { + // Builtin bitcode is linked into every PTU and plain bitcode only once. + 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.CreateDevice(OffloadType::CUDA); + ASSERT_THAT_EXPECTED(DeviceCI, llvm::Succeeded()); + // 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.CreateHost(OffloadType::CUDA); + ASSERT_THAT_EXPECTED(HostCI, llvm::Succeeded()); + auto Interp = Interpreter::createWithDevice( + std::move(*HostCI), std::move(*DeviceCI), OffloadType::CUDA); + 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 : + {"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
