diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h index afb0bbc98079d..b3d64458d777c 100644 --- a/clang/include/clang/Interpreter/Interpreter.h +++ b/clang/include/clang/Interpreter/Interpreter.h @@ -41,34 +41,8 @@ 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. @@ -77,9 +51,6 @@ 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(); @@ -88,9 +59,6 @@ class Interpreter { ~Interpreter(); static llvm::Expected> create(std::unique_ptr CI); - static llvm::Expected> - createWithCUDA(std::unique_ptr CI, - std::unique_ptr DCI); const CompilerInstance *getCompilerInstance() const; llvm::Expected getExecutionEngine(); diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index c30a08a5722dc..1f429e4305790 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -24,7 +24,6 @@ #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; @@ -722,9 +721,8 @@ 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()) { - auto VFS = CGM.getFileSystem(); - auto CudaGpuBinaryOrErr = - VFS->getBufferForFile(CudaGpuBinaryFileName, -1, false); + llvm::ErrorOr> CudaGpuBinaryOrErr = + llvm::MemoryBuffer::getFileOrSTDIN(CudaGpuBinaryFileName); 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 784ff77c61727..29adf88acd704 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -264,7 +264,6 @@ 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 @@ -294,7 +293,6 @@ 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 803369009dfe4..5cd29d3657879 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6255,10 +6255,6 @@ 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 3594f4c66e677..e3e953c34c59f 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; + const CodeGenOptions CodeGenOpts; // Intentionally copied in. unsigned HandlingTopLevelDecls; diff --git a/clang/lib/Interpreter/CMakeLists.txt b/clang/lib/Interpreter/CMakeLists.txt index b2c4690163944..721864c0cc1ea 100644 --- a/clang/lib/Interpreter/CMakeLists.txt +++ b/clang/lib/Interpreter/CMakeLists.txt @@ -1,7 +1,6 @@ set(LLVM_LINK_COMPONENTS core native - MC Option OrcJit OrcShared @@ -15,7 +14,6 @@ add_clang_library(clangInterpreter IncrementalExecutor.cpp IncrementalParser.cpp Interpreter.cpp - DeviceOffload.cpp DEPENDS intrinsics_gen diff --git a/clang/lib/Interpreter/DeviceOffload.cpp b/clang/lib/Interpreter/DeviceOffload.cpp deleted file mode 100644 index 70f50e371a9ca..0000000000000 --- a/clang/lib/Interpreter/DeviceOffload.cpp +++ /dev/null @@ -1,176 +0,0 @@ -//===---------- 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( - std::unique_ptr Instance, IncrementalParser &HostParser, - llvm::LLVMContext &LLVMCtx, - llvm::IntrusiveRefCntPtr FS, - llvm::Error &Err) - : IncrementalParser(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 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 deleted file mode 100644 index ae76aff7244ba..0000000000000 --- a/clang/lib/Interpreter/DeviceOffload.h +++ /dev/null @@ -1,51 +0,0 @@ -//===----------- 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_OFFLOAD_H -#define LLVM_CLANG_LIB_INTERPRETER_OFFLOAD_H - -#include "IncrementalParser.h" -#include "llvm/Support/FileSystem.h" -#include "llvm/Support/VirtualFileSystem.h" - -namespace clang { - -class IncrementalCUDADeviceParser : public IncrementalParser { -public: - IncrementalCUDADeviceParser( - 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_OFFLOAD_H diff --git a/clang/lib/Interpreter/IncrementalParser.cpp b/clang/lib/Interpreter/IncrementalParser.cpp index 8af63625398f5..2b932623a5fea 100644 --- a/clang/lib/Interpreter/IncrementalParser.cpp +++ b/clang/lib/Interpreter/IncrementalParser.cpp @@ -122,15 +122,6 @@ class IncrementalAction : public WrapperFrontendAction { } }; -CodeGenerator *IncrementalParser::getCodeGen() const { - FrontendAction *WrappedAct = Act->getWrapped(); - if (!WrappedAct->hasIRSupport()) - return nullptr; - return static_cast(WrappedAct)->getCodeGenerator(); -} - -IncrementalParser::IncrementalParser() {} - IncrementalParser::IncrementalParser(std::unique_ptr Instance, llvm::LLVMContext &LLVMCtx, llvm::Error &Err) @@ -144,21 +135,6 @@ IncrementalParser::IncrementalParser(std::unique_ptr Instance, 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() { @@ -229,6 +205,14 @@ 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(); @@ -283,7 +267,7 @@ IncrementalParser::Parse(llvm::StringRef input) { "Lexer must be EOF when starting incremental parse!"); } - if (CodeGenerator *CG = getCodeGen()) { + if (CodeGenerator *CG = getCodeGen(Act.get())) { std::unique_ptr M(CG->ReleaseModule()); CG->StartModule("incr_module_" + std::to_string(PTUs.size()), M->getContext()); @@ -313,7 +297,7 @@ void IncrementalParser::CleanUpPTU(PartialTranslationUnit &PTU) { } llvm::StringRef IncrementalParser::GetMangledName(GlobalDecl GD) const { - CodeGenerator *CG = getCodeGen(); + CodeGenerator *CG = getCodeGen(Act.get()); assert(CG); return CG->GetMangledName(GD); } diff --git a/clang/lib/Interpreter/IncrementalParser.h b/clang/lib/Interpreter/IncrementalParser.h index 3427cde286857..8e45d6b5931bc 100644 --- a/clang/lib/Interpreter/IncrementalParser.h +++ b/clang/lib/Interpreter/IncrementalParser.h @@ -29,7 +29,6 @@ class LLVMContext; namespace clang { class ASTConsumer; -class CodeGenerator; class CompilerInstance; class IncrementalAction; class Parser; @@ -38,7 +37,6 @@ class Parser; /// changes between the subsequent incremental input. /// class IncrementalParser { -protected: /// Long-lived, incremental parsing action. std::unique_ptr Act; @@ -58,20 +56,17 @@ class IncrementalParser { /// of code. std::list PTUs; - IncrementalParser(); - public: IncrementalParser(std::unique_ptr Instance, llvm::LLVMContext &LLVMCtx, llvm::Error &Err); - virtual ~IncrementalParser(); + ~IncrementalParser(); - CompilerInstance *getCI() { return CI.get(); } - CodeGenerator *getCodeGen() const; + const CompilerInstance *getCI() const { return CI.get(); } /// 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. - virtual llvm::Expected Parse(llvm::StringRef Input); + 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 a9836f6f96b04..24fb9da69a8bc 100644 --- a/clang/lib/Interpreter/Interpreter.cpp +++ b/clang/lib/Interpreter/Interpreter.cpp @@ -15,11 +15,9 @@ #include "IncrementalExecutor.h" #include "IncrementalParser.h" -#include "DeviceOffload.h" #include "clang/AST/ASTContext.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" @@ -141,6 +139,7 @@ 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"); @@ -173,54 +172,6 @@ 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); @@ -249,34 +200,6 @@ 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 E; - - llvm::Error Err = llvm::Error::success(); - auto DeviceParser = std::make_unique( - 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(); } @@ -292,13 +215,6 @@ llvm::Expected Interpreter::getExecutionEngine() { 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 E; - } return IncrParser->Parse(Code); } @@ -363,7 +279,7 @@ Interpreter::getSymbolAddressFromLinkerName(llvm::StringRef Name) const { llvm::Error Interpreter::Undo(unsigned N) { std::list &PTUs = IncrParser->getPTUs(); - if (N >= PTUs.size()) + if (N > PTUs.size()) return llvm::make_error("Operation failed. " "Too many undos", std::error_code()); diff --git a/clang/test/Interpreter/CUDA/device-function-template.cu b/clang/test/Interpreter/CUDA/device-function-template.cu deleted file mode 100644 index f0077a2c51470..0000000000000 --- a/clang/test/Interpreter/CUDA/device-function-template.cu +++ /dev/null @@ -1,24 +0,0 @@ -// 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 deleted file mode 100644 index 396f8f0f93e0c..0000000000000 --- a/clang/test/Interpreter/CUDA/device-function.cu +++ /dev/null @@ -1,24 +0,0 @@ -// 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 deleted file mode 100644 index 8e44e34032704..0000000000000 --- a/clang/test/Interpreter/CUDA/host-and-device.cu +++ /dev/null @@ -1,27 +0,0 @@ -// 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 deleted file mode 100644 index 9991572462ad5..0000000000000 --- a/clang/test/Interpreter/CUDA/lit.local.cfg +++ /dev/null @@ -1,2 +0,0 @@ -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 deleted file mode 100644 index 852cc04f6de68..0000000000000 --- a/clang/test/Interpreter/CUDA/memory.cu +++ /dev/null @@ -1,23 +0,0 @@ -// 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 deleted file mode 100644 index ef9d68df464dd..0000000000000 --- a/clang/test/Interpreter/CUDA/sanity.cu +++ /dev/null @@ -1,11 +0,0 @@ -// 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 739ecf698e598..e9bfaf2e96774 100644 --- a/clang/test/lit.cfg.py +++ b/clang/test/lit.cfg.py @@ -87,41 +87,9 @@ 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, - input=testcode) - - except OSError: - print('could not exec clang-repl') - return False - - if clang_repl_cmd.returncode == 0: - if clang_repl_cmd.stdout.find(b"success") != -1: - return True - - print('could not run clang-repl with cuda') - 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 1552d65eb5332..33faf3fab58f0 100644 --- a/clang/tools/clang-repl/ClangRepl.cpp +++ b/clang/tools/clang-repl/ClangRepl.cpp @@ -23,10 +23,6 @@ #include "llvm/Support/TargetSelect.h" // llvm::Initialize* #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"), @@ -94,36 +90,9 @@ int main(int argc, const char **argv) { return 0; } - clang::IncrementalCompilerBuilder CB; - CB.SetCompilerArgs(ClangArgv); - - std::unique_ptr DeviceCI; - if (CudaEnabled) { - // initialize NVPTX backend - LLVMInitializeNVPTXTargetInfo(); - LLVMInitializeNVPTXTarget(); - LLVMInitializeNVPTXTargetMC(); - LLVMInitializeNVPTXAsmPrinter(); - - 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. - std::unique_ptr CI; - if (CudaEnabled) { - CI = ExitOnErr(CB.CreateCudaHost()); - } else { - CI = ExitOnErr(CB.CreateCpp()); - } + auto CI = ExitOnErr(clang::IncrementalCompilerBuilder::create(ClangArgv)); // Set an error handler, so that any LLVM backend diagnostics go through our // error handler. @@ -132,23 +101,8 @@ 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 70e10b1e53bd9..c82d11de20e0d 100644 --- a/clang/unittests/Interpreter/ExceptionTests/InterpreterExceptionTest.cpp +++ b/clang/unittests/Interpreter/ExceptionTests/InterpreterExceptionTest.cpp @@ -38,9 +38,7 @@ createInterpreter(const Args &ExtraArgs = {}, DiagnosticConsumer *Client = nullptr) { Args ClangArgs = {"-Xclang", "-emit-llvm-only"}; ClangArgs.insert(ClangArgs.end(), ExtraArgs.begin(), ExtraArgs.end()); - auto CB = clang::IncrementalCompilerBuilder(); - CB.SetCompilerArgs(ClangArgs); - auto CI = cantFail(CB.CreateCpp()); + auto CI = cantFail(clang::IncrementalCompilerBuilder::create(ClangArgs)); 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 f43b3ddac68f9..1f6df2aa226c4 100644 --- a/clang/unittests/Interpreter/IncrementalProcessingTest.cpp +++ b/clang/unittests/Interpreter/IncrementalProcessingTest.cpp @@ -52,9 +52,7 @@ const Function *getGlobalInit(llvm::Module *M) { TEST(IncrementalProcessing, EmitCXXGlobalInitFunc) { std::vector ClangArgv = {"-Xclang", "-emit-llvm-only"}; - auto CB = clang::IncrementalCompilerBuilder(); - CB.SetCompilerArgs(ClangArgv); - auto CI = cantFail(CB.CreateCpp()); + auto CI = llvm::cantFail(IncrementalCompilerBuilder::create(ClangArgv)); 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 5e03eeaf4daef..d555911a89451 100644 --- a/clang/unittests/Interpreter/InterpreterTest.cpp +++ b/clang/unittests/Interpreter/InterpreterTest.cpp @@ -40,9 +40,7 @@ createInterpreter(const Args &ExtraArgs = {}, DiagnosticConsumer *Client = nullptr) { Args ClangArgs = {"-Xclang", "-emit-llvm-only"}; ClangArgs.insert(ClangArgs.end(), ExtraArgs.begin(), ExtraArgs.end()); - auto CB = clang::IncrementalCompilerBuilder(); - CB.SetCompilerArgs(ClangArgs); - auto CI = cantFail(CB.CreateCpp()); + auto CI = cantFail(clang::IncrementalCompilerBuilder::create(ClangArgs)); if (Client) CI->getDiagnostics().setClient(Client, /*ShouldOwnClient=*/false); return cantFail(clang::Interpreter::create(std::move(CI)));