https://github.com/aaronj0 updated https://github.com/llvm/llvm-project/pull/229788
>From a3ca63b60174c4921d6f8b12065ca8025476ecd1 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 2edd0988c39a77..b90a1f3f1bc707 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" @@ -172,6 +173,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; }); + } + 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 039c919fe1e1e0..2a12ae5065e4ff 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 6e1475575a685b6fc41292ea3bd33f28668eacff 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 f24b40f7cba236..e3152b6ed1fc1d 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 = false); /// 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/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index b90a1f3f1bc707..1e757dcbdfb8d3 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" @@ -176,12 +175,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; return linkInModules(*M, LinkModules, CodeGenOpts, LangOpts, TargetOpts, DiagConsumer); 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/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp index 2a12ae5065e4ff..6df5504029e5e1 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
