Author: Aaron Jomy Date: 2026-10-21T12:54:21Z New Revision: caa2f91ebd9755a984a0d671d4b4eeafc1288a3d
URL: https://github.com/llvm/llvm-project/commit/caa2f91ebd9755a984a0d671d4b4eeafc1288a3d DIFF: https://github.com/llvm/llvm-project/commit/caa2f91ebd9755a984a0d671d4b4eeafc1288a3d.diff LOG: [clang-repl] Link builtin bitcode into every PTU (#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 links 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. loadLinkModules now takes a flag to open only those entries, so plain -mlink-bitcode-file inputs are linked once and never reopened. Add a lit test that calls sinf and expf from a kernel (needs a GPU) and a unittest that links one bitcode library as builtin and an empty one as a plain bitcode file, deletes the plain one once the interpreter is up, and checks that two device PTUs define the builtin function rather than declare it. Added: clang/test/Interpreter/CUDA/libdevice-math.cu Modified: clang/include/clang/CodeGen/ModuleLinker.h clang/lib/CIR/FrontendAction/CIRGenAction.cpp clang/lib/CodeGen/CodeGenAction.cpp clang/lib/CodeGen/ModuleLinker.cpp clang/unittests/Interpreter/CMakeLists.txt clang/unittests/Interpreter/DeviceOffloadTest.cpp Removed: ################################################################################ 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..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); } @@ -372,7 +380,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 +459,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) 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
