diff --git a/patches/llvm/clang16-2-CUDA.patch b/patches/llvm/clang16-2-CUDA.patch new file mode 100644 index 00000000..ffaed5b9 --- /dev/null +++ b/patches/llvm/clang16-2-CUDA.patch @@ -0,0 +1,969 @@ +diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h +index e68021845..43573fb1a 100644 +--- a/clang/include/clang/Interpreter/Interpreter.h ++++ b/clang/include/clang/Interpreter/Interpreter.h +@@ -42,8 +42,34 @@ class IncrementalParser; + /// Create a pre-configured \c CompilerInstance for incremental processing. + class IncrementalCompilerBuilder { + public: ++ IncrementalCompilerBuilder() {} ++ ++ void SetCompilerArgs(const std::vector &Args) { ++ UserArgs = Args; ++ } ++ ++ // General C++ ++ llvm::Expected> CreateCpp(); ++ ++ // Offload options ++ void SetOffloadArch(llvm::StringRef Arch) { OffloadArch = Arch; }; ++ ++ // CUDA specific ++ void SetCudaSDK(llvm::StringRef path) { CudaSDKPath = path; }; ++ ++ llvm::Expected> CreateCudaHost(); ++ llvm::Expected> CreateCudaDevice(); ++ ++private: + static llvm::Expected> + create(std::vector &ClangArgv); ++ ++ llvm::Expected> createCuda(bool device); ++ ++ std::vector UserArgs; ++ ++ llvm::StringRef OffloadArch; ++ llvm::StringRef CudaSDKPath; + }; + + /// Provides top-level interfaces for incremental compilation and execution. +@@ -52,6 +78,9 @@ class Interpreter { + std::unique_ptr IncrParser; + std::unique_ptr IncrExecutor; + ++ // An optional parser for CUDA offloading ++ std::unique_ptr DeviceParser; ++ + Interpreter(std::unique_ptr CI, llvm::Error &Err); + + llvm::Error CreateExecutor(); +@@ -66,6 +95,9 @@ public: + ~Interpreter(); + static llvm::Expected> + create(std::unique_ptr CI); ++ static llvm::Expected> ++ createWithCUDA(std::unique_ptr CI, ++ std::unique_ptr DCI); + const ASTContext &getASTContext() const; + ASTContext &getASTContext(); + const CompilerInstance *getCompilerInstance() const; +diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp +index bb887df3e..9cb8ae33b 100644 +--- a/clang/lib/CodeGen/CGCUDANV.cpp ++++ b/clang/lib/CodeGen/CGCUDANV.cpp +@@ -24,6 +24,7 @@ + #include "llvm/IR/DerivedTypes.h" + #include "llvm/IR/ReplaceConstant.h" + #include "llvm/Support/Format.h" ++#include "llvm/Support/VirtualFileSystem.h" + + using namespace clang; + using namespace CodeGen; +@@ -721,8 +722,9 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() { + // handle so CUDA runtime can figure out what to call on the GPU side. + std::unique_ptr CudaGpuBinary = nullptr; + if (!CudaGpuBinaryFileName.empty()) { +- llvm::ErrorOr> CudaGpuBinaryOrErr = +- llvm::MemoryBuffer::getFileOrSTDIN(CudaGpuBinaryFileName); ++ auto VFS = CGM.getFileSystem(); ++ auto CudaGpuBinaryOrErr = ++ VFS->getBufferForFile(CudaGpuBinaryFileName, -1, false); + if (std::error_code EC = CudaGpuBinaryOrErr.getError()) { + CGM.getDiags().Report(diag::err_cannot_open_file) + << CudaGpuBinaryFileName << EC.message(); +diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp +index 2b2192678..8ea66845e 100644 +--- a/clang/lib/CodeGen/CodeGenAction.cpp ++++ b/clang/lib/CodeGen/CodeGenAction.cpp +@@ -263,6 +263,7 @@ namespace clang { + // Links each entry in LinkModules into our module. Returns true on error. + bool LinkInModules() { + for (auto &LM : LinkModules) { ++ assert(LM.Module && "LinkModule does not actually have a module"); + if (LM.PropagateAttrs) + for (Function &F : *LM.Module) { + // Skip intrinsics. Keep consistent with how intrinsics are created +@@ -291,6 +292,7 @@ namespace clang { + if (Err) + return true; + } ++ LinkModules.clear(); + return false; // success + } + +diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp +index 12d602fed..978e4d404 100644 +--- a/clang/lib/CodeGen/CodeGenModule.cpp ++++ b/clang/lib/CodeGen/CodeGenModule.cpp +@@ -6228,6 +6228,10 @@ void CodeGenModule::EmitLinkageSpec(const LinkageSpecDecl *LSD) { + } + + void CodeGenModule::EmitTopLevelStmt(const TopLevelStmtDecl *D) { ++ // Device code should not be at top level. ++ if (LangOpts.CUDA && LangOpts.CUDAIsDevice) ++ return; ++ + std::unique_ptr &CurCGF = + GlobalTopLevelStmtBlockInFlight.first; + +diff --git a/clang/lib/CodeGen/ModuleBuilder.cpp b/clang/lib/CodeGen/ModuleBuilder.cpp +index e3e953c34..3594f4c66 100644 +--- a/clang/lib/CodeGen/ModuleBuilder.cpp ++++ b/clang/lib/CodeGen/ModuleBuilder.cpp +@@ -36,7 +36,7 @@ namespace { + IntrusiveRefCntPtr FS; // Only used for debug info. + const HeaderSearchOptions &HeaderSearchOpts; // Only used for debug info. + const PreprocessorOptions &PreprocessorOpts; // Only used for debug info. +- const CodeGenOptions CodeGenOpts; // Intentionally copied in. ++ const CodeGenOptions &CodeGenOpts; + + unsigned HandlingTopLevelDecls; + +diff --git a/clang/lib/Interpreter/CMakeLists.txt b/clang/lib/Interpreter/CMakeLists.txt +index 565e824bf..32f1b7c37 100644 +--- a/clang/lib/Interpreter/CMakeLists.txt ++++ b/clang/lib/Interpreter/CMakeLists.txt +@@ -1,6 +1,7 @@ + set(LLVM_LINK_COMPONENTS + core + native ++ MC + Option + OrcJit + Support +@@ -9,6 +10,7 @@ set(LLVM_LINK_COMPONENTS + ) + + add_clang_library(clangInterpreter ++ DeviceOffload.cpp + IncrementalExecutor.cpp + IncrementalParser.cpp + Interpreter.cpp +diff --git a/clang/lib/Interpreter/DeviceOffload.cpp b/clang/lib/Interpreter/DeviceOffload.cpp +new file mode 100644 +index 000000000..8e39af6ab +--- /dev/null ++++ b/clang/lib/Interpreter/DeviceOffload.cpp +@@ -0,0 +1,176 @@ ++//===---------- DeviceOffload.cpp - Device Offloading------------*- C++ -*-===// ++// ++// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. ++// See https://llvm.org/LICENSE.txt for license information. ++// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception ++// ++//===----------------------------------------------------------------------===// ++// ++// This file implements offloading to CUDA devices. ++// ++//===----------------------------------------------------------------------===// ++ ++#include "DeviceOffload.h" ++ ++#include "clang/Basic/TargetOptions.h" ++#include "clang/CodeGen/ModuleBuilder.h" ++#include "clang/Frontend/CompilerInstance.h" ++ ++#include "llvm/IR/LegacyPassManager.h" ++#include "llvm/MC/TargetRegistry.h" ++#include "llvm/Target/TargetMachine.h" ++ ++namespace clang { ++ ++IncrementalCUDADeviceParser::IncrementalCUDADeviceParser( ++ Interpreter &Interp, std::unique_ptr Instance, ++ IncrementalParser &HostParser, llvm::LLVMContext &LLVMCtx, ++ llvm::IntrusiveRefCntPtr FS, ++ llvm::Error &Err) ++ : IncrementalParser(Interp, std::move(Instance), LLVMCtx, Err), ++ HostParser(HostParser), VFS(FS) { ++ if (Err) ++ return; ++ StringRef Arch = CI->getTargetOpts().CPU; ++ if (!Arch.starts_with("sm_") || Arch.substr(3).getAsInteger(10, SMVersion)) { ++ Err = llvm::joinErrors(std::move(Err), llvm::make_error( ++ "Invalid CUDA architecture", ++ llvm::inconvertibleErrorCode())); ++ return; ++ } ++} ++ ++llvm::Expected ++IncrementalCUDADeviceParser::Parse(llvm::StringRef Input) { ++ auto PTU = IncrementalParser::Parse(Input); ++ if (!PTU) ++ return PTU.takeError(); ++ ++ auto PTX = GeneratePTX(); ++ if (!PTX) ++ return PTX.takeError(); ++ ++ auto Err = GenerateFatbinary(); ++ if (Err) ++ return std::move(Err); ++ ++ std::string FatbinFileName = ++ "/incr_module_" + std::to_string(PTUs.size()) + ".fatbin"; ++ VFS->addFile(FatbinFileName, 0, ++ llvm::MemoryBuffer::getMemBuffer( ++ llvm::StringRef(FatbinContent.data(), FatbinContent.size()), ++ "", false)); ++ ++ HostParser.getCI()->getCodeGenOpts().CudaGpuBinaryFileName = FatbinFileName; ++ ++ FatbinContent.clear(); ++ ++ return PTU; ++} ++ ++llvm::Expected IncrementalCUDADeviceParser::GeneratePTX() { ++ auto &PTU = PTUs.back(); ++ std::string Error; ++ ++ const llvm::Target *Target = llvm::TargetRegistry::lookupTarget( ++ PTU.TheModule->getTargetTriple(), Error); ++ if (!Target) ++ return llvm::make_error(std::move(Error), ++ std::error_code()); ++ llvm::TargetOptions TO = llvm::TargetOptions(); ++ llvm::TargetMachine *TargetMachine = Target->createTargetMachine( ++ PTU.TheModule->getTargetTriple(), getCI()->getTargetOpts().CPU, "", TO, ++ llvm::Reloc::Model::PIC_); ++ PTU.TheModule->setDataLayout(TargetMachine->createDataLayout()); ++ ++ PTXCode.clear(); ++ llvm::raw_svector_ostream dest(PTXCode); ++ ++ llvm::legacy::PassManager PM; ++ if (TargetMachine->addPassesToEmitFile(PM, dest, nullptr, ++ llvm::CGFT_AssemblyFile)) { ++ return llvm::make_error( ++ "NVPTX backend cannot produce PTX code.", ++ llvm::inconvertibleErrorCode()); ++ } ++ ++ if (!PM.run(*PTU.TheModule)) ++ return llvm::make_error("Failed to emit PTX code.", ++ llvm::inconvertibleErrorCode()); ++ ++ PTXCode += '\0'; ++ while (PTXCode.size() % 8) ++ PTXCode += '\0'; ++ return PTXCode.str(); ++} ++ ++llvm::Error IncrementalCUDADeviceParser::GenerateFatbinary() { ++ enum FatBinFlags { ++ AddressSize64 = 0x01, ++ HasDebugInfo = 0x02, ++ ProducerCuda = 0x04, ++ HostLinux = 0x10, ++ HostMac = 0x20, ++ HostWindows = 0x40 ++ }; ++ ++ struct FatBinInnerHeader { ++ uint16_t Kind; // 0x00 ++ uint16_t unknown02; // 0x02 ++ uint32_t HeaderSize; // 0x04 ++ uint32_t DataSize; // 0x08 ++ uint32_t unknown0c; // 0x0c ++ uint32_t CompressedSize; // 0x10 ++ uint32_t SubHeaderSize; // 0x14 ++ uint16_t VersionMinor; // 0x18 ++ uint16_t VersionMajor; // 0x1a ++ uint32_t CudaArch; // 0x1c ++ uint32_t unknown20; // 0x20 ++ uint32_t unknown24; // 0x24 ++ uint32_t Flags; // 0x28 ++ uint32_t unknown2c; // 0x2c ++ uint32_t unknown30; // 0x30 ++ uint32_t unknown34; // 0x34 ++ uint32_t UncompressedSize; // 0x38 ++ uint32_t unknown3c; // 0x3c ++ uint32_t unknown40; // 0x40 ++ uint32_t unknown44; // 0x44 ++ FatBinInnerHeader(uint32_t DataSize, uint32_t CudaArch, uint32_t Flags) ++ : Kind(1 /*PTX*/), unknown02(0x0101), HeaderSize(sizeof(*this)), ++ DataSize(DataSize), unknown0c(0), CompressedSize(0), ++ SubHeaderSize(HeaderSize - 8), VersionMinor(2), VersionMajor(4), ++ CudaArch(CudaArch), unknown20(0), unknown24(0), Flags(Flags), ++ unknown2c(0), unknown30(0), unknown34(0), UncompressedSize(0), ++ unknown3c(0), unknown40(0), unknown44(0) {} ++ }; ++ ++ struct FatBinHeader { ++ uint32_t Magic; // 0x00 ++ uint16_t Version; // 0x04 ++ uint16_t HeaderSize; // 0x06 ++ uint32_t DataSize; // 0x08 ++ uint32_t unknown0c; // 0x0c ++ public: ++ FatBinHeader(uint32_t DataSize) ++ : Magic(0xba55ed50), Version(1), HeaderSize(sizeof(*this)), ++ DataSize(DataSize), unknown0c(0) {} ++ }; ++ ++ FatBinHeader OuterHeader(sizeof(FatBinInnerHeader) + PTXCode.size()); ++ FatbinContent.append((char *)&OuterHeader, ++ ((char *)&OuterHeader) + OuterHeader.HeaderSize); ++ ++ FatBinInnerHeader InnerHeader(PTXCode.size(), SMVersion, ++ FatBinFlags::AddressSize64 | ++ FatBinFlags::HostLinux); ++ FatbinContent.append((char *)&InnerHeader, ++ ((char *)&InnerHeader) + InnerHeader.HeaderSize); ++ ++ FatbinContent.append(PTXCode.begin(), PTXCode.end()); ++ ++ return llvm::Error::success(); ++} ++ ++IncrementalCUDADeviceParser::~IncrementalCUDADeviceParser() {} ++ ++} // namespace clang +diff --git a/clang/lib/Interpreter/DeviceOffload.h b/clang/lib/Interpreter/DeviceOffload.h +new file mode 100644 +index 000000000..ce4f218c9 +--- /dev/null ++++ b/clang/lib/Interpreter/DeviceOffload.h +@@ -0,0 +1,51 @@ ++//===----------- DeviceOffload.h - Device Offloading ------------*- C++ -*-===// ++// ++// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. ++// See https://llvm.org/LICENSE.txt for license information. ++// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception ++// ++//===----------------------------------------------------------------------===// ++// ++// This file implements classes required for offloading to CUDA devices. ++// ++//===----------------------------------------------------------------------===// ++ ++#ifndef LLVM_CLANG_LIB_INTERPRETER_DEVICE_OFFLOAD_H ++#define LLVM_CLANG_LIB_INTERPRETER_DEVICE_OFFLOAD_H ++ ++#include "IncrementalParser.h" ++#include "llvm/Support/FileSystem.h" ++#include "llvm/Support/VirtualFileSystem.h" ++ ++namespace clang { ++ ++class IncrementalCUDADeviceParser : public IncrementalParser { ++public: ++ IncrementalCUDADeviceParser( ++ Interpreter &Interp, std::unique_ptr Instance, ++ IncrementalParser &HostParser, llvm::LLVMContext &LLVMCtx, ++ llvm::IntrusiveRefCntPtr VFS, ++ llvm::Error &Err); ++ ++ llvm::Expected ++ Parse(llvm::StringRef Input) override; ++ ++ // Generate PTX for the last PTU ++ llvm::Expected GeneratePTX(); ++ ++ // Generate fatbinary contents in memory ++ llvm::Error GenerateFatbinary(); ++ ++ ~IncrementalCUDADeviceParser(); ++ ++protected: ++ IncrementalParser &HostParser; ++ int SMVersion; ++ llvm::SmallString<1024> PTXCode; ++ llvm::SmallVector FatbinContent; ++ llvm::IntrusiveRefCntPtr VFS; ++}; ++ ++} // namespace clang ++ ++#endif // LLVM_CLANG_LIB_INTERPRETER_DEVICE_OFFLOAD_H +diff --git a/clang/lib/Interpreter/IncrementalParser.cpp b/clang/lib/Interpreter/IncrementalParser.cpp +index f892eeb81..a3ae1aa8a 100644 +--- a/clang/lib/Interpreter/IncrementalParser.cpp ++++ b/clang/lib/Interpreter/IncrementalParser.cpp +@@ -194,6 +194,15 @@ public: + } + }; + ++CodeGenerator *IncrementalParser::getCodeGen() const { ++ FrontendAction *WrappedAct = Act->getWrapped(); ++ if (!WrappedAct->hasIRSupport()) ++ return nullptr; ++ return static_cast(WrappedAct)->getCodeGenerator(); ++} ++ ++IncrementalParser::IncrementalParser() {} ++ + IncrementalParser::IncrementalParser(Interpreter &Interp, + std::unique_ptr Instance, + llvm::LLVMContext &LLVMCtx, +@@ -211,6 +220,21 @@ IncrementalParser::IncrementalParser(Interpreter &Interp, + P.reset( + new Parser(CI->getPreprocessor(), CI->getSema(), /*SkipBodies=*/false)); + P->Initialize(); ++ ++ // An initial PTU is needed as CUDA includes some headers automatically ++ auto PTU = ParseOrWrapTopLevelDecl(); ++ if (auto E = PTU.takeError()) { ++ consumeError(std::move(E)); // FIXME ++ return; // PTU.takeError(); ++ } ++ ++ if (CodeGenerator *CG = getCodeGen()) { ++ std::unique_ptr M(CG->ReleaseModule()); ++ CG->StartModule("incr_module_" + std::to_string(PTUs.size()), ++ M->getContext()); ++ PTU->TheModule = std::move(M); ++ assert(PTU->TheModule && "Failed to create initial PTU"); ++ } + } + + IncrementalParser::~IncrementalParser() { +@@ -281,14 +305,6 @@ IncrementalParser::ParseOrWrapTopLevelDecl() { + return LastPTU; + } + +-static CodeGenerator *getCodeGen(FrontendAction *Act) { +- IncrementalAction *IncrAct = static_cast(Act); +- FrontendAction *WrappedAct = IncrAct->getWrapped(); +- if (!WrappedAct->hasIRSupport()) +- return nullptr; +- return static_cast(WrappedAct)->getCodeGenerator(); +-} +- + llvm::Expected + IncrementalParser::Parse(llvm::StringRef input) { + Preprocessor &PP = CI->getPreprocessor(); +@@ -351,7 +367,7 @@ IncrementalParser::Parse(llvm::StringRef input) { + + std::unique_ptr IncrementalParser::GenModule() { + static unsigned ID = 0; +- if (CodeGenerator *CG = getCodeGen(Act.get())) { ++ if (CodeGenerator *CG = getCodeGen()) { + std::unique_ptr M(CG->ReleaseModule()); + CG->StartModule("incr_module_" + std::to_string(ID++), M->getContext()); + return M; +@@ -378,7 +394,7 @@ void IncrementalParser::CleanUpPTU(PartialTranslationUnit &PTU) { + } + + llvm::StringRef IncrementalParser::GetMangledName(GlobalDecl GD) const { +- CodeGenerator *CG = getCodeGen(Act.get()); ++ CodeGenerator *CG = getCodeGen(); + assert(CG); + return CG->GetMangledName(GD); + } +diff --git a/clang/lib/Interpreter/IncrementalParser.h b/clang/lib/Interpreter/IncrementalParser.h +index 99e37588d..def5750d1 100644 +--- a/clang/lib/Interpreter/IncrementalParser.h ++++ b/clang/lib/Interpreter/IncrementalParser.h +@@ -28,6 +28,7 @@ class LLVMContext; + + namespace clang { + class ASTConsumer; ++class CodeGenerator; + class CompilerInstance; + class IncrementalAction; + class Interpreter; +@@ -36,6 +37,7 @@ class Parser; + /// changes between the subsequent incremental input. + /// + class IncrementalParser { ++protected: + /// Long-lived, incremental parsing action. + std::unique_ptr Act; + +@@ -55,18 +57,21 @@ class IncrementalParser { + /// of code. + std::list PTUs; + ++ IncrementalParser(); ++ + public: + IncrementalParser(Interpreter &Interp, + std::unique_ptr Instance, + llvm::LLVMContext &LLVMCtx, llvm::Error &Err); +- ~IncrementalParser(); ++ virtual ~IncrementalParser(); + +- const CompilerInstance *getCI() const { return CI.get(); } ++ CompilerInstance *getCI() { return CI.get(); } ++ CodeGenerator *getCodeGen() const; + + /// Parses incremental input by creating an in-memory file. + ///\returns a \c PartialTranslationUnit which holds information about the + /// \c TranslationUnitDecl and \c llvm::Module corresponding to the input. +- llvm::Expected Parse(llvm::StringRef Input); ++ virtual llvm::Expected Parse(llvm::StringRef Input); + + /// Uses the CodeGenModule mangled name cache and avoids recomputing. + ///\returns the mangled name of a \c GD. +diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp +index 4391bd008..74d428a86 100644 +--- a/clang/lib/Interpreter/Interpreter.cpp ++++ b/clang/lib/Interpreter/Interpreter.cpp +@@ -13,6 +13,7 @@ + + #include "clang/Interpreter/Interpreter.h" + ++#include "DeviceOffload.h" + #include "IncrementalExecutor.h" + #include "IncrementalParser.h" + +@@ -22,6 +23,7 @@ + #include "clang/AST/TypeVisitor.h" + #include "clang/Basic/DiagnosticSema.h" + #include "clang/Basic/TargetInfo.h" ++#include "clang/CodeGen/CodeGenAction.h" + #include "clang/CodeGen/ModuleBuilder.h" + #include "clang/CodeGen/ObjectFilePCHContainerOperations.h" + #include "clang/Driver/Compilation.h" +@@ -146,7 +148,6 @@ IncrementalCompilerBuilder::create(std::vector &ClangArgv) { + // action and use other actions in incremental mode. + // FIXME: Print proper driver diagnostics if the driver flags are wrong. + // We do C++ by default; append right after argv[0] if no "-x" given +- ClangArgv.insert(ClangArgv.end(), "-xc++"); + ClangArgv.insert(ClangArgv.end(), "-Xclang"); + ClangArgv.insert(ClangArgv.end(), "-fincremental-extensions"); + ClangArgv.insert(ClangArgv.end(), "-c"); +@@ -179,6 +180,54 @@ IncrementalCompilerBuilder::create(std::vector &ClangArgv) { + return CreateCI(**ErrOrCC1Args); + } + ++llvm::Expected> ++IncrementalCompilerBuilder::CreateCpp() { ++ std::vector Argv; ++ Argv.reserve(5 + 1 + UserArgs.size()); ++ Argv.push_back("-xc++"); ++ Argv.insert(Argv.end(), UserArgs.begin(), UserArgs.end()); ++ ++ return IncrementalCompilerBuilder::create(Argv); ++} ++ ++llvm::Expected> ++IncrementalCompilerBuilder::createCuda(bool device) { ++ std::vector Argv; ++ Argv.reserve(5 + 4 + UserArgs.size()); ++ ++ Argv.push_back("-xcuda"); ++ if (device) ++ Argv.push_back("--cuda-device-only"); ++ else ++ Argv.push_back("--cuda-host-only"); ++ ++ std::string SDKPathArg = "--cuda-path="; ++ if (!CudaSDKPath.empty()) { ++ SDKPathArg += CudaSDKPath; ++ Argv.push_back(SDKPathArg.c_str()); ++ } ++ ++ std::string ArchArg = "--offload-arch="; ++ if (!OffloadArch.empty()) { ++ ArchArg += OffloadArch; ++ Argv.push_back(ArchArg.c_str()); ++ } ++ ++ Argv.insert(Argv.end(), UserArgs.begin(), UserArgs.end()); ++ ++ return IncrementalCompilerBuilder::create(Argv); ++} ++ ++llvm::Expected> ++IncrementalCompilerBuilder::CreateCudaDevice() { ++ return IncrementalCompilerBuilder::createCuda(true); ++} ++ ++llvm::Expected> ++IncrementalCompilerBuilder::CreateCudaHost() { ++ return IncrementalCompilerBuilder::createCuda(false); ++} ++ + Interpreter::Interpreter(std::unique_ptr CI, + llvm::Error &Err) { + llvm::ErrorAsOutParameter EAO(&Err); +@@ -239,6 +288,34 @@ Interpreter::create(std::unique_ptr CI) { + return std::move(Interp); + } + ++llvm::Expected> ++Interpreter::createWithCUDA(std::unique_ptr CI, ++ std::unique_ptr DCI) { ++ // avoid writing fat binary to disk using an in-memory virtual file system ++ llvm::IntrusiveRefCntPtr IMVFS = ++ std::make_unique(); ++ llvm::IntrusiveRefCntPtr OverlayVFS = ++ std::make_unique( ++ llvm::vfs::getRealFileSystem()); ++ OverlayVFS->pushOverlay(IMVFS); ++ CI->createFileManager(OverlayVFS); ++ ++ auto Interp = Interpreter::create(std::move(CI)); ++ if (auto E = Interp.takeError()) ++ return std::move(E); ++ ++ llvm::Error Err = llvm::Error::success(); ++ auto DeviceParser = std::make_unique( ++ **Interp, std::move(DCI), *(*Interp)->IncrParser.get(), ++ *(*Interp)->TSCtx->getContext(), IMVFS, Err); ++ if (Err) ++ return std::move(Err); ++ ++ (*Interp)->DeviceParser = std::move(DeviceParser); ++ ++ return Interp; ++} ++ + const CompilerInstance *Interpreter::getCompilerInstance() const { + return IncrParser->getCI(); + } +@@ -268,6 +345,14 @@ size_t Interpreter::getEffectivePTUSize() const { + + llvm::Expected + Interpreter::Parse(llvm::StringRef Code) { ++ // If we have a device parser, parse it first. ++ // The generated code will be included in the host compilation ++ if (DeviceParser) { ++ auto DevicePTU = DeviceParser->Parse(Code); ++ if (auto E = DevicePTU.takeError()) ++ return std::move(E); ++ } ++ + // Tell the interpreter sliently ignore unused expressions since value + // printing could cause it. + getCompilerInstance()->getDiagnostics().setSeverity( +diff --git a/clang/test/Interpreter/CUDA/device-function-template.cu b/clang/test/Interpreter/CUDA/device-function-template.cu +new file mode 100644 +index 000000000..f0077a2c5 +--- /dev/null ++++ b/clang/test/Interpreter/CUDA/device-function-template.cu +@@ -0,0 +1,24 @@ ++// Tests device function templates ++// RUN: cat %s | clang-repl --cuda | FileCheck %s ++ ++extern "C" int printf(const char*, ...); ++ ++template __device__ inline T sum(T a, T b) { return a + b; } ++__global__ void test_kernel(int* value) { *value = sum(40, 2); } ++ ++int var; ++int* devptr = nullptr; ++printf("cudaMalloc: %d\n", cudaMalloc((void **) &devptr, sizeof(int))); ++// 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(int), cudaMemcpyDeviceToHost)); ++// CHECK-NEXT: cudaMemcpy: 0 ++ ++printf("Value: %d\n", var); ++// CHECK-NEXT: Value: 42 ++ ++%quit +diff --git a/clang/test/Interpreter/CUDA/device-function.cu b/clang/test/Interpreter/CUDA/device-function.cu +new file mode 100644 +index 000000000..396f8f0f9 +--- /dev/null ++++ b/clang/test/Interpreter/CUDA/device-function.cu +@@ -0,0 +1,24 @@ ++// Tests __device__ function calls ++// RUN: cat %s | clang-repl --cuda | FileCheck %s ++ ++extern "C" int printf(const char*, ...); ++ ++__device__ inline void test_device(int* value) { *value = 42; } ++__global__ void test_kernel(int* value) { test_device(value); } ++ ++int var; ++int* devptr = nullptr; ++printf("cudaMalloc: %d\n", cudaMalloc((void **) &devptr, sizeof(int))); ++// 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(int), cudaMemcpyDeviceToHost)); ++// CHECK-NEXT: cudaMemcpy: 0 ++ ++printf("Value: %d\n", var); ++// CHECK-NEXT: Value: 42 ++ ++%quit +diff --git a/clang/test/Interpreter/CUDA/host-and-device.cu b/clang/test/Interpreter/CUDA/host-and-device.cu +new file mode 100644 +index 000000000..8e44e3403 +--- /dev/null ++++ b/clang/test/Interpreter/CUDA/host-and-device.cu +@@ -0,0 +1,27 @@ ++// Checks that a function is available in both __host__ and __device__ ++// RUN: cat %s | clang-repl --cuda | FileCheck %s ++ ++extern "C" int printf(const char*, ...); ++ ++__host__ __device__ inline int sum(int a, int b){ return a + b; } ++__global__ void kernel(int * output){ *output = sum(40,2); } ++ ++printf("Host sum: %d\n", sum(41,1)); ++// CHECK: Host sum: 42 ++ ++int var = 0; ++int * deviceVar; ++printf("cudaMalloc: %d\n", cudaMalloc((void **) &deviceVar, sizeof(int))); ++// CHECK-NEXT: cudaMalloc: 0 ++ ++kernel<<<1,1>>>(deviceVar); ++printf("CUDA Error: %d\n", cudaGetLastError()); ++// CHECK-NEXT: CUDA Error: 0 ++ ++printf("cudaMemcpy: %d\n", cudaMemcpy(&var, deviceVar, sizeof(int), cudaMemcpyDeviceToHost)); ++// CHECK-NEXT: cudaMemcpy: 0 ++ ++printf("var: %d\n", var); ++// CHECK-NEXT: var: 42 ++ ++%quit +diff --git a/clang/test/Interpreter/CUDA/lit.local.cfg b/clang/test/Interpreter/CUDA/lit.local.cfg +new file mode 100644 +index 000000000..999157246 +--- /dev/null ++++ b/clang/test/Interpreter/CUDA/lit.local.cfg +@@ -0,0 +1,2 @@ ++if 'host-supports-cuda' not in config.available_features: ++ config.unsupported = True +diff --git a/clang/test/Interpreter/CUDA/memory.cu b/clang/test/Interpreter/CUDA/memory.cu +new file mode 100644 +index 000000000..852cc04f6 +--- /dev/null ++++ b/clang/test/Interpreter/CUDA/memory.cu +@@ -0,0 +1,23 @@ ++// Tests cudaMemcpy and writes from kernel ++// RUN: cat %s | clang-repl --cuda | FileCheck %s ++ ++extern "C" int printf(const char*, ...); ++ ++__global__ void test_func(int* value) { *value = 42; } ++ ++int var; ++int* devptr = nullptr; ++printf("cudaMalloc: %d\n", cudaMalloc((void **) &devptr, sizeof(int))); ++// CHECK: cudaMalloc: 0 ++ ++test_func<<<1,1>>>(devptr); ++printf("CUDA Error: %d\n", cudaGetLastError()); ++// CHECK-NEXT: CUDA Error: 0 ++ ++printf("cudaMemcpy: %d\n", cudaMemcpy(&var, devptr, sizeof(int), cudaMemcpyDeviceToHost)); ++// CHECK-NEXT: cudaMemcpy: 0 ++ ++printf("Value: %d\n", var); ++// CHECK-NEXT: Value: 42 ++ ++%quit +diff --git a/clang/test/Interpreter/CUDA/sanity.cu b/clang/test/Interpreter/CUDA/sanity.cu +new file mode 100644 +index 000000000..ef9d68df4 +--- /dev/null ++++ b/clang/test/Interpreter/CUDA/sanity.cu +@@ -0,0 +1,11 @@ ++// RUN: cat %s | clang-repl --cuda | FileCheck %s ++ ++extern "C" int printf(const char*, ...); ++ ++__global__ void test_func() {} ++ ++test_func<<<1,1>>>(); ++printf("CUDA Error: %d", cudaGetLastError()); ++// CHECK: CUDA Error: 0 ++ ++%quit +diff --git a/clang/test/lit.cfg.py b/clang/test/lit.cfg.py +index cc55c3c44..fe237fadc 100644 +--- a/clang/test/lit.cfg.py ++++ b/clang/test/lit.cfg.py +@@ -86,9 +86,39 @@ def have_host_jit_feature_support(feature_name): + + return 'true' in clang_repl_out + ++def have_host_clang_repl_cuda(): ++ clang_repl_exe = lit.util.which('clang-repl', config.clang_tools_dir) ++ ++ if not clang_repl_exe: ++ return False ++ ++ testcode = b'\n'.join([ ++ b"__global__ void test_func() {}", ++ b"test_func<<<1,1>>>();", ++ b"extern \"C\" int puts(const char *s);", ++ b"puts(cudaGetLastError() ? \"failure\" : \"success\");", ++ b"%quit" ++ ]) ++ try: ++ clang_repl_cmd = subprocess.run([clang_repl_exe, '--cuda'], ++ stdout=subprocess.PIPE, ++ stderr=subprocess.PIPE, ++ input=testcode) ++ except OSError: ++ return False ++ ++ if clang_repl_cmd.returncode == 0: ++ if clang_repl_cmd.stdout.find(b"success") != -1: ++ return True ++ ++ return False ++ + if have_host_jit_feature_support('jit'): + config.available_features.add('host-supports-jit') + ++ if have_host_clang_repl_cuda(): ++ config.available_features.add('host-supports-cuda') ++ + if config.clang_staticanalyzer: + config.available_features.add('staticanalyzer') + tools.append('clang-check') +diff --git a/clang/tools/clang-repl/ClangRepl.cpp b/clang/tools/clang-repl/ClangRepl.cpp +index 33faf3fab..19733e193 100644 +--- a/clang/tools/clang-repl/ClangRepl.cpp ++++ b/clang/tools/clang-repl/ClangRepl.cpp +@@ -20,9 +20,13 @@ + #include "llvm/Support/CommandLine.h" + #include "llvm/Support/ManagedStatic.h" // llvm_shutdown + #include "llvm/Support/Signals.h" +-#include "llvm/Support/TargetSelect.h" // llvm::Initialize* ++#include "llvm/Support/TargetSelect.h" + #include + ++static llvm::cl::opt CudaEnabled("cuda", llvm::cl::Hidden); ++static llvm::cl::opt CudaPath("cuda-path", llvm::cl::Hidden); ++static llvm::cl::opt OffloadArch("offload-arch", llvm::cl::Hidden); ++ + static llvm::cl::list + ClangArgs("Xcc", + llvm::cl::desc("Argument to pass to the CompilerInvocation"), +@@ -76,8 +80,11 @@ int main(int argc, const char **argv) { + std::vector ClangArgv(ClangArgs.size()); + std::transform(ClangArgs.begin(), ClangArgs.end(), ClangArgv.begin(), + [](const std::string &s) -> const char * { return s.data(); }); +- llvm::InitializeNativeTarget(); +- llvm::InitializeNativeTargetAsmPrinter(); ++ // Initialize all targets (required for device offloading) ++ llvm::InitializeAllTargetInfos(); ++ llvm::InitializeAllTargets(); ++ llvm::InitializeAllTargetMCs(); ++ llvm::InitializeAllAsmPrinters(); + + if (OptHostSupportsJit) { + auto J = llvm::orc::LLJITBuilder().create(); +@@ -90,9 +97,30 @@ int main(int argc, const char **argv) { + return 0; + } + ++ clang::IncrementalCompilerBuilder CB; ++ CB.SetCompilerArgs(ClangArgv); ++ ++ std::unique_ptr DeviceCI; ++ if (CudaEnabled) { ++ if (!CudaPath.empty()) ++ CB.SetCudaSDK(CudaPath); ++ ++ if (OffloadArch.empty()) { ++ OffloadArch = "sm_35"; ++ } ++ CB.SetOffloadArch(OffloadArch); ++ ++ DeviceCI = ExitOnErr(CB.CreateCudaDevice()); ++ } ++ + // FIXME: Investigate if we could use runToolOnCodeWithArgs from tooling. It + // can replace the boilerplate code for creation of the compiler instance. +- auto CI = ExitOnErr(clang::IncrementalCompilerBuilder::create(ClangArgv)); ++ std::unique_ptr CI; ++ if (CudaEnabled) { ++ CI = ExitOnErr(CB.CreateCudaHost()); ++ } else { ++ CI = ExitOnErr(CB.CreateCpp()); ++ } + + // Set an error handler, so that any LLVM backend diagnostics go through our + // error handler. +@@ -101,8 +129,23 @@ int main(int argc, const char **argv) { + + // Load any requested plugins. + CI->LoadRequestedPlugins(); ++ if (CudaEnabled) ++ DeviceCI->LoadRequestedPlugins(); ++ ++ std::unique_ptr Interp; ++ if (CudaEnabled) { ++ Interp = ExitOnErr( ++ clang::Interpreter::createWithCUDA(std::move(CI), std::move(DeviceCI))); ++ ++ if (CudaPath.empty()) { ++ ExitOnErr(Interp->LoadDynamicLibrary("libcudart.so")); ++ } else { ++ auto CudaRuntimeLibPath = CudaPath + "/lib/libcudart.so"; ++ ExitOnErr(Interp->LoadDynamicLibrary(CudaRuntimeLibPath.c_str())); ++ } ++ } else ++ Interp = ExitOnErr(clang::Interpreter::create(std::move(CI))); + +- auto Interp = ExitOnErr(clang::Interpreter::create(std::move(CI))); + for (const std::string &input : OptInputs) { + if (auto Err = Interp->ParseAndExecute(input)) + llvm::logAllUnhandledErrors(std::move(Err), llvm::errs(), "error: "); +diff --git a/clang/unittests/Interpreter/ExceptionTests/InterpreterExceptionTest.cpp b/clang/unittests/Interpreter/ExceptionTests/InterpreterExceptionTest.cpp +index 6d0433a98..63bb69038 100644 +--- a/clang/unittests/Interpreter/ExceptionTests/InterpreterExceptionTest.cpp ++++ b/clang/unittests/Interpreter/ExceptionTests/InterpreterExceptionTest.cpp +@@ -38,7 +38,9 @@ createInterpreter(const Args &ExtraArgs = {}, + DiagnosticConsumer *Client = nullptr) { + Args ClangArgs = {"-Xclang", "-emit-llvm-only"}; + ClangArgs.insert(ClangArgs.end(), ExtraArgs.begin(), ExtraArgs.end()); +- auto CI = cantFail(clang::IncrementalCompilerBuilder::create(ClangArgs)); ++ auto CB = clang::IncrementalCompilerBuilder(); ++ CB.SetCompilerArgs(ClangArgs); ++ auto CI = cantFail(CB.CreateCpp()); + if (Client) + CI->getDiagnostics().setClient(Client, /*ShouldOwnClient=*/false); + return cantFail(clang::Interpreter::create(std::move(CI))); +diff --git a/clang/unittests/Interpreter/IncrementalProcessingTest.cpp b/clang/unittests/Interpreter/IncrementalProcessingTest.cpp +index b7ad468e1..6d477c9ab 100644 +--- a/clang/unittests/Interpreter/IncrementalProcessingTest.cpp ++++ b/clang/unittests/Interpreter/IncrementalProcessingTest.cpp +@@ -52,7 +52,9 @@ const Function *getGlobalInit(llvm::Module *M) { + + TEST(IncrementalProcessing, EmitCXXGlobalInitFunc) { + std::vector ClangArgv = {"-Xclang", "-emit-llvm-only"}; +- auto CI = llvm::cantFail(IncrementalCompilerBuilder::create(ClangArgv)); ++ auto CB = clang::IncrementalCompilerBuilder(); ++ CB.SetCompilerArgs(ClangArgv); ++ auto CI = cantFail(CB.CreateCpp()); + auto Interp = llvm::cantFail(Interpreter::create(std::move(CI))); + + std::array PTUs; +diff --git a/clang/unittests/Interpreter/InterpreterTest.cpp b/clang/unittests/Interpreter/InterpreterTest.cpp +index 330fd18ab..338003cd9 100644 +--- a/clang/unittests/Interpreter/InterpreterTest.cpp ++++ b/clang/unittests/Interpreter/InterpreterTest.cpp +@@ -46,7 +46,9 @@ createInterpreter(const Args &ExtraArgs = {}, + DiagnosticConsumer *Client = nullptr) { + Args ClangArgs = {"-Xclang", "-emit-llvm-only"}; + ClangArgs.insert(ClangArgs.end(), ExtraArgs.begin(), ExtraArgs.end()); +- auto CI = cantFail(clang::IncrementalCompilerBuilder::create(ClangArgs)); ++ auto CB = clang::IncrementalCompilerBuilder(); ++ CB.SetCompilerArgs(ClangArgs); ++ auto CI = cantFail(CB.CreateCpp()); + if (Client) + CI->getDiagnostics().setClient(Client, /*ShouldOwnClient=*/false); + return cantFail(clang::Interpreter::create(std::move(CI))); diff --git a/patches/llvm/clang16-D141215-Value.patch b/patches/llvm/clang16-D141215-Value.patch deleted file mode 100644 index 76762368..00000000 --- a/patches/llvm/clang16-D141215-Value.patch +++ /dev/null @@ -1,1838 +0,0 @@ -diff --git a/clang/include/clang/AST/Decl.h b/clang/include/clang/AST/Decl.h -index 863f6ac57f2a..2e343405cfbe 100644 ---- a/clang/include/clang/AST/Decl.h -+++ b/clang/include/clang/AST/Decl.h -@@ -4308,6 +4308,7 @@ class TopLevelStmtDecl : public Decl { - friend class ASTDeclWriter; - - Stmt *Statement = nullptr; -+ bool IsSemiMissing = false; - - TopLevelStmtDecl(DeclContext *DC, SourceLocation L, Stmt *S) - : Decl(TopLevelStmt, DC, L), Statement(S) {} -@@ -4321,6 +4322,12 @@ public: - SourceRange getSourceRange() const override LLVM_READONLY; - Stmt *getStmt() { return Statement; } - const Stmt *getStmt() const { return Statement; } -+ void setStmt(Stmt *S) { -+ assert(IsSemiMissing && "Operation supported for printing values only!"); -+ Statement = S; -+ } -+ bool isValuePrinting() const { return IsSemiMissing; } -+ void setValuePrinting(bool Missing = true) { IsSemiMissing = Missing; } - - static bool classof(const Decl *D) { return classofKind(D->getKind()); } - static bool classofKind(Kind K) { return K == TopLevelStmt; } -diff --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def -index 96feae991ccb..7526298557a8 100644 ---- a/clang/include/clang/Basic/TokenKinds.def -+++ b/clang/include/clang/Basic/TokenKinds.def -@@ -936,6 +936,9 @@ ANNOTATION(module_end) - // into the name of a header unit. - ANNOTATION(header_unit) - -+// Annotation for end of input in clang-repl. -+ANNOTATION(repl_input_end) -+ - #undef PRAGMA_ANNOTATION - #undef ANNOTATION - #undef TESTING_KEYWORD -diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h -index fd22af976613..020c2a88c185 100644 ---- a/clang/include/clang/Interpreter/Interpreter.h -+++ b/clang/include/clang/Interpreter/Interpreter.h -@@ -14,13 +14,15 @@ - #ifndef LLVM_CLANG_INTERPRETER_INTERPRETER_H - #define LLVM_CLANG_INTERPRETER_INTERPRETER_H - --#include "clang/Interpreter/PartialTranslationUnit.h" -- -+#include "clang/AST/Decl.h" - #include "clang/AST/GlobalDecl.h" -+#include "clang/Interpreter/PartialTranslationUnit.h" -+#include "clang/Interpreter/Value.h" - -+#include "llvm/ADT/DenseMap.h" - #include "llvm/ExecutionEngine/JITSymbol.h" -+#include "llvm/ExecutionEngine/Orc/Shared/ExecutorAddress.h" - #include "llvm/Support/Error.h" -- - #include - #include - -@@ -52,22 +54,24 @@ class Interpreter { - - Interpreter(std::unique_ptr CI, llvm::Error &Err); - -+ llvm::Error CreateExecutor(); -+ unsigned InitPTUSize = 0; -+ -+ Value LastValue; -+ - public: - ~Interpreter(); - static llvm::Expected> - create(std::unique_ptr CI); -+ const ASTContext &getASTContext() const; -+ ASTContext &getASTContext(); - const CompilerInstance *getCompilerInstance() const; -- const llvm::orc::LLJIT *getExecutionEngine() const; -+ llvm::Expected getExecutionEngine(); -+ - llvm::Expected Parse(llvm::StringRef Code); - llvm::Error Execute(PartialTranslationUnit &T); -- llvm::Error ParseAndExecute(llvm::StringRef Code) { -- auto PTU = Parse(Code); -- if (!PTU) -- return PTU.takeError(); -- if (PTU->TheModule) -- return Execute(*PTU); -- return llvm::Error::success(); -- } -+ llvm::Error ParseAndExecute(llvm::StringRef Code, Value *V = nullptr); -+ llvm::Expected CompileDtorCall(CXXRecordDecl *CXXRD); - - /// Undo N previous incremental inputs. - llvm::Error Undo(unsigned N = 1); -@@ -75,16 +79,33 @@ public: - /// \returns the \c JITTargetAddress of a \c GlobalDecl. This interface uses - /// the CodeGenModule's internal mangling cache to avoid recomputing the - /// mangled name. -- llvm::Expected getSymbolAddress(GlobalDecl GD) const; -+ llvm::Expected getSymbolAddress(GlobalDecl GD) const; - - /// \returns the \c JITTargetAddress of a given name as written in the IR. -- llvm::Expected -+ llvm::Expected - getSymbolAddress(llvm::StringRef IRName) const; - - /// \returns the \c JITTargetAddress of a given name as written in the object - /// file. -- llvm::Expected -+ llvm::Expected - getSymbolAddressFromLinkerName(llvm::StringRef LinkerName) const; -+ -+ size_t getEffectivePTUSize() const; -+ -+ enum InterfaceKind { NoAlloc, WithAlloc, CopyArray }; -+ -+ const llvm::SmallVectorImpl &getValuePrintingInfo() const { -+ return ValuePrintingInfo; -+ } -+ -+ Expr *SynthesizeExpr(Expr *E); -+ -+private: -+ bool FindRuntimeInterface(); -+ -+ llvm::DenseMap Dtors; -+ -+ llvm::SmallVector ValuePrintingInfo; - }; - } // namespace clang - -diff --git a/clang/include/clang/Interpreter/Value.h b/clang/include/clang/Interpreter/Value.h -new file mode 100644 -index 000000000000..f70eb214b3d5 ---- /dev/null -+++ b/clang/include/clang/Interpreter/Value.h -@@ -0,0 +1,174 @@ -+//===--- Value.h - Incremental Compiation and Execution---*- C++ -*-===// -+// -+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -+// See https://llvm.org/LICENSE.txt for license information. -+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -+// -+//===----------------------------------------------------------------------===// -+// -+// This file defines the class that used to represent a value in incremental -+// C++. -+// -+//===----------------------------------------------------------------------===// -+ -+#ifndef LLVM_CLANG_INTERPRETER_VALUE_H -+#define LLVM_CLANG_INTERPRETER_VALUE_H -+ -+#include -+ -+namespace llvm { -+class raw_ostream; -+ -+} // namespace llvm -+ -+namespace clang { -+ -+class ASTContext; -+class Interpreter; -+class QualType; -+ -+#if __has_attribute(visibility) && \ -+ (!(defined(_WIN32) || defined(__CYGWIN__)) || \ -+ (defined(__MINGW32__) && defined(__clang__))) -+#if defined(LLVM_BUILD_LLVM_DYLIB) || defined(LLVM_BUILD_SHARED_LIBS) -+#define REPL_EXTERNAL_VISIBILITY __attribute__((visibility("default"))) -+#else -+#define REPL_EXTERNAL_VISIBILITY -+#endif -+#else -+#if defined(_WIN32) -+#define REPL_EXTERNAL_VISIBILITY __declspec(dllexport) -+#endif -+#endif -+ -+#define REPL_BUILTIN_TYPES \ -+ X(bool, Bool) \ -+ X(char, Char_S) \ -+ X(signed char, SChar) \ -+ X(unsigned char, UChar) \ -+ X(short, Short) \ -+ X(unsigned short, UShort) \ -+ X(int, Int) \ -+ X(unsigned int, UInt) \ -+ X(long, Long) \ -+ X(unsigned long, ULong) \ -+ X(long long, LongLong) \ -+ X(unsigned long long, ULongLong) \ -+ X(float, Float) \ -+ X(double, Double) \ -+ X(long double, LongDouble) -+ -+class REPL_EXTERNAL_VISIBILITY Value { -+ union Storage { -+#define X(type, name) type m_##name; -+ REPL_BUILTIN_TYPES -+#undef X -+ void *m_Ptr; -+ }; -+ -+public: -+ enum Kind { -+#define X(type, name) K_##name, -+ REPL_BUILTIN_TYPES -+#undef X -+ -+ K_Void, -+ K_PtrOrObj, -+ K_Unspecified -+ }; -+ -+ Value() = default; -+ Value(Interpreter *In, void *Ty); -+ Value(const Value &RHS); -+ Value(Value &&RHS) noexcept; -+ Value &operator=(const Value &RHS); -+ Value &operator=(Value &&RHS) noexcept; -+ ~Value(); -+ -+ void printType(llvm::raw_ostream &Out) const; -+ void printData(llvm::raw_ostream &Out) const; -+ void print(llvm::raw_ostream &Out) const; -+ void dump() const; -+ void clear(); -+ -+ ASTContext &getASTContext(); -+ const ASTContext &getASTContext() const; -+ Interpreter &getInterpreter(); -+ const Interpreter &getInterpreter() const; -+ QualType getType() const; -+ -+ bool isValid() const { return ValueKind != K_Unspecified; } -+ bool isVoid() const { return ValueKind == K_Void; } -+ bool isManuallyAlloc() const { return IsManuallyAlloc; } -+ Kind getKind() const { return ValueKind; } -+ void setKind(Kind K) { ValueKind = K; } -+ void setOpaqueType(void *Ty) { OpaqueType = Ty; } -+ -+ void *getPtr() const; -+ void setPtr(void *Ptr) { Data.m_Ptr = Ptr; } -+ -+ bool isPointerOrObjectType() const { return ValueKind == K_PtrOrObj; } -+ -+#define X(type, name) \ -+ void set##name(type Val) { Data.m_##name = Val; } \ -+ type get##name() const { return Data.m_##name; } -+ REPL_BUILTIN_TYPES -+#undef X -+ -+ // Allow castAs to be partially specialized. -+ template struct CastFwd { -+ static T cast(const Value &V) { -+ if (V.isPointerOrObjectType()) -+ return (T)(uintptr_t)V.getAs(); -+ if (!V.isValid() || V.isVoid()) { -+ return T(); -+ } -+ return V.getAs(); -+ } -+ }; -+ -+ template struct CastFwd { -+ static T *cast(const Value &V) { -+ if (V.isPointerOrObjectType()) -+ return (T *)(uintptr_t)V.getAs(); -+ return nullptr; -+ } -+ }; -+ -+ /// \brief Get the value with cast. -+ // -+ /// Get the value cast to T. This is similar to reinterpret_cast(value), -+ /// casting the value of builtins (except void), enums and pointers. -+ /// Values referencing an object are treated as pointers to the object. -+ template T castAs() const { return CastFwd::cast(*this); } -+ -+ /// \brief Get to the value with type checking casting the underlying -+ /// stored value to T. -+ template T getAs() const { -+ switch (ValueKind) { -+ default: -+ return T(); -+#define X(type, name) \ -+ case Value::K_##name: \ -+ return (T)Data.m_##name; -+ REPL_BUILTIN_TYPES -+#undef X -+ } -+ } -+ -+private: -+ Interpreter *Interp = nullptr; -+ void *OpaqueType = nullptr; -+ Storage Data; -+ Kind ValueKind = K_Unspecified; -+ bool IsManuallyAlloc = false; -+}; -+ -+template <> inline void *Value::getAs() const { -+ if (isPointerOrObjectType()) -+ return Data.m_Ptr; -+ return (void *)getAs(); -+} -+ -+} // namespace clang -+#endif -diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h -index 6f9581b9ea1f..6b73f43a1fb7 100644 ---- a/clang/include/clang/Parse/Parser.h -+++ b/clang/include/clang/Parse/Parser.h -@@ -18,6 +18,7 @@ - #include "clang/Basic/OpenMPKinds.h" - #include "clang/Basic/OperatorPrecedence.h" - #include "clang/Basic/Specifiers.h" -+#include "clang/Basic/TokenKinds.h" - #include "clang/Lex/CodeCompletionHandler.h" - #include "clang/Lex/Preprocessor.h" - #include "clang/Sema/DeclSpec.h" -@@ -692,7 +693,8 @@ private: - bool isEofOrEom() { - tok::TokenKind Kind = Tok.getKind(); - return Kind == tok::eof || Kind == tok::annot_module_begin || -- Kind == tok::annot_module_end || Kind == tok::annot_module_include; -+ Kind == tok::annot_module_end || Kind == tok::annot_module_include || -+ Kind == tok::annot_repl_input_end; - } - - /// Checks if the \p Level is valid for use in a fold expression. -diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp -index 12d602fed693..8952920a09ad 100644 ---- a/clang/lib/CodeGen/CodeGenModule.cpp -+++ b/clang/lib/CodeGen/CodeGenModule.cpp -@@ -7186,8 +7186,14 @@ void CodeGenModule::printPostfixForExternalizedDecl(llvm::raw_ostream &OS, - } - - void CodeGenModule::moveLazyEmissionStates(CodeGenModule *NewBuilder) { -- assert(DeferredDeclsToEmit.empty() && -- "Should have emitted all decls deferred to emit."); -+ // FIXME: Re-enable the assertions once we fix regular codegen to not leave -+ // weak references behind. -+ // The code example also leaves entries in WeakRefReferences in regular clang. -+ // #include -+ // auto p = std::make_shared(42); -+ // -+ // assert(DeferredDeclsToEmit.empty() && -+ // "Should have emitted all decls deferred to emit."); - assert(NewBuilder->DeferredDecls.empty() && - "Newly created module should not have deferred decls"); - NewBuilder->DeferredDecls = std::move(DeferredDecls); -diff --git a/clang/lib/Frontend/PrintPreprocessedOutput.cpp b/clang/lib/Frontend/PrintPreprocessedOutput.cpp -index ffa85e523c03..1b262d9e6f7c 100644 ---- a/clang/lib/Frontend/PrintPreprocessedOutput.cpp -+++ b/clang/lib/Frontend/PrintPreprocessedOutput.cpp -@@ -663,7 +663,8 @@ void PrintPPOutputPPCallbacks::HandleWhitespaceBeforeTok(const Token &Tok, - // them. - if (Tok.is(tok::eof) || - (Tok.isAnnotation() && !Tok.is(tok::annot_header_unit) && -- !Tok.is(tok::annot_module_begin) && !Tok.is(tok::annot_module_end))) -+ !Tok.is(tok::annot_module_begin) && !Tok.is(tok::annot_module_end) && -+ !Tok.is(tok::annot_repl_input_end))) - return; - - // EmittedDirectiveOnThisLine takes priority over RequireSameLine. -@@ -819,6 +820,9 @@ static void PrintPreprocessedTokens(Preprocessor &PP, Token &Tok, - // -traditional-cpp the lexer keeps /all/ whitespace, including comments. - PP.Lex(Tok); - continue; -+ } else if (Tok.is(tok::annot_repl_input_end)) { -+ PP.Lex(Tok); -+ continue; - } else if (Tok.is(tok::eod)) { - // Don't print end of directive tokens, since they are typically newlines - // that mess up our line tracking. These come from unknown pre-processor -diff --git a/clang/lib/Interpreter/CMakeLists.txt b/clang/lib/Interpreter/CMakeLists.txt -index c49f22fddd8e..565e824bf0c9 100644 ---- a/clang/lib/Interpreter/CMakeLists.txt -+++ b/clang/lib/Interpreter/CMakeLists.txt -@@ -12,6 +12,8 @@ add_clang_library(clangInterpreter - IncrementalExecutor.cpp - IncrementalParser.cpp - Interpreter.cpp -+ InterpreterUtils.cpp -+ Value.cpp - - DEPENDS - intrinsics_gen -diff --git a/clang/lib/Interpreter/IncrementalExecutor.cpp b/clang/lib/Interpreter/IncrementalExecutor.cpp -index 37d230b61f76..b4b6bc8b40fe 100644 ---- a/clang/lib/Interpreter/IncrementalExecutor.cpp -+++ b/clang/lib/Interpreter/IncrementalExecutor.cpp -@@ -86,15 +86,12 @@ llvm::Error IncrementalExecutor::runCtors() const { - return Jit->initialize(Jit->getMainJITDylib()); - } - --llvm::Expected -+llvm::Expected - IncrementalExecutor::getSymbolAddress(llvm::StringRef Name, - SymbolNameKind NameKind) const { - auto Sym = (NameKind == LinkerName) ? Jit->lookupLinkerMangled(Name) - : Jit->lookup(Name); -- -- if (!Sym) -- return Sym.takeError(); -- return Sym->getValue(); -+ return std::move(Sym); - } - - } // end namespace clang -diff --git a/clang/lib/Interpreter/IncrementalExecutor.h b/clang/lib/Interpreter/IncrementalExecutor.h -index 54d37c76326b..8bead233d448 100644 ---- a/clang/lib/Interpreter/IncrementalExecutor.h -+++ b/clang/lib/Interpreter/IncrementalExecutor.h -@@ -51,9 +51,9 @@ public: - llvm::Error removeModule(PartialTranslationUnit &PTU); - llvm::Error runCtors() const; - llvm::Error cleanUp(); -- llvm::Expected -+ llvm::Expected - getSymbolAddress(llvm::StringRef Name, SymbolNameKind NameKind) const; -- llvm::orc::LLJIT *getExecutionEngine() const { return Jit.get(); } -+ llvm::orc::LLJIT &GetExecutionEngine() { return *Jit; } - }; - - } // end namespace clang -diff --git a/clang/lib/Interpreter/IncrementalParser.cpp b/clang/lib/Interpreter/IncrementalParser.cpp -index 373e2844b4e4..259b9875e6be 100644 ---- a/clang/lib/Interpreter/IncrementalParser.cpp -+++ b/clang/lib/Interpreter/IncrementalParser.cpp -@@ -11,7 +11,6 @@ - //===----------------------------------------------------------------------===// - - #include "IncrementalParser.h" -- - #include "clang/AST/DeclContextInternals.h" - #include "clang/CodeGen/BackendUtil.h" - #include "clang/CodeGen/CodeGenAction.h" -@@ -19,9 +18,9 @@ - #include "clang/Frontend/CompilerInstance.h" - #include "clang/Frontend/FrontendAction.h" - #include "clang/FrontendTool/Utils.h" -+#include "clang/Interpreter/Interpreter.h" - #include "clang/Parse/Parser.h" - #include "clang/Sema/Sema.h" -- - #include "llvm/Option/ArgList.h" - #include "llvm/Support/CrashRecoveryContext.h" - #include "llvm/Support/Error.h" -@@ -31,6 +30,79 @@ - - namespace clang { - -+class IncrementalASTConsumer final : public ASTConsumer { -+ Interpreter &Interp; -+ std::unique_ptr Consumer; -+ -+public: -+ IncrementalASTConsumer(Interpreter &InterpRef, std::unique_ptr C) -+ : Interp(InterpRef), Consumer(std::move(C)) {} -+ -+ bool HandleTopLevelDecl(DeclGroupRef DGR) override final { -+ if (DGR.isNull()) -+ return true; -+ if (!Consumer) -+ return true; -+ -+ for (Decl *D : DGR) -+ if (auto *TSD = llvm::dyn_cast(D); -+ TSD && TSD->isValuePrinting()) -+ TSD->setStmt(Interp.SynthesizeExpr(cast(TSD->getStmt()))); -+ -+ return Consumer->HandleTopLevelDecl(DGR); -+ } -+ void HandleTranslationUnit(ASTContext &Ctx) override final { -+ Consumer->HandleTranslationUnit(Ctx); -+ } -+ void HandleInlineFunctionDefinition(FunctionDecl *D) override final { -+ Consumer->HandleInlineFunctionDefinition(D); -+ } -+ void HandleInterestingDecl(DeclGroupRef D) override final { -+ Consumer->HandleInterestingDecl(D); -+ } -+ void HandleTagDeclDefinition(TagDecl *D) override final { -+ Consumer->HandleTagDeclDefinition(D); -+ } -+ void HandleTagDeclRequiredDefinition(const TagDecl *D) override final { -+ Consumer->HandleTagDeclRequiredDefinition(D); -+ } -+ void HandleCXXImplicitFunctionInstantiation(FunctionDecl *D) override final { -+ Consumer->HandleCXXImplicitFunctionInstantiation(D); -+ } -+ void HandleTopLevelDeclInObjCContainer(DeclGroupRef D) override final { -+ Consumer->HandleTopLevelDeclInObjCContainer(D); -+ } -+ void HandleImplicitImportDecl(ImportDecl *D) override final { -+ Consumer->HandleImplicitImportDecl(D); -+ } -+ void CompleteTentativeDefinition(VarDecl *D) override final { -+ Consumer->CompleteTentativeDefinition(D); -+ } -+ void CompleteExternalDeclaration(VarDecl *D) override final { -+ Consumer->CompleteExternalDeclaration(D); -+ } -+ void AssignInheritanceModel(CXXRecordDecl *RD) override final { -+ Consumer->AssignInheritanceModel(RD); -+ } -+ void HandleCXXStaticMemberVarInstantiation(VarDecl *D) override final { -+ Consumer->HandleCXXStaticMemberVarInstantiation(D); -+ } -+ void HandleVTable(CXXRecordDecl *RD) override final { -+ Consumer->HandleVTable(RD); -+ } -+ ASTMutationListener *GetASTMutationListener() override final { -+ return Consumer->GetASTMutationListener(); -+ } -+ ASTDeserializationListener *GetASTDeserializationListener() override final { -+ return Consumer->GetASTDeserializationListener(); -+ } -+ void PrintStats() override final { Consumer->PrintStats(); } -+ bool shouldSkipFunctionBody(Decl *D) override final { -+ return Consumer->shouldSkipFunctionBody(D); -+ } -+ static bool classof(const clang::ASTConsumer *) { return true; } -+}; -+ - /// A custom action enabling the incremental processing functionality. - /// - /// The usual \p FrontendAction expects one call to ExecuteAction and once it -@@ -122,7 +194,8 @@ public: - } - }; - --IncrementalParser::IncrementalParser(std::unique_ptr Instance, -+IncrementalParser::IncrementalParser(Interpreter &Interp, -+ std::unique_ptr Instance, - llvm::LLVMContext &LLVMCtx, - llvm::Error &Err) - : CI(std::move(Instance)) { -@@ -131,6 +204,9 @@ IncrementalParser::IncrementalParser(std::unique_ptr Instance, - if (Err) - return; - CI->ExecuteAction(*Act); -+ std::unique_ptr IncrConsumer = -+ std::make_unique(Interp, CI->takeASTConsumer()); -+ CI->setASTConsumer(std::move(IncrConsumer)); - Consumer = &CI->getASTConsumer(); - P.reset( - new Parser(CI->getPreprocessor(), CI->getSema(), /*SkipBodies=*/false)); -@@ -158,8 +234,8 @@ IncrementalParser::ParseOrWrapTopLevelDecl() { - LastPTU.TUPart = C.getTranslationUnitDecl(); - - // Skip previous eof due to last incremental input. -- if (P->getCurToken().is(tok::eof)) { -- P->ConsumeToken(); -+ if (P->getCurToken().is(tok::annot_repl_input_end)) { -+ P->ConsumeAnyToken(); - // FIXME: Clang does not call ExitScope on finalizing the regular TU, we - // might want to do that around HandleEndOfTranslationUnit. - P->ExitScope(); -@@ -259,25 +335,30 @@ IncrementalParser::Parse(llvm::StringRef input) { - Token Tok; - do { - PP.Lex(Tok); -- } while (Tok.isNot(tok::eof)); -+ } while (Tok.isNot(tok::annot_repl_input_end)); - } - - Token AssertTok; - PP.Lex(AssertTok); -- assert(AssertTok.is(tok::eof) && -+ assert(AssertTok.is(tok::annot_repl_input_end) && - "Lexer must be EOF when starting incremental parse!"); - -- if (CodeGenerator *CG = getCodeGen(Act.get())) { -- std::unique_ptr M(CG->ReleaseModule()); -- CG->StartModule("incr_module_" + std::to_string(PTUs.size()), -- M->getContext()); -- -+ if (std::unique_ptr M = GenModule()) - PTU->TheModule = std::move(M); -- } - - return PTU; - } - -+std::unique_ptr IncrementalParser::GenModule() { -+ static unsigned ID = 0; -+ if (CodeGenerator *CG = getCodeGen(Act.get())) { -+ std::unique_ptr M(CG->ReleaseModule()); -+ CG->StartModule("incr_module_" + std::to_string(ID++), M->getContext()); -+ return M; -+ } -+ return nullptr; -+} -+ - void IncrementalParser::CleanUpPTU(PartialTranslationUnit &PTU) { - TranslationUnitDecl *MostRecentTU = PTU.TUPart; - TranslationUnitDecl *FirstTU = MostRecentTU->getFirstDecl(); -diff --git a/clang/lib/Interpreter/IncrementalParser.h b/clang/lib/Interpreter/IncrementalParser.h -index 8e45d6b5931b..99e37588df9d 100644 ---- a/clang/lib/Interpreter/IncrementalParser.h -+++ b/clang/lib/Interpreter/IncrementalParser.h -@@ -16,7 +16,6 @@ - #include "clang/Interpreter/PartialTranslationUnit.h" - - #include "clang/AST/GlobalDecl.h" -- - #include "llvm/ADT/ArrayRef.h" - #include "llvm/ADT/StringRef.h" - #include "llvm/Support/Error.h" -@@ -31,8 +30,8 @@ namespace clang { - class ASTConsumer; - class CompilerInstance; - class IncrementalAction; -+class Interpreter; - class Parser; -- - /// Provides support for incremental compilation. Keeps track of the state - /// changes between the subsequent incremental input. - /// -@@ -57,7 +56,8 @@ class IncrementalParser { - std::list PTUs; - - public: -- IncrementalParser(std::unique_ptr Instance, -+ IncrementalParser(Interpreter &Interp, -+ std::unique_ptr Instance, - llvm::LLVMContext &LLVMCtx, llvm::Error &Err); - ~IncrementalParser(); - -@@ -76,6 +76,8 @@ public: - - std::list &getPTUs() { return PTUs; } - -+ std::unique_ptr GenModule(); -+ - private: - llvm::Expected ParseOrWrapTopLevelDecl(); - }; -diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp -index a6f5fdc6eefc..e8a07eceb678 100644 ---- a/clang/lib/Interpreter/Interpreter.cpp -+++ b/clang/lib/Interpreter/Interpreter.cpp -@@ -16,7 +16,11 @@ - #include "IncrementalExecutor.h" - #include "IncrementalParser.h" - -+#include "InterpreterUtils.h" - #include "clang/AST/ASTContext.h" -+#include "clang/AST/Mangle.h" -+#include "clang/AST/TypeVisitor.h" -+#include "clang/Basic/DiagnosticSema.h" - #include "clang/Basic/TargetInfo.h" - #include "clang/CodeGen/ModuleBuilder.h" - #include "clang/CodeGen/ObjectFilePCHContainerOperations.h" -@@ -27,12 +31,15 @@ - #include "clang/Driver/Tool.h" - #include "clang/Frontend/CompilerInstance.h" - #include "clang/Frontend/TextDiagnosticBuffer.h" -+#include "clang/Interpreter/Value.h" - #include "clang/Lex/PreprocessorOptions.h" -- -+#include "clang/Sema/Lookup.h" -+#include "llvm/ExecutionEngine/JITSymbol.h" -+#include "llvm/ExecutionEngine/Orc/LLJIT.h" - #include "llvm/IR/Module.h" - #include "llvm/Support/Errc.h" --#include "llvm/Support/Host.h" -- -+#include "llvm/Support/raw_ostream.h" -+#include "llvm/TargetParser/Host.h" - using namespace clang; - - // FIXME: Figure out how to unify with namespace init_convenience from -@@ -176,7 +183,7 @@ Interpreter::Interpreter(std::unique_ptr CI, - llvm::ErrorAsOutParameter EAO(&Err); - auto LLVMCtx = std::make_unique(); - TSCtx = std::make_unique(std::move(LLVMCtx)); -- IncrParser = std::make_unique(std::move(CI), -+ IncrParser = std::make_unique(*this, std::move(CI), - *TSCtx->getContext(), Err); - } - -@@ -189,6 +196,28 @@ Interpreter::~Interpreter() { - } - } - -+// These better to put in a runtime header but we can't. This is because we -+// can't find the precise resource directory in unittests so we have to hard -+// code them. -+const char *const Runtimes = R"( -+ void* operator new(__SIZE_TYPE__, void* __p) noexcept; -+ void *__clang_Interpreter_SetValueWithAlloc(void*, void*, void*); -+ void __clang_Interpreter_SetValueNoAlloc(void*, void*, void*, void*); -+ void __clang_Interpreter_SetValueNoAlloc(void*, void*, void*, float); -+ void __clang_Interpreter_SetValueNoAlloc(void*, void*, void*, double); -+ void __clang_Interpreter_SetValueNoAlloc(void*, void*, void*, long double); -+ void __clang_Interpreter_SetValueNoAlloc(void*,void*,void*,unsigned long long); -+ template -+ void __clang_Interpreter_SetValueCopyArr(T* Src, void* Placement, unsigned long Size) { -+ for (auto Idx = 0; Idx < Size; ++Idx) -+ new ((void*)(((T*)Placement) + Idx)) T(Src[Idx]); -+ } -+ template -+ void __clang_Interpreter_SetValueCopyArr(const T (*Src)[N], void* Placement, unsigned long Size) { -+ __clang_Interpreter_SetValueCopyArr(Src[0], Placement, Size); -+ } -+)"; -+ - llvm::Expected> - Interpreter::create(std::unique_ptr CI) { - llvm::Error Err = llvm::Error::success(); -@@ -196,6 +225,14 @@ Interpreter::create(std::unique_ptr CI) { - std::unique_ptr(new Interpreter(std::move(CI), Err)); - if (Err) - return std::move(Err); -+ if (llvm::Error Err = Interp->ParseAndExecute(Runtimes)) -+ return std::move(Err); -+ -+ Interp->ValuePrintingInfo.resize(3); -+ // FIXME: This is a ugly hack. Undo command checks its availability by looking -+ // at the size of the PTU list. However we have parsed something in the -+ // beginning of the REPL so we have to mark them as 'Irrevocable'. -+ Interp->InitPTUSize = Interp->IncrParser->getPTUs().size(); - return std::move(Interp); - } - -@@ -203,19 +240,50 @@ const CompilerInstance *Interpreter::getCompilerInstance() const { - return IncrParser->getCI(); - } - --const llvm::orc::LLJIT *Interpreter::getExecutionEngine() const { -- if (IncrExecutor) -- return IncrExecutor->getExecutionEngine(); -- return nullptr; -+llvm::Expected Interpreter::getExecutionEngine() { -+ if (!IncrExecutor) { -+ if (auto Err = CreateExecutor()) -+ return std::move(Err); -+ } -+ -+ return IncrExecutor->GetExecutionEngine(); -+} -+ -+ASTContext &Interpreter::getASTContext() { -+ return getCompilerInstance()->getASTContext(); -+} -+ -+const ASTContext &Interpreter::getASTContext() const { -+ return getCompilerInstance()->getASTContext(); -+} -+ -+size_t Interpreter::getEffectivePTUSize() const { -+ std::list &PTUs = IncrParser->getPTUs(); -+ assert(PTUs.size() >= InitPTUSize && "empty PTU list?"); -+ return PTUs.size() - InitPTUSize; - } - - llvm::Expected - Interpreter::Parse(llvm::StringRef Code) { -+ // Tell the interpreter sliently ignore unused expressions since value -+ // printing could cause it. -+ getCompilerInstance()->getDiagnostics().setSeverity( -+ clang::diag::warn_unused_expr, diag::Severity::Ignored, SourceLocation()); - return IncrParser->Parse(Code); - } - -+llvm::Error Interpreter::CreateExecutor() { -+ const clang::TargetInfo &TI = -+ getCompilerInstance()->getASTContext().getTargetInfo(); -+ llvm::Error Err = llvm::Error::success(); -+ auto Executor = std::make_unique(*TSCtx, Err, TI); -+ if (!Err) -+ IncrExecutor = std::move(Executor); -+ -+ return Err; -+} -+ - llvm::Error Interpreter::Execute(PartialTranslationUnit &T) { -- assert(T.TheModule); - if (!IncrExecutor) { - const clang::TargetInfo &TI = - getCompilerInstance()->getASTContext().getTargetInfo(); -@@ -235,7 +303,26 @@ llvm::Error Interpreter::Execute(PartialTranslationUnit &T) { - return llvm::Error::success(); - } - --llvm::Expected -+llvm::Error Interpreter::ParseAndExecute(llvm::StringRef Code, Value *V) { -+ -+ auto PTU = Parse(Code); -+ if (!PTU) -+ return PTU.takeError(); -+ if (PTU->TheModule) -+ if (llvm::Error Err = Execute(*PTU)) -+ return Err; -+ -+ if (LastValue.isValid()) { -+ if (!V) { -+ LastValue.dump(); -+ LastValue.clear(); -+ } else -+ *V = std::move(LastValue); -+ } -+ return llvm::Error::success(); -+} -+ -+llvm::Expected - Interpreter::getSymbolAddress(GlobalDecl GD) const { - if (!IncrExecutor) - return llvm::make_error("Operation failed. " -@@ -245,7 +332,7 @@ Interpreter::getSymbolAddress(GlobalDecl GD) const { - return getSymbolAddress(MangledName); - } - --llvm::Expected -+llvm::Expected - Interpreter::getSymbolAddress(llvm::StringRef IRName) const { - if (!IncrExecutor) - return llvm::make_error("Operation failed. " -@@ -255,7 +342,7 @@ Interpreter::getSymbolAddress(llvm::StringRef IRName) const { - return IncrExecutor->getSymbolAddress(IRName, IncrementalExecutor::IRName); - } - --llvm::Expected -+llvm::Expected - Interpreter::getSymbolAddressFromLinkerName(llvm::StringRef Name) const { - if (!IncrExecutor) - return llvm::make_error("Operation failed. " -@@ -268,7 +355,7 @@ Interpreter::getSymbolAddressFromLinkerName(llvm::StringRef Name) const { - llvm::Error Interpreter::Undo(unsigned N) { - - std::list &PTUs = IncrParser->getPTUs(); -- if (N > PTUs.size()) -+ if (N > getEffectivePTUSize()) - return llvm::make_error("Operation failed. " - "Too many undos", - std::error_code()); -@@ -283,3 +370,305 @@ llvm::Error Interpreter::Undo(unsigned N) { - } - return llvm::Error::success(); - } -+ -+llvm::Expected -+Interpreter::CompileDtorCall(CXXRecordDecl *CXXRD) { -+ assert(CXXRD && "Cannot compile a destructor for a nullptr"); -+ if (auto Dtor = Dtors.find(CXXRD); Dtor != Dtors.end()) -+ return Dtor->getSecond(); -+ -+ if (CXXRD->hasIrrelevantDestructor()) -+ return llvm::orc::ExecutorAddr{}; -+ -+ CXXDestructorDecl *DtorRD = -+ getCompilerInstance()->getSema().LookupDestructor(CXXRD); -+ -+ llvm::StringRef Name = -+ IncrParser->GetMangledName(GlobalDecl(DtorRD, Dtor_Base)); -+ auto AddrOrErr = getSymbolAddress(Name); -+ if (!AddrOrErr) -+ return AddrOrErr.takeError(); -+ -+ Dtors[CXXRD] = *AddrOrErr; -+ return AddrOrErr; -+} -+ -+static constexpr llvm::StringRef MagicRuntimeInterface[] = { -+ "__clang_Interpreter_SetValueNoAlloc", -+ "__clang_Interpreter_SetValueWithAlloc", -+ "__clang_Interpreter_SetValueCopyArr"}; -+ -+bool Interpreter::FindRuntimeInterface() { -+ if (llvm::all_of(ValuePrintingInfo, [](Expr *E) { return E != nullptr; })) -+ return true; -+ -+ Sema &S = getCompilerInstance()->getSema(); -+ ASTContext &Ctx = S.getASTContext(); -+ -+ auto LookupInterface = [&](Expr *&Interface, llvm::StringRef Name) { -+ LookupResult R(S, &Ctx.Idents.get(Name), SourceLocation(), -+ Sema::LookupOrdinaryName, Sema::ForVisibleRedeclaration); -+ S.LookupQualifiedName(R, Ctx.getTranslationUnitDecl()); -+ if (R.empty()) -+ return false; -+ -+ CXXScopeSpec CSS; -+ Interface = S.BuildDeclarationNameExpr(CSS, R, /*ADL=*/false).get(); -+ return true; -+ }; -+ -+ if (!LookupInterface(ValuePrintingInfo[NoAlloc], -+ MagicRuntimeInterface[NoAlloc])) -+ return false; -+ if (!LookupInterface(ValuePrintingInfo[WithAlloc], -+ MagicRuntimeInterface[WithAlloc])) -+ return false; -+ if (!LookupInterface(ValuePrintingInfo[CopyArray], -+ MagicRuntimeInterface[CopyArray])) -+ return false; -+ return true; -+} -+ -+namespace { -+ -+class RuntimeInterfaceBuilder -+ : public TypeVisitor { -+ clang::Interpreter &Interp; -+ ASTContext &Ctx; -+ Sema &S; -+ Expr *E; -+ llvm::SmallVector Args; -+ -+public: -+ RuntimeInterfaceBuilder(clang::Interpreter &In, ASTContext &C, Sema &SemaRef, -+ Expr *VE, ArrayRef FixedArgs) -+ : Interp(In), Ctx(C), S(SemaRef), E(VE) { -+ // The Interpreter* parameter and the out parameter `OutVal`. -+ for (Expr *E : FixedArgs) -+ Args.push_back(E); -+ -+ // Get rid of ExprWithCleanups. -+ if (auto *EWC = llvm::dyn_cast_if_present(E)) -+ E = EWC->getSubExpr(); -+ } -+ -+ExprResult getCall() { -+ QualType Ty = E->getType(); -+ QualType DesugaredTy = Ty.getDesugaredType(Ctx); -+ -+ // For lvalue struct, we treat it as a reference. -+ if (DesugaredTy->isRecordType() && E->isLValue()) { -+ DesugaredTy = Ctx.getLValueReferenceType(DesugaredTy); -+ Ty = Ctx.getLValueReferenceType(Ty); -+ } -+ -+ Expr *TypeArg = -+ CStyleCastPtrExpr(S, Ctx.VoidPtrTy, (uintptr_t)Ty.getAsOpaquePtr()); -+ // The QualType parameter `OpaqueType`, represented as `void*`. -+ Args.push_back(TypeArg); -+ -+ // We push the last parameter based on the type of the Expr. Note we need -+ // special care for rvalue struct. -+ Interpreter::InterfaceKind Kind = Visit(&*DesugaredTy); -+ switch (Kind) { -+ case Interpreter::InterfaceKind::WithAlloc: -+ case Interpreter::InterfaceKind::CopyArray: { -+ // __clang_Interpreter_SetValueWithAlloc. -+ ExprResult AllocCall = S.ActOnCallExpr( -+ /*Scope=*/nullptr, -+ Interp.getValuePrintingInfo()[Interpreter::InterfaceKind::WithAlloc], -+ E->getBeginLoc(), Args, E->getEndLoc()); -+ assert(!AllocCall.isInvalid() && "Can't create runtime interface call!"); -+ -+ TypeSourceInfo *TSI = Ctx.getTrivialTypeSourceInfo(Ty, SourceLocation()); -+ -+ // Force CodeGen to emit destructor. -+ if (auto *RD = Ty->getAsCXXRecordDecl()) { -+ auto *Dtor = S.LookupDestructor(RD); -+ Dtor->addAttr(UsedAttr::CreateImplicit(Ctx)); -+ Interp.getCompilerInstance()->getASTConsumer().HandleTopLevelDecl( -+ DeclGroupRef(Dtor)); -+ } -+ -+ // __clang_Interpreter_SetValueCopyArr. -+ if (Kind == Interpreter::InterfaceKind::CopyArray) { -+ const auto *ConstantArrTy = -+ cast(DesugaredTy.getTypePtr()); -+ size_t ArrSize = Ctx.getConstantArrayElementCount(ConstantArrTy); -+ Expr *ArrSizeExpr = IntegerLiteralExpr(Ctx, ArrSize); -+ Expr *Args[] = {E, AllocCall.get(), ArrSizeExpr}; -+ return S.ActOnCallExpr( -+ /*Scope *=*/nullptr, -+ Interp -+ .getValuePrintingInfo()[Interpreter::InterfaceKind::CopyArray], -+ SourceLocation(), Args, SourceLocation()); -+ } -+ Expr *Args[] = {AllocCall.get()}; -+ ExprResult CXXNewCall = S.BuildCXXNew( -+ E->getSourceRange(), -+ /*UseGlobal=*/true, /*PlacementLParen=*/SourceLocation(), Args, -+ /*PlacementRParen=*/SourceLocation(), -+ /*TypeIdParens=*/SourceRange(), TSI->getType(), TSI, std::nullopt, -+ E->getSourceRange(), E); -+ -+ assert(!CXXNewCall.isInvalid() && -+ "Can't create runtime placement new call!"); -+ -+ return S.ActOnFinishFullExpr(CXXNewCall.get(), -+ /*DiscardedValue=*/false); -+ } -+ // __clang_Interpreter_SetValueNoAlloc. -+ case Interpreter::InterfaceKind::NoAlloc: { -+ return S.ActOnCallExpr( -+ /*Scope=*/nullptr, -+ Interp.getValuePrintingInfo()[Interpreter::InterfaceKind::NoAlloc], -+ E->getBeginLoc(), Args, E->getEndLoc()); -+ } -+ default: llvm_unreachable("Unknown InterfaceKind."); -+ } -+ } -+ -+ Interpreter::InterfaceKind VisitRecordType(const RecordType *Ty) { -+ return Interpreter::InterfaceKind::WithAlloc; -+ } -+ -+ Interpreter::InterfaceKind -+ VisitMemberPointerType(const MemberPointerType *Ty) { -+ llvm_unreachable("Not implemented yet"); -+ } -+ Interpreter::InterfaceKind -+ VisitConstantArrayType(const ConstantArrayType *Ty) { -+ return Interpreter::InterfaceKind::CopyArray; -+ } -+ -+ Interpreter::InterfaceKind VisitPointerType(const PointerType *Ty) { -+ TypeSourceInfo *TSI = Ctx.getTrivialTypeSourceInfo(Ctx.VoidPtrTy); -+ ExprResult CastedExpr = -+ S.BuildCStyleCastExpr(SourceLocation(), TSI, SourceLocation(), E); -+ assert(!CastedExpr.isInvalid() && "Can not create cstyle cast expression"); -+ Args.push_back(CastedExpr.get()); -+ return Interpreter::InterfaceKind::NoAlloc; -+ } -+ -+ Interpreter::InterfaceKind VisitReferenceType(const ReferenceType *Ty) { -+ ExprResult AddrOfE = S.CreateBuiltinUnaryOp(SourceLocation(), UO_AddrOf, E); -+ assert(!AddrOfE.isInvalid() && "Can not create unary expression"); -+ Args.push_back(AddrOfE.get()); -+ return Interpreter::InterfaceKind::NoAlloc; -+ } -+ -+ Interpreter::InterfaceKind VisitBuiltinType(const BuiltinType *Ty) { -+ if (Ty->isNullPtrType()) -+ Args.push_back(E); -+ else if (Ty->isFloatingType()) -+ Args.push_back(E); -+ else if (Ty->isVoidType()) -+ Args.push_back(E); -+ else if (Ty->isIntegralOrEnumerationType()) -+ HandleIntegralOrEnumType(Ty); -+ -+ return Interpreter::InterfaceKind::NoAlloc; -+ } -+ -+ Interpreter::InterfaceKind VisitEnumType(const EnumType *Ty) { -+ HandleIntegralOrEnumType(Ty); -+ return Interpreter::InterfaceKind::NoAlloc; -+ } -+ -+private: -+ // Force cast these types to uint64 to reduce the number of overloads of -+ // `__clang_Interpreter_SetValueNoAlloc`. -+ void HandleIntegralOrEnumType(const Type *Ty) { -+ TypeSourceInfo *TSI = Ctx.getTrivialTypeSourceInfo(Ctx.UnsignedLongLongTy); -+ ExprResult CastedExpr = -+ S.BuildCStyleCastExpr(SourceLocation(), TSI, SourceLocation(), E); -+ assert(!CastedExpr.isInvalid() && "Cannot create cstyle cast expr"); -+ Args.push_back(CastedExpr.get()); -+ } -+}; -+} // namespace -+ -+// This synthesizes a call expression to a speciall -+// function that is responsible for generating the Value. -+// In general, we transform: -+// clang-repl> x -+// To: -+// // 1. If x is a built-in type like int, float. -+// __clang_Interpreter_SetValueNoAlloc(ThisInterp, OpaqueValue, xQualType, x); -+// // 2. If x is a struct, and a lvalue. -+// __clang_Interpreter_SetValueNoAlloc(ThisInterp, OpaqueValue, xQualType, -+// &x); -+// // 3. If x is a struct, but a rvalue. -+// new (__clang_Interpreter_SetValueWithAlloc(ThisInterp, OpaqueValue, -+// xQualType)) (x); -+ -+Expr *Interpreter::SynthesizeExpr(Expr *E) { -+ Sema &S = getCompilerInstance()->getSema(); -+ ASTContext &Ctx = S.getASTContext(); -+ -+ if (!FindRuntimeInterface()) -+ llvm_unreachable("We can't find the runtime iterface for pretty print!"); -+ -+ // Create parameter `ThisInterp`. -+ auto *ThisInterp = CStyleCastPtrExpr(S, Ctx.VoidPtrTy, (uintptr_t)this); -+ -+ // Create parameter `OutVal`. -+ auto *OutValue = CStyleCastPtrExpr(S, Ctx.VoidPtrTy, (uintptr_t)&LastValue); -+ -+ // Build `__clang_Interpreter_SetValue*` call. -+ RuntimeInterfaceBuilder Builder(*this, Ctx, S, E, {ThisInterp, OutValue}); -+ -+ ExprResult Result = Builder.getCall(); -+ assert(!Result.isInvalid() && "Failed to generate the CallExpr!"); -+ return Result.get(); -+} -+ -+// Temporary rvalue struct that need special care. -+REPL_EXTERNAL_VISIBILITY void * -+__clang_Interpreter_SetValueWithAlloc(void *This, void *OutVal, -+ void *OpaqueType) { -+ Value &VRef = *(Value *)OutVal; -+ VRef = Value(static_cast(This), OpaqueType); -+ return VRef.getPtr(); -+} -+ -+// Pointers, lvalue struct that can take as a reference. -+REPL_EXTERNAL_VISIBILITY void -+__clang_Interpreter_SetValueNoAlloc(void *This, void *OutVal, void *OpaqueType, -+ void *Val) { -+ Value &VRef = *(Value *)OutVal; -+ VRef = Value(static_cast(This), OpaqueType); -+ VRef.setPtr(Val); -+} -+ -+REPL_EXTERNAL_VISIBILITY void -+__clang_Interpreter_SetValueNoAlloc(void *This, void *OutVal, void *OpaqueType, -+ unsigned long long Val) { -+ Value &VRef = *(Value *)OutVal; -+ VRef = Value(static_cast(This), OpaqueType); -+ VRef.setULongLong(Val); -+} -+ -+REPL_EXTERNAL_VISIBILITY void -+__clang_Interpreter_SetValueNoAlloc(void *This, void *OutVal, void *OpaqueType, -+ float Val) { -+ Value &VRef = *(Value *)OutVal; -+ VRef = Value(static_cast(This), OpaqueType); -+ VRef.setFloat(Val); -+} -+ -+REPL_EXTERNAL_VISIBILITY void -+__clang_Interpreter_SetValueNoAlloc(void *This, void *OutVal, void *OpaqueType, -+ double Val) { -+ Value &VRef = *(Value *)OutVal; -+ VRef = Value(static_cast(This), OpaqueType); -+ VRef.setDouble(Val); -+} -+ -+REPL_EXTERNAL_VISIBILITY void -+__clang_Interpreter_SetValueNoAlloc(void *This, void *OutVal, void *OpaqueType, -+ long double Val) { -+ Value &VRef = *(Value *)OutVal; -+ VRef = Value(static_cast(This), OpaqueType); -+ VRef.setLongDouble(Val); -+} -diff --git a/clang/lib/Interpreter/InterpreterUtils.cpp b/clang/lib/Interpreter/InterpreterUtils.cpp -new file mode 100644 -index 000000000000..1ec2bf54987b ---- /dev/null -+++ b/clang/lib/Interpreter/InterpreterUtils.cpp -@@ -0,0 +1,111 @@ -+//===--- InterpreterUtils.cpp - Incremental Utils --------*- C++ -*-===// -+// -+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -+// See https://llvm.org/LICENSE.txt for license information. -+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -+// -+//===----------------------------------------------------------------------===// -+// -+// This file implements some common utils used in the incremental library. -+// -+//===----------------------------------------------------------------------===// -+ -+#include "InterpreterUtils.h" -+ -+namespace clang { -+ -+IntegerLiteral *IntegerLiteralExpr(ASTContext &C, uintptr_t Ptr) { -+ const llvm::APInt Addr(8 * sizeof(void *), Ptr); -+ return IntegerLiteral::Create(C, Addr, C.getUIntPtrType(), SourceLocation()); -+} -+ -+Expr *CStyleCastPtrExpr(Sema &S, QualType Ty, Expr *E) { -+ ASTContext &Ctx = S.getASTContext(); -+ if (!Ty->isPointerType()) -+ Ty = Ctx.getPointerType(Ty); -+ -+ TypeSourceInfo *TSI = Ctx.getTrivialTypeSourceInfo(Ty, SourceLocation()); -+ Expr *Result = -+ S.BuildCStyleCastExpr(SourceLocation(), TSI, SourceLocation(), E).get(); -+ assert(Result && "Cannot create CStyleCastPtrExpr"); -+ return Result; -+} -+ -+Expr *CStyleCastPtrExpr(Sema &S, QualType Ty, uintptr_t Ptr) { -+ ASTContext &Ctx = S.getASTContext(); -+ return CStyleCastPtrExpr(S, Ty, IntegerLiteralExpr(Ctx, Ptr)); -+} -+ -+Sema::DeclGroupPtrTy CreateDGPtrFrom(Sema &S, Decl *D) { -+ SmallVector DeclsInGroup; -+ DeclsInGroup.push_back(D); -+ Sema::DeclGroupPtrTy DeclGroupPtr = S.BuildDeclaratorGroup(DeclsInGroup); -+ return DeclGroupPtr; -+} -+ -+NamespaceDecl *LookupNamespace(Sema &S, llvm::StringRef Name, -+ const DeclContext *Within) { -+ DeclarationName DName = &S.Context.Idents.get(Name); -+ LookupResult R(S, DName, SourceLocation(), -+ Sema::LookupNestedNameSpecifierName); -+ R.suppressDiagnostics(); -+ if (!Within) -+ S.LookupName(R, S.TUScope); -+ else { -+ if (const clang::TagDecl *TD = dyn_cast(Within)) -+ if (!TD->getDefinition()) -+ // No definition, no lookup result. -+ return nullptr; -+ -+ S.LookupQualifiedName(R, const_cast(Within)); -+ } -+ -+ if (R.empty()) -+ return nullptr; -+ -+ R.resolveKind(); -+ -+ return dyn_cast(R.getFoundDecl()); -+} -+ -+NamedDecl *LookupNamed(Sema &S, llvm::StringRef Name, -+ const DeclContext *Within) { -+ DeclarationName DName = &S.Context.Idents.get(Name); -+ LookupResult R(S, DName, SourceLocation(), Sema::LookupOrdinaryName, -+ Sema::ForVisibleRedeclaration); -+ -+ R.suppressDiagnostics(); -+ -+ if (!Within) -+ S.LookupName(R, S.TUScope); -+ else { -+ const DeclContext *PrimaryWithin = nullptr; -+ if (const auto *TD = dyn_cast(Within)) -+ PrimaryWithin = llvm::dyn_cast_or_null(TD->getDefinition()); -+ else -+ PrimaryWithin = Within->getPrimaryContext(); -+ -+ // No definition, no lookup result. -+ if (!PrimaryWithin) -+ return nullptr; -+ -+ S.LookupQualifiedName(R, const_cast(PrimaryWithin)); -+ } -+ -+ if (R.empty()) -+ return nullptr; -+ R.resolveKind(); -+ -+ if (R.isSingleResult()) -+ return llvm::dyn_cast(R.getFoundDecl()); -+ -+ return nullptr; -+} -+ -+std::string GetFullTypeName(ASTContext &Ctx, QualType QT) { -+ PrintingPolicy Policy(Ctx.getPrintingPolicy()); -+ Policy.SuppressScope = false; -+ Policy.AnonymousTagLocations = false; -+ return QT.getAsString(Policy); -+} -+} // namespace clang -diff --git a/clang/lib/Interpreter/InterpreterUtils.h b/clang/lib/Interpreter/InterpreterUtils.h -new file mode 100644 -index 000000000000..61f4cf7e1239 ---- /dev/null -+++ b/clang/lib/Interpreter/InterpreterUtils.h -@@ -0,0 +1,55 @@ -+//===--- InterpreterUtils.h - Incremental Utils --------*- C++ -*-===// -+// -+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -+// See https://llvm.org/LICENSE.txt for license information. -+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -+// -+//===----------------------------------------------------------------------===// -+// -+// This file implements some common utils used in the incremental library. -+// -+//===----------------------------------------------------------------------===// -+ -+#ifndef LLVM_CLANG_INTERPRETER_UTILS_H -+#define LLVM_CLANG_INTERPRETER_UTILS_H -+ -+#include "clang/AST/ASTContext.h" -+#include "clang/AST/Mangle.h" -+#include "clang/AST/TypeVisitor.h" -+#include "clang/Basic/TargetInfo.h" -+#include "clang/CodeGen/ModuleBuilder.h" -+#include "clang/CodeGen/ObjectFilePCHContainerOperations.h" -+#include "clang/Driver/Compilation.h" -+#include "clang/Driver/Driver.h" -+#include "clang/Driver/Job.h" -+#include "clang/Driver/Options.h" -+#include "clang/Driver/Tool.h" -+#include "clang/Frontend/CompilerInstance.h" -+#include "clang/Frontend/TextDiagnosticBuffer.h" -+#include "clang/Lex/PreprocessorOptions.h" -+ -+#include "clang/Sema/Lookup.h" -+#include "llvm/IR/Module.h" -+#include "llvm/Support/Errc.h" -+#include "llvm/TargetParser/Host.h" -+ -+// TODO: create a sub namespace `repl`. -+namespace clang { -+IntegerLiteral *IntegerLiteralExpr(ASTContext &C, uintptr_t Ptr); -+ -+Expr *CStyleCastPtrExpr(Sema &S, QualType Ty, Expr *E); -+ -+Expr *CStyleCastPtrExpr(Sema &S, QualType Ty, uintptr_t Ptr); -+ -+Sema::DeclGroupPtrTy CreateDGPtrFrom(Sema &S, Decl *D); -+ -+NamespaceDecl *LookupNamespace(Sema &S, llvm::StringRef Name, -+ const DeclContext *Within = nullptr); -+ -+NamedDecl *LookupNamed(Sema &S, llvm::StringRef Name, -+ const DeclContext *Within); -+ -+std::string GetFullTypeName(ASTContext &Ctx, QualType QT); -+} // namespace clang -+ -+#endif -diff --git a/clang/lib/Interpreter/Value.cpp b/clang/lib/Interpreter/Value.cpp -new file mode 100644 -index 000000000000..34a77344bc86 ---- /dev/null -+++ b/clang/lib/Interpreter/Value.cpp -@@ -0,0 +1,260 @@ -+//===--- Interpreter.h - Incremental Compiation and Execution---*- C++ -*-===// -+// -+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -+// See https://llvm.org/LICENSE.txt for license information. -+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -+// -+//===----------------------------------------------------------------------===// -+// -+// This file defines the class that used to represent a value in incremental -+// C++. -+// -+//===----------------------------------------------------------------------===// -+ -+#include "clang/Interpreter/Value.h" -+#include "clang/AST/ASTContext.h" -+#include "clang/AST/Type.h" -+#include "clang/Interpreter/Interpreter.h" -+#include "llvm/ADT/StringExtras.h" -+#include "llvm/Support/ErrorHandling.h" -+#include "llvm/Support/raw_os_ostream.h" -+#include -+#include -+#include -+ -+using namespace clang; -+ -+namespace { -+ -+// This is internal buffer maintained by Value, used to hold temporaries. -+class ValueStorage { -+public: -+ using DtorFunc = void (*)(void *); -+ -+ static unsigned char *CreatePayload(void *DtorF, size_t AllocSize, -+ size_t ElementsSize) { -+ if (AllocSize < sizeof(Canary)) -+ AllocSize = sizeof(Canary); -+ unsigned char *Buf = -+ new unsigned char[ValueStorage::getPayloadOffset() + AllocSize]; -+ ValueStorage *VS = new (Buf) ValueStorage(DtorF, AllocSize, ElementsSize); -+ std::memcpy(VS->getPayload(), Canary, sizeof(Canary)); -+ return VS->getPayload(); -+ } -+ -+ unsigned char *getPayload() { return Storage; } -+ const unsigned char *getPayload() const { return Storage; } -+ -+ static unsigned getPayloadOffset() { -+ static ValueStorage Dummy(nullptr, 0, 0); -+ return Dummy.getPayload() - reinterpret_cast(&Dummy); -+ } -+ -+ static ValueStorage *getFromPayload(void *Payload) { -+ ValueStorage *R = reinterpret_cast( -+ (unsigned char *)Payload - getPayloadOffset()); -+ return R; -+ } -+ -+ void Retain() { ++RefCnt; } -+ -+ void Release() { -+ assert(RefCnt > 0 && "Can't release if reference count is already zero"); -+ if (--RefCnt == 0) { -+ // We hace a non-trivial dtor. -+ if (Dtor && IsAlive()) { -+ assert(Elements && "We at least should have 1 element in Value"); -+ size_t Stride = AllocSize / Elements; -+ for (size_t Idx = 0; Idx < Elements; ++Idx) -+ (*Dtor)(getPayload() + Idx * Stride); -+ } -+ delete[] reinterpret_cast(this); -+ } -+ } -+ -+ bool IsAlive() const { -+ return std::memcmp(getPayload(), Canary, sizeof(Canary)) != 0; -+ } -+ -+private: -+ ValueStorage(void *DtorF, size_t AllocSize, size_t ElementsNum) -+ : RefCnt(1), Dtor(reinterpret_cast(DtorF)), -+ AllocSize(AllocSize), Elements(ElementsNum) {} -+ -+ mutable unsigned RefCnt; -+ DtorFunc Dtor = nullptr; -+ size_t AllocSize = 0; -+ size_t Elements = 0; -+ unsigned char Storage[1]; -+ -+ static constexpr unsigned char Canary[8] = {0x4c, 0x37, 0xad, 0x8f, -+ 0x2d, 0x23, 0x95, 0x91}; -+}; -+} // namespace -+ -+static Value::Kind ConvertQualTypeToKind(const ASTContext &Ctx, QualType QT) { -+ if (Ctx.hasSameType(QT, Ctx.VoidTy)) -+ return Value::K_Void; -+ -+ if (const auto *ET = dyn_cast(QT.getTypePtr())) -+ QT = ET->getDecl()->getIntegerType(); -+ -+ if (!QT->isBuiltinType() || QT->castAs()->isNullPtrType()) -+ return Value::K_PtrOrObj; -+ -+ switch (QT->getAs()->getKind()) { -+ default: -+ assert(false && "Type not supported"); -+ return Value::K_Unspecified; -+#define X(type, name) \ -+ case BuiltinType::name: \ -+ return Value::K_##name; -+ REPL_BUILTIN_TYPES -+#undef X -+ } -+} -+ -+Value::Value(Interpreter *In, void *Ty) : Interp(In), OpaqueType(Ty) { -+ setKind(ConvertQualTypeToKind(getASTContext(), getType())); -+ if (ValueKind == K_PtrOrObj) { -+ QualType Canon = getType().getCanonicalType(); -+ if ((Canon->isPointerType() || Canon->isObjectType() || -+ Canon->isReferenceType()) && -+ (Canon->isRecordType() || Canon->isConstantArrayType() || -+ Canon->isMemberPointerType())) { -+ IsManuallyAlloc = true; -+ // Compile dtor function. -+ Interpreter &Interp = getInterpreter(); -+ void *DtorF = nullptr; -+ size_t ElementsSize = 1; -+ QualType DtorTy = getType(); -+ -+ if (const auto *ArrTy = -+ llvm::dyn_cast(DtorTy.getTypePtr())) { -+ DtorTy = ArrTy->getElementType(); -+ llvm::APInt ArrSize(sizeof(size_t) * 8, 1); -+ do { -+ ArrSize *= ArrTy->getSize(); -+ ArrTy = llvm::dyn_cast( -+ ArrTy->getElementType().getTypePtr()); -+ } while (ArrTy); -+ ElementsSize = static_cast(ArrSize.getZExtValue()); -+ } -+ if (const auto *RT = DtorTy->getAs()) { -+ if (CXXRecordDecl *CXXRD = -+ llvm::dyn_cast(RT->getDecl())) { -+ if (llvm::Expected Addr = -+ Interp.CompileDtorCall(CXXRD)) -+ DtorF = reinterpret_cast(Addr->getValue()); -+ else -+ llvm::logAllUnhandledErrors(Addr.takeError(), llvm::errs()); -+ } -+ } -+ -+ size_t AllocSize = -+ getASTContext().getTypeSizeInChars(getType()).getQuantity(); -+ unsigned char *Payload = -+ ValueStorage::CreatePayload(DtorF, AllocSize, ElementsSize); -+ setPtr((void *)Payload); -+ } -+ } -+} -+ -+Value::Value(const Value &RHS) -+ : Interp(RHS.Interp), OpaqueType(RHS.OpaqueType), Data(RHS.Data), -+ ValueKind(RHS.ValueKind), IsManuallyAlloc(RHS.IsManuallyAlloc) { -+ if (IsManuallyAlloc) -+ ValueStorage::getFromPayload(getPtr())->Retain(); -+} -+ -+Value::Value(Value &&RHS) noexcept { -+ Interp = std::exchange(RHS.Interp, nullptr); -+ OpaqueType = std::exchange(RHS.OpaqueType, nullptr); -+ Data = RHS.Data; -+ ValueKind = std::exchange(RHS.ValueKind, K_Unspecified); -+ IsManuallyAlloc = std::exchange(RHS.IsManuallyAlloc, false); -+ -+ if (IsManuallyAlloc) -+ ValueStorage::getFromPayload(getPtr())->Release(); -+} -+ -+Value &Value::operator=(const Value &RHS) { -+ if (IsManuallyAlloc) -+ ValueStorage::getFromPayload(getPtr())->Release(); -+ -+ Interp = RHS.Interp; -+ OpaqueType = RHS.OpaqueType; -+ Data = RHS.Data; -+ ValueKind = RHS.ValueKind; -+ IsManuallyAlloc = RHS.IsManuallyAlloc; -+ -+ if (IsManuallyAlloc) -+ ValueStorage::getFromPayload(getPtr())->Retain(); -+ -+ return *this; -+} -+ -+Value &Value::operator=(Value &&RHS) noexcept { -+ if (IsManuallyAlloc) -+ ValueStorage::getFromPayload(getPtr())->Release(); -+ -+ Interp = std::exchange(RHS.Interp, nullptr); -+ OpaqueType = std::exchange(RHS.OpaqueType, nullptr); -+ ValueKind = std::exchange(RHS.ValueKind, K_Unspecified); -+ IsManuallyAlloc = std::exchange(RHS.IsManuallyAlloc, false); -+ -+ Data = RHS.Data; -+ -+ return *this; -+} -+ -+void Value::clear() { -+ if (IsManuallyAlloc) -+ ValueStorage::getFromPayload(getPtr())->Release(); -+ ValueKind = K_Unspecified; -+ OpaqueType = nullptr; -+ Interp = nullptr; -+ IsManuallyAlloc = false; -+} -+ -+Value::~Value() { clear(); } -+ -+void *Value::getPtr() const { -+ assert(ValueKind == K_PtrOrObj); -+ return Data.m_Ptr; -+} -+ -+QualType Value::getType() const { -+ return QualType::getFromOpaquePtr(OpaqueType); -+} -+ -+Interpreter &Value::getInterpreter() { -+ assert(Interp != nullptr && -+ "Can't get interpreter from a default constructed value"); -+ return *Interp; -+} -+ -+const Interpreter &Value::getInterpreter() const { -+ assert(Interp != nullptr && -+ "Can't get interpreter from a default constructed value"); -+ return *Interp; -+} -+ -+ASTContext &Value::getASTContext() { return getInterpreter().getASTContext(); } -+ -+const ASTContext &Value::getASTContext() const { -+ return getInterpreter().getASTContext(); -+} -+ -+void Value::dump() const { print(llvm::outs()); } -+ -+void Value::printType(llvm::raw_ostream &Out) const { -+ Out << "Not implement yet.\n"; -+} -+void Value::printData(llvm::raw_ostream &Out) const { -+ Out << "Not implement yet.\n"; -+} -+void Value::print(llvm::raw_ostream &Out) const { -+ assert(OpaqueType != nullptr && "Can't print default Value"); -+ Out << "Not implement yet.\n"; -+} -diff --git a/clang/lib/Lex/PPLexerChange.cpp b/clang/lib/Lex/PPLexerChange.cpp -index 66168467ecf5..0822f83b58df 100644 ---- a/clang/lib/Lex/PPLexerChange.cpp -+++ b/clang/lib/Lex/PPLexerChange.cpp -@@ -526,13 +526,19 @@ bool Preprocessor::HandleEndOfFile(Token &Result, bool isEndOfMacro) { - return LeavingSubmodule; - } - } -- - // If this is the end of the main file, form an EOF token. - assert(CurLexer && "Got EOF but no current lexer set!"); - const char *EndPos = getCurLexerEndPos(); - Result.startToken(); - CurLexer->BufferPtr = EndPos; -- CurLexer->FormTokenWithChars(Result, EndPos, tok::eof); -+ -+ if (isIncrementalProcessingEnabled()) { -+ CurLexer->FormTokenWithChars(Result, EndPos, tok::annot_repl_input_end); -+ Result.setAnnotationEndLoc(Result.getLocation()); -+ Result.setAnnotationValue(nullptr); -+ } else { -+ CurLexer->FormTokenWithChars(Result, EndPos, tok::eof); -+ } - - if (isCodeCompletionEnabled()) { - // Inserting the code-completion point increases the source buffer by 1, -diff --git a/clang/lib/Parse/ParseCXXInlineMethods.cpp b/clang/lib/Parse/ParseCXXInlineMethods.cpp -index 3a7f5426d4a7..57a3dfba4f1d 100644 ---- a/clang/lib/Parse/ParseCXXInlineMethods.cpp -+++ b/clang/lib/Parse/ParseCXXInlineMethods.cpp -@@ -836,6 +836,7 @@ bool Parser::ConsumeAndStoreUntil(tok::TokenKind T1, tok::TokenKind T2, - case tok::annot_module_begin: - case tok::annot_module_end: - case tok::annot_module_include: -+ case tok::annot_repl_input_end: - // Ran out of tokens. - return false; - -@@ -1242,6 +1243,7 @@ bool Parser::ConsumeAndStoreInitializer(CachedTokens &Toks, - case tok::annot_module_begin: - case tok::annot_module_end: - case tok::annot_module_include: -+ case tok::annot_repl_input_end: - // Ran out of tokens. - return false; - -diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp -index e6812ac72c88..3e64e6f074fc 100644 ---- a/clang/lib/Parse/ParseDecl.cpp -+++ b/clang/lib/Parse/ParseDecl.cpp -@@ -2030,6 +2030,7 @@ void Parser::SkipMalformedDecl() { - case tok::annot_module_begin: - case tok::annot_module_end: - case tok::annot_module_include: -+ case tok::annot_repl_input_end: - return; - - default: -@@ -5394,6 +5395,13 @@ Parser::DeclGroupPtrTy Parser::ParseTopLevelStmtDecl() { - - SmallVector DeclsInGroup; - DeclsInGroup.push_back(Actions.ActOnTopLevelStmtDecl(R.get())); -+ -+ if (Tok.is(tok::annot_repl_input_end) && -+ Tok.getAnnotationValue() != nullptr) { -+ ConsumeAnnotationToken(); -+ cast(DeclsInGroup.back())->setValuePrinting(); -+ } -+ - // Currently happens for things like -fms-extensions and use `__if_exists`. - for (Stmt *S : Stmts) - DeclsInGroup.push_back(Actions.ActOnTopLevelStmtDecl(S)); -diff --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp -index 1c8441fafc48..446fc7ea6385 100644 ---- a/clang/lib/Parse/ParseStmt.cpp -+++ b/clang/lib/Parse/ParseStmt.cpp -@@ -543,9 +543,22 @@ StmtResult Parser::ParseExprStatement(ParsedStmtContext StmtCtx) { - return ParseCaseStatement(StmtCtx, /*MissingCase=*/true, Expr); - } - -- // Otherwise, eat the semicolon. -- ExpectAndConsumeSemi(diag::err_expected_semi_after_expr); -- return handleExprStmt(Expr, StmtCtx); -+ Token *CurTok = nullptr; -+ // If we're parsing an ExprStmt and the last semicolon is missing and the -+ // incremental externsion is enabled and we're reaching the end, consider we -+ // want to do value printing. Note we shouldn't eat the token since the -+ // callback need it. -+ if (PP.isIncrementalProcessingEnabled() && Tok.is(tok::annot_repl_input_end)) -+ CurTok = &Tok; -+ else -+ // Otherwise, eat the semicolon. -+ ExpectAndConsumeSemi(diag::err_expected_semi_after_expr); -+ -+ StmtResult R = handleExprStmt(Expr, StmtCtx); -+ if (!R.isInvalid() && CurTok) -+ CurTok->setAnnotationValue(R.get()); -+ -+ return R; - } - - /// ParseSEHTryBlockCommon -diff --git a/clang/lib/Parse/Parser.cpp b/clang/lib/Parse/Parser.cpp -index 6db3dc3156fd..9682f12ed211 100644 ---- a/clang/lib/Parse/Parser.cpp -+++ b/clang/lib/Parse/Parser.cpp -@@ -319,6 +319,7 @@ bool Parser::SkipUntil(ArrayRef Toks, SkipUntilFlags Flags) { - case tok::annot_module_begin: - case tok::annot_module_end: - case tok::annot_module_include: -+ case tok::annot_repl_input_end: - // Stop before we change submodules. They generally indicate a "good" - // place to pick up parsing again (except in the special case where - // we're trying to skip to EOF). -@@ -614,8 +615,8 @@ bool Parser::ParseTopLevelDecl(DeclGroupPtrTy &Result, - - // Skip over the EOF token, flagging end of previous input for incremental - // processing -- if (PP.isIncrementalProcessingEnabled() && Tok.is(tok::eof)) -- ConsumeToken(); -+ if (PP.isIncrementalProcessingEnabled() && Tok.is(tok::annot_repl_input_end)) -+ ConsumeAnnotationToken(); - - Result = nullptr; - switch (Tok.getKind()) { -@@ -695,6 +696,7 @@ bool Parser::ParseTopLevelDecl(DeclGroupPtrTy &Result, - return false; - - case tok::eof: -+ case tok::annot_repl_input_end: - // Check whether -fmax-tokens= was reached. - if (PP.getMaxTokens() != 0 && PP.getTokenCount() > PP.getMaxTokens()) { - PP.Diag(Tok.getLocation(), diag::warn_max_tokens_total) -diff --git a/clang/tools/clang-repl/CMakeLists.txt b/clang/tools/clang-repl/CMakeLists.txt -index b51a18c10cdc..15d7f9439ff5 100644 ---- a/clang/tools/clang-repl/CMakeLists.txt -+++ b/clang/tools/clang-repl/CMakeLists.txt -@@ -12,6 +12,7 @@ add_clang_tool(clang-repl - ) - - clang_target_link_libraries(clang-repl PRIVATE -+ clangAST - clangBasic - clangFrontend - clangInterpreter -diff --git a/clang/unittests/Interpreter/CMakeLists.txt b/clang/unittests/Interpreter/CMakeLists.txt -index 1a099dbbfe59..698494b9897f 100644 ---- a/clang/unittests/Interpreter/CMakeLists.txt -+++ b/clang/unittests/Interpreter/CMakeLists.txt -@@ -22,3 +22,5 @@ target_link_libraries(ClangReplInterpreterTests PUBLIC - if(NOT WIN32) - add_subdirectory(ExceptionTests) - endif() -+ -+export_executable_symbols(ClangReplInterpreterTests) -diff --git a/clang/unittests/Interpreter/InterpreterTest.cpp b/clang/unittests/Interpreter/InterpreterTest.cpp -index d4900a0e4de8..4e737a0b313a 100644 ---- a/clang/unittests/Interpreter/InterpreterTest.cpp -+++ b/clang/unittests/Interpreter/InterpreterTest.cpp -@@ -17,6 +17,7 @@ - #include "clang/AST/Mangle.h" - #include "clang/Frontend/CompilerInstance.h" - #include "clang/Frontend/TextDiagnosticPrinter.h" -+#include "clang/Interpreter/Value.h" - #include "clang/Sema/Lookup.h" - #include "clang/Sema/Sema.h" - -@@ -33,6 +34,10 @@ using namespace clang; - #define CLANG_INTERPRETER_NO_SUPPORT_EXEC - #endif - -+int Global = 42; -+int getGlobal() { return Global; } -+void setGlobal(int val) { Global = val; } -+ - namespace { - using Args = std::vector; - static std::unique_ptr -@@ -276,8 +281,7 @@ TEST(IncrementalProcessing, InstantiateTemplate) { - std::vector Args = {"-fno-delayed-template-parsing"}; - std::unique_ptr Interp = createInterpreter(Args); - -- llvm::cantFail(Interp->Parse("void* operator new(__SIZE_TYPE__, void* __p);" -- "extern \"C\" int printf(const char*,...);" -+ llvm::cantFail(Interp->Parse("extern \"C\" int printf(const char*,...);" - "class A {};" - "struct B {" - " template" -@@ -314,4 +318,55 @@ TEST(IncrementalProcessing, InstantiateTemplate) { - free(NewA); - } - -+TEST(InterpreterTest, Value) { -+ std::unique_ptr Interp = createInterpreter(); -+ -+ Value V1; -+ llvm::cantFail(Interp->ParseAndExecute("int x = 42;")); -+ llvm::cantFail(Interp->ParseAndExecute("x", &V1)); -+ EXPECT_TRUE(V1.isValid()); -+ EXPECT_EQ(V1.getInt(), 42); -+ EXPECT_TRUE(V1.getType()->isIntegerType()); -+ EXPECT_EQ(V1.getKind(), Value::K_Int); -+ EXPECT_FALSE(V1.isManuallyAlloc()); -+ EXPECT_FALSE(V1.isPointerOrObjectType()); -+ -+ Value V2; -+ llvm::cantFail(Interp->ParseAndExecute("double y = 3.14;")); -+ llvm::cantFail(Interp->ParseAndExecute("y", &V2)); -+ EXPECT_TRUE(V2.isValid()); -+ EXPECT_EQ(V2.getDouble(), 3.14); -+ EXPECT_TRUE(V2.getType()->isFloatingType()); -+ EXPECT_EQ(V2.getKind(), Value::K_Double); -+ EXPECT_FALSE(V2.isManuallyAlloc()); -+ EXPECT_FALSE(V2.isPointerOrObjectType()); -+ -+ Value V3; -+ llvm::cantFail(Interp->ParseAndExecute( -+ "struct S { int* p; S() { p = new int(42); } ~S() { delete p; }};")); -+ llvm::cantFail(Interp->ParseAndExecute("S{}", &V3)); -+ EXPECT_TRUE(V3.isValid()); -+ EXPECT_TRUE(V3.getType()->isRecordType()); -+ EXPECT_EQ(V3.getKind(), Value::K_PtrOrObj); -+ EXPECT_TRUE(V3.isManuallyAlloc()); -+ EXPECT_TRUE(V3.isPointerOrObjectType()); -+ -+ Value V4; -+ llvm::cantFail(Interp->ParseAndExecute("int getGlobal();")); -+ llvm::cantFail(Interp->ParseAndExecute("void setGlobal(int);")); -+ llvm::cantFail(Interp->ParseAndExecute("getGlobal()", &V4)); -+ EXPECT_EQ(V4.getInt(), 42); -+ EXPECT_TRUE(V4.getType()->isIntegerType()); -+ -+ Value V5; -+ // Change the global from the compiled code. -+ setGlobal(43); -+ llvm::cantFail(Interp->ParseAndExecute("getGlobal()", &V5)); -+ EXPECT_EQ(V5.getInt(), 43); -+ EXPECT_TRUE(V5.getType()->isIntegerType()); -+ -+ // Change the global from the interpreted code. -+ llvm::cantFail(Interp->ParseAndExecute("setGlobal(44);")); -+ EXPECT_EQ(getGlobal(), 44); -+} - } // end anonymous namespace diff --git a/patches/llvm/clang16-D148435-WeakRefs.patch b/patches/llvm/clang16-D148435-WeakRefs.patch deleted file mode 100644 index e04f43c2..00000000 --- a/patches/llvm/clang16-D148435-WeakRefs.patch +++ /dev/null @@ -1,33 +0,0 @@ -diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp -index 12d602fed693..e73102544361 100644 ---- a/clang/lib/CodeGen/CodeGenModule.cpp -+++ b/clang/lib/CodeGen/CodeGenModule.cpp -@@ -7202,7 +7202,6 @@ void CodeGenModule::moveLazyEmissionStates(CodeGenModule *NewBuilder) { - "Newly created module should not have manglings"); - NewBuilder->Manglings = std::move(Manglings); - -- assert(WeakRefReferences.empty() && "Not all WeakRefRefs have been applied"); - NewBuilder->WeakRefReferences = std::move(WeakRefReferences); - - NewBuilder->TBAA = std::move(TBAA); -diff --git a/clang/test/Interpreter/execute-weak.cpp b/clang/test/Interpreter/execute-weak.cpp -index 5b343512c545..66f2214ab03c 100644 ---- a/clang/test/Interpreter/execute-weak.cpp -+++ b/clang/test/Interpreter/execute-weak.cpp -@@ -2,11 +2,16 @@ - // RUN: clang-repl "int i = 10;" 'extern "C" int printf(const char*,...);' \ - // RUN: 'auto r1 = printf("i = %d\n", i);' | FileCheck --check-prefix=CHECK-DRIVER %s - // CHECK-DRIVER: i = 10 -+// - // UNSUPPORTED: system-aix, system-windows - // RUN: cat %s | clang-repl | FileCheck %s -+ - extern "C" int printf(const char *, ...); - int __attribute__((weak)) bar() { return 42; } - auto r4 = printf("bar() = %d\n", bar()); - // CHECK: bar() = 42 - -+int a = 12; -+static __typeof(a) b __attribute__((__weakref__("a"))); -+int c = b; - %quit