diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000000..6d727c70be --- /dev/null +++ b/CMakeLists.txt @@ -0,0 +1,42 @@ +cmake_minimum_required(VERSION 2.8.8) + +project(hipify) + +find_package(LLVM REQUIRED PATHS ${LLVM_DIR} NO_DEFAULT_PATH) + +message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}") +message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}") + +include_directories(${LLVM_INCLUDE_DIRS}) +link_directories(${LLVM_LIBRARY_DIRS}) +add_definitions(${LLVM_DEFINITIONS}) + +# Now build our tools +add_executable(hipify src/Cuda2Hip.cpp) + +# Link against LLVM and CLANG tools libraries +target_link_libraries(hipify + clangASTMatchers + clangFrontend + clangTooling + clangParse + clangSerialization + clangSema + clangEdit + clangLex + clangAnalysis + clangDriver + clangAST + clangToolingCore + clangRewrite + clangBasic + LLVMSupport + LLVMMCParser + LLVMMC + LLVMBitReader + LLVMOption + LLVMCore) + +set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11 -pthread -D __STDC_LIMIT_MACROS -D __STDC_CONSTANT_MACROS -D __STDC_FORMAT_MACROS -fno-rtti -fvisibility-inlines-hidden") diff --git a/README.md b/README.md index 91165785bb..a813325e65 100644 --- a/README.md +++ b/README.md @@ -1,2 +1,10 @@ # HIPIFY -The CLANG-based HIPIFY tool - translates CUDA to HIP +The CLANG-based HIPIFY tool - translates CUDA to [HIP](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/blob/master/README.md) + +# How to Build +- Get LLVM and CLANG sources and build CLANG, then install it to a **LLVM_INSTALL_PATH** folder +- make build dir for hipify and run the folloing commands there: +``` +> cmake -DLLVM_DIR="LLVM_INSTALL_PATH" path_to_hipify_src +> make +``` diff --git a/src/Cuda2Hip.cpp b/src/Cuda2Hip.cpp new file mode 100644 index 0000000000..c4c8f58810 --- /dev/null +++ b/src/Cuda2Hip.cpp @@ -0,0 +1,526 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/** + * @file Cuda2Hip.cpp + * + * This file is compiled and linked into clang based hipify tool. + */ +#include "clang/ASTMatchers/ASTMatchers.h" +#include "clang/ASTMatchers/ASTMatchFinder.h" +#include "clang/Basic/SourceManager.h" +#include "clang/Frontend/FrontendActions.h" +#include "clang/Lex/Lexer.h" +#include "clang/Tooling/CommonOptionsParser.h" +#include "clang/Tooling/Refactoring.h" +#include "clang/Tooling/Tooling.h" +#include "llvm/Support/CommandLine.h" +#include "llvm/Support/MemoryBuffer.h" +#include "llvm/Support/Signals.h" +#include "llvm/Support/Debug.h" +#include "clang/Frontend/TextDiagnosticPrinter.h" +#include "clang/Rewrite/Core/Rewriter.h" +#include "clang/Lex/MacroInfo.h" +#include "clang/Frontend/CompilerInstance.h" +#include "clang/Lex/Preprocessor.h" +#include "clang/Lex/PPCallbacks.h" + +#include +#include + +using namespace clang; +using namespace clang::ast_matchers; +using namespace clang::tooling; +using namespace llvm; + +#define DEBUG_TYPE "cuda2hip" + +namespace { + struct hipName { + hipName() { + // defines + cuda2hipRename["__CUDACC__"] = "__HIPCC__"; + + // includes + cuda2hipRename["cuda_runtime.h"] = "hip_runtime.h"; + cuda2hipRename["cuda_runtime_api.h"] = "hip_runtime_api.h"; + + // Error codes and return types: + cuda2hipRename["cudaError_t"] = "hipError_t"; + cuda2hipRename["cudaError"] = "hipError_t"; + cuda2hipRename["cudaSuccess"] = "hipSuccess"; + + cuda2hipRename["cudaErrorUnknown"] = "hipErrorUnknown"; + cuda2hipRename["cudaErrorMemoryAllocation"] = "hipErrorMemoryAllocation"; + cuda2hipRename["cudaErrorMemoryFree"] = "hipErrorMemoryFree"; + cuda2hipRename["cudaErrorUnknownSymbol"] = "hipErrorUnknownSymbol"; + cuda2hipRename["cudaErrorOutOfResources"] = "hipErrorOutOfResources"; + cuda2hipRename["cudaErrorInvalidValue"] = "hipErrorInvalidValue"; + cuda2hipRename["cudaErrorInvalidResourceHandle"] = "hipErrorInvalidResourceHandle"; + cuda2hipRename["cudaErrorInvalidDevice"] = "hipErrorInvalidDevice"; + cuda2hipRename["cudaErrorNoDevice"] = "hipErrorNoDevice"; + cuda2hipRename["cudaErrorNotReady"] = "hipErrorNotReady"; + cuda2hipRename["cudaErrorUnknown"] = "hipErrorUnknown"; + + // error APIs: + cuda2hipRename["cudaGetLastError"] = "hipGetLastError"; + cuda2hipRename["cudaPeekAtLastError"] = "hipPeekAtLastError"; + cuda2hipRename["cudaGetErrorName"] = "hipGetErrorName"; + cuda2hipRename["cudaGetErrorString"] = "hipGetErrorString"; + + // Memcpy + cuda2hipRename["cudaMemcpy"] = "hipMemcpy"; + cuda2hipRename["cudaMemcpyHostToHost"] = "hipMemcpyHostToHost"; + cuda2hipRename["cudaMemcpyHostToDevice"] = "hipMemcpyHostToDevice"; + cuda2hipRename["cudaMemcpyDeviceToHost"] = "hipMemcpyDeviceToHost"; + cuda2hipRename["cudaMemcpyDeviceToDevice"] = "hipMemcpyDeviceToDevice"; + cuda2hipRename["cudaMemcpyDefault"] = "hipMemcpyDefault"; + cuda2hipRename["cudaMemcpyToSymbol"] = "hipMemcpyToSymbol"; + cuda2hipRename["cudaMemset"] = "hipMemset"; + cuda2hipRename["cudaMemsetAsync"] = "hipMemsetAsync"; + cuda2hipRename["cudaMemcpyAsync"] = "hipMemcpyAsync"; + cuda2hipRename["cudaMemGetInfo"] = "hipMemGetInfo"; + cuda2hipRename["cudaMemcpyKind"] = "hipMemcpyKind"; + + // Memory management : + cuda2hipRename["cudaMalloc"] = "hipMalloc"; + cuda2hipRename["cudaMallocHost"] = "hipMallocHost"; + cuda2hipRename["cudaFree"] = "hipFree"; + cuda2hipRename["cudaFreeHost"] = "hipFreeHost"; + + // Coordinate Indexing and Dimensions: + cuda2hipRename["threadIdx.x"] = "hipThreadIdx_x"; + cuda2hipRename["threadIdx.y"] = "hipThreadIdx_y"; + cuda2hipRename["threadIdx.z"] = "hipThreadIdx_z"; + + cuda2hipRename["blockIdx.x"] = "hipBlockIdx_x"; + cuda2hipRename["blockIdx.y"] = "hipBlockIdx_y"; + cuda2hipRename["blockIdx.z"] = "hipBlockIdx_z"; + + cuda2hipRename["blockDim.x"] = "hipBlockDim_x"; + cuda2hipRename["blockDim.y"] = "hipBlockDim_y"; + cuda2hipRename["blockDim.z"] = "hipBlockDim_z"; + + cuda2hipRename["gridDim.x"] = "hipGridDim_x"; + cuda2hipRename["gridDim.y"] = "hipGridDim_y"; + cuda2hipRename["gridDim.z"] = "hipGridDim_z"; + + cuda2hipRename["blockIdx.x"] = "hipBlockIdx_x"; + cuda2hipRename["blockIdx.y"] = "hipBlockIdx_y"; + cuda2hipRename["blockIdx.z"] = "hipBlockIdx_z"; + + cuda2hipRename["blockDim.x"] = "hipBlockDim_x"; + cuda2hipRename["blockDim.y"] = "hipBlockDim_y"; + cuda2hipRename["blockDim.z"] = "hipBlockDim_z"; + + cuda2hipRename["gridDim.x"] = "hipGridDim_x"; + cuda2hipRename["gridDim.y"] = "hipGridDim_y"; + cuda2hipRename["gridDim.z"] = "hipGridDim_z"; + + + cuda2hipRename["warpSize"] = "hipWarpSize"; + + // Events + cuda2hipRename["cudaEvent_t"] = "hipEvent_t"; + cuda2hipRename["cudaEventCreate"] = "hipEventCreate"; + cuda2hipRename["cudaEventCreateWithFlags"] = "hipEventCreateWithFlags"; + cuda2hipRename["cudaEventDestroy"] = "hipEventDestroy"; + cuda2hipRename["cudaEventRecord"] = "hipEventRecord"; + cuda2hipRename["cudaEventElapsedTime"] = "hipEventElapsedTime"; + cuda2hipRename["cudaEventSynchronize"] = "hipEventSynchronize"; + + // Streams + cuda2hipRename["cudaStream_t"] = "hipStream_t"; + cuda2hipRename["cudaStreamCreate"] = "hipStreamCreate"; + cuda2hipRename["cudaStreamCreateWithFlags"] = "hipStreamCreateWithFlags"; + cuda2hipRename["cudaStreamDestroy"] = "hipStreamDestroy"; + cuda2hipRename["cudaStreamWaitEvent"] = "hipStreamWaitEven"; + cuda2hipRename["cudaStreamSynchronize"] = "hipStreamSynchronize"; + cuda2hipRename["cudaStreamDefault"] = "hipStreamDefault"; + cuda2hipRename["cudaStreamNonBlocking"] = "hipStreamNonBlocking"; + + // Other synchronization + cuda2hipRename["cudaDeviceSynchronize"] = "hipDeviceSynchronize"; + cuda2hipRename["cudaThreadSynchronize"] = "hipDeviceSynchronize"; // translate deprecated cudaThreadSynchronize + cuda2hipRename["cudaDeviceReset"] = "hipDeviceReset"; + cuda2hipRename["cudaThreadExit"] = "hipDeviceReset"; // translate deprecated cudaThreadExit + cuda2hipRename["cudaSetDevice"] = "hipSetDevice"; + cuda2hipRename["cudaGetDevice"] = "hipGetDevice"; + + // Device + cuda2hipRename["cudaDeviceProp"] = "hipDeviceProp_t"; + cuda2hipRename["cudaGetDeviceProperties"] = "hipDeviceGetProperties"; + + // Cache config + cuda2hipRename["cudaDeviceSetCacheConfig"] = "hipDeviceSetCacheConfig"; + cuda2hipRename["cudaThreadSetCacheConfig"] = "hipDeviceSetCacheConfig"; // translate deprecated + cuda2hipRename["cudaDeviceGetCacheConfig"] = "hipDeviceGetCacheConfig"; + cuda2hipRename["cudaThreadGetCacheConfig"] = "hipDeviceGetCacheConfig"; // translate deprecated + cuda2hipRename["cudaFuncCache"] = "hipFuncCache"; + cuda2hipRename["cudaFuncCachePreferNone"] = "hipFuncCachePreferNone"; + cuda2hipRename["cudaFuncCachePreferShared"] = "hipFuncCachePreferShared"; + cuda2hipRename["cudaFuncCachePreferL1"] = "hipFuncCachePreferL1"; + cuda2hipRename["cudaFuncCachePreferEqual"] = "hipFuncCachePreferEqual"; + // function + cuda2hipRename["cudaFuncSetCacheConfig"] = "hipFuncSetCacheConfig"; + + cuda2hipRename["cudaDriverGetVersion"] = "hipDriverGetVersion"; + + // Peer2Peer + cuda2hipRename["cudaDeviceCanAccessPeer"] = "hipDeviceCanAccessPeer"; + cuda2hipRename["cudaDeviceDisablePeerAccess"] = "hipDeviceDisablePeerAccess"; + cuda2hipRename["cudaDeviceEnablePeerAccess"] = "hipDeviceEnablePeerAccess"; + cuda2hipRename["cudaMemcpyPeerAsync"] = "hipMemcpyPeerAsync"; + cuda2hipRename["cudaMemcpyPeer"] = "hipMemcpyPeer"; + + // Shared mem: + cuda2hipRename["cudaDeviceSetSharedMemConfig"] = "hipDeviceSetSharedMemConfig"; + cuda2hipRename["cudaThreadSetSharedMemConfig"] = "hipDeviceSetSharedMemConfig"; // translate deprecated + cuda2hipRename["cudaDeviceGetSharedMemConfig"] = "hipDeviceGetSharedMemConfig"; + cuda2hipRename["cudaThreadGetSharedMemConfig"] = "hipDeviceGetSharedMemConfig"; // translate deprecated + cuda2hipRename["cudaSharedMemConfig"] = "hipSharedMemConfig"; + cuda2hipRename["cudaSharedMemBankSizeDefault"] = "hipSharedMemBankSizeDefault"; + cuda2hipRename["cudaSharedMemBankSizeFourByte"] = "hipSharedMemBankSizeFourByte"; + cuda2hipRename["cudaSharedMemBankSizeEightByte"] = "hipSharedMemBankSizeEightByte"; + + cuda2hipRename["cudaGetDeviceCount"] = "hipGetDeviceCount"; + + // Profiler + //cuda2hipRename["cudaProfilerInitialize"] = "hipProfilerInitialize"; // see if these are called anywhere. + cuda2hipRename["cudaProfilerStart"] = "hipProfilerStart"; + cuda2hipRename["cudaProfilerStop"] = "hipProfilerStop"; + + cuda2hipRename["cudaChannelFormatDesc"] = "hipChannelFormatDesc"; + cuda2hipRename["cudaFilterModePoint"] = "hipFilterModePoint"; + cuda2hipRename["cudaReadModeElementType"] = "hipReadModeElementType"; + + cuda2hipRename["cudaCreateChannelDesc"] = "hipCreateChannelDesc"; + cuda2hipRename["cudaBindTexture"] = "hipBindTexture"; + cuda2hipRename["cudaUnbindTexture"] = "hipUnbindTexture"; + } + DenseMap cuda2hipRename; + }; + + struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { + HipifyPPCallbacks(Replacements * R) + : SeenEnd(false), _sm(nullptr), Replace(R) + { + } + + virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override + { + Preprocessor &PP = CI.getPreprocessor(); + SourceManager & SM = CI.getSourceManager(); + setSourceManager(&SM); + PP.addPPCallbacks(std::unique_ptr(this)); + PP.Retain(); + return true; + } + + virtual void InclusionDirective( + SourceLocation hash_loc, + const Token &include_token, + StringRef file_name, + bool is_angled, + CharSourceRange filename_range, + const FileEntry *file, + StringRef search_path, + StringRef relative_path, + const Module *imported) override + { + if (_sm->isWrittenInMainFile(hash_loc)) { + if (is_angled) { + if (N.cuda2hipRename.count(file_name)) { + std::string repName = N.cuda2hipRename[file_name]; + llvm::errs() << "\nInclude file found: " << file_name << "\n"; + llvm::errs() << "\nSourceLocation:"; filename_range.getBegin().dump(*_sm); + llvm::errs() << "\nWill be replaced with " << repName << "\n"; + SourceLocation sl = filename_range.getBegin(); + SourceLocation sle = filename_range.getEnd(); + const char* B = _sm->getCharacterData(sl); + const char* E = _sm->getCharacterData(sle); + repName = "<" + repName + ">"; + Replacement Rep(*_sm, sl, E - B, repName); + Replace->insert(Rep); + } + } + } + } + + virtual void MacroDefined(const Token &MacroNameTok, + const MacroDirective *MD) override + { + if (_sm->isWrittenInMainFile(MD->getLocation()) && + MD->getKind() == MacroDirective::MD_Define) + { + for (auto T : MD->getMacroInfo()->tokens()) + { + if (T.isAnyIdentifier()) { + StringRef name = T.getIdentifierInfo()->getName(); + if (N.cuda2hipRename.count(name)) { + StringRef repName = N.cuda2hipRename[name]; + llvm::errs() << "\nIdentifier " << name + << " found in definition of macro " << MacroNameTok.getIdentifierInfo()->getName() << "\n"; + llvm::errs() << "\nwill be replaced with: " << repName << "\n"; + SourceLocation sl = T.getLocation(); + llvm::errs() << "\nSourceLocation: ";sl.dump(*_sm); + llvm::errs() << "\n"; + Replacement Rep(*_sm, sl, name.size(), repName); + Replace->insert(Rep); + } + } + } + } + } + + void EndOfMainFile() override + { + + } + + bool SeenEnd; + void setSourceManager(SourceManager * sm) { _sm = sm; } + + private: + + SourceManager * _sm; + Replacements * Replace; + struct hipName N; + }; + +class Cuda2HipCallback : public MatchFinder::MatchCallback { + public: + Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {} + + void run(const MatchFinder::MatchResult &Result) override { + + SourceManager * SM = Result.SourceManager; + + if (const CallExpr * call = Result.Nodes.getNodeAs("cudaCall")) + { + const FunctionDecl * funcDcl = call->getDirectCallee(); + std::string name = funcDcl->getDeclName().getAsString(); + if (N.cuda2hipRename.count(name)) { + std::string repName = N.cuda2hipRename[name]; + SourceLocation sl = call->getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); + Replace->insert(Rep); + } + } + + if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs("cudaLaunchKernel")) + { + LangOptions DefaultLangOptions; + + const FunctionDecl * kernelDecl = launchKernel->getDirectCallee(); + + const ParmVarDecl * pvdFirst = kernelDecl->getParamDecl(0); + const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1); + SourceLocation kernelArgListStart(pvdFirst->getLocStart()); + SourceLocation kernelArgListEnd(pvdLast->getLocEnd()); + SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); + size_t replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); + std::string outs(SM->getCharacterData(kernelArgListStart), replacementLength); + llvm::outs() << "initial paramlist: " << outs.c_str() << "\n"; + outs = "hipLaunchParm lp, " + outs; + llvm::outs() << "new paramlist: " << outs.c_str() << "\n"; + Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, outs); + Replace->insert(Rep0); + + + std::string name = kernelDecl->getDeclName().getAsString(); + std::string repName = "hipLaunchKernel(HIP_KERNEL_NAME(" + name + "), "; + + + const CallExpr * config = launchKernel->getConfig(); + llvm::outs() << "\nKernel config arguments:\n"; + for (unsigned argno = 0; argno < config->getNumArgs(); argno++) + { + const Expr * arg = config->getArg(argno); + if (!isa(arg)) { + std::string typeCtor = ""; + const ParmVarDecl * pvd = config->getDirectCallee()->getParamDecl(argno); + + SourceLocation sl(arg->getLocStart()); + SourceLocation el(arg->getLocEnd()); + SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); + std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl)); + llvm::outs() << "args[ " << argno << "]" << outs.c_str() << " <" << pvd->getType().getAsString() << ">\n"; + if (pvd->getType().getAsString().compare("dim3") == 0) + repName += " dim3(" + outs + "),"; + else + repName += " " + outs + ","; + } else + repName += " 0,"; + } + + for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++) + { + const Expr * arg = launchKernel->getArg(argno); + SourceLocation sl(arg->getLocStart()); + SourceLocation el(arg->getLocEnd()); + SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); + std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl)); + llvm::outs() << outs.c_str() << "\n"; + repName += " " + outs + ","; + } + repName.pop_back(); + repName += ")"; + size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) - + SM->getCharacterData(launchKernel->getLocStart()); + Replacement Rep(*SM, launchKernel->getLocStart(), length, repName); + Replace->insert(Rep); + } + + if (const MemberExpr * threadIdx = Result.Nodes.getNodeAs("cudaBuiltin")) + { + if (const OpaqueValueExpr * refBase = dyn_cast(threadIdx->getBase())) { + if (const DeclRefExpr * declRef = dyn_cast(refBase->getSourceExpr())) { + std::string name = declRef->getDecl()->getNameAsString(); + std::string memberName = threadIdx->getMemberDecl()->getNameAsString(); + size_t pos = memberName.find_first_not_of("__fetch_builtin_"); + memberName = memberName.substr(pos, memberName.length() - pos); + name += "." + memberName; + std::string repName = N.cuda2hipRename[name]; + SourceLocation sl = threadIdx->getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); + Replace->insert(Rep); + } + } + } + + if (const DeclRefExpr * cudaEnumConstant = Result.Nodes.getNodeAs("cudaEnumConstant")) + { + std::string name = cudaEnumConstant->getDecl()->getNameAsString(); + std::string repName = N.cuda2hipRename[name]; + SourceLocation sl = cudaEnumConstant->getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); + Replace->insert(Rep); + } + + if (const VarDecl * cudaStructVar = Result.Nodes.getNodeAs("cudaStructVar")) + { + std::string name = + cudaStructVar->getType()->getAsStructureType()->getDecl()->getNameAsString(); + std::string repName = N.cuda2hipRename[name]; + SourceLocation sl = cudaStructVar->getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); + Replace->insert(Rep); + } + } + + private: + Replacements *Replace; + struct hipName N; +}; + +} // end anonymous namespace + +// Set up the command line options +static cl::extrahelp CommonHelp(CommonOptionsParser::HelpMessage); +static cl::OptionCategory ToolTemplateCategory("CUDA to HIP source translator options"); + +int main(int argc, const char **argv) { + + llvm::sys::PrintStackTraceOnErrorSignal(); + + int Result; + + CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory); + std::vector savedSources; + for (auto I : OptionsParser.getSourcePathList()) + { + size_t pos = I.find(".cu"); + if (pos != std::string::npos) + { + std::string dst = I.substr(0, pos) + ".hip.cu"; + std::ifstream source(I, std::ios::binary); + std::ofstream dest(dst, std::ios::binary); + dest << source.rdbuf(); + source.close(); + dest.close(); + savedSources.push_back(dst); + } + } + + RefactoringTool Tool(OptionsParser.getCompilations(), savedSources); + ast_matchers::MatchFinder Finder; + Cuda2HipCallback Callback(&Tool.getReplacements()); + HipifyPPCallbacks PPCallbacks(&Tool.getReplacements()); + Finder.addMatcher(callExpr(isExpansionInMainFile(), callee(functionDecl(matchesName("cuda.*")))).bind("cudaCall"), &Callback); + Finder.addMatcher(cudaKernelCallExpr().bind("cudaLaunchKernel"), &Callback); + Finder.addMatcher(memberExpr(isExpansionInMainFile(), hasObjectExpression(hasType(cxxRecordDecl(matchesName("__cuda_builtin_"))))).bind("cudaBuiltin"), &Callback); + Finder.addMatcher(declRefExpr(isExpansionInMainFile(), to(enumConstantDecl(matchesName("cuda.*")))).bind("cudaEnumConstant"), &Callback); + Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(cxxRecordDecl(matchesName("cuda.*")))).bind("cudaStructVar"), &Callback); + + auto action = newFrontendActionFactory(&Finder, &PPCallbacks); + + std::vector compilationStages; + compilationStages.push_back("--cuda-host-only"); + compilationStages.push_back("--cuda-device-only"); + + for (auto Stage : compilationStages) + { + Tool.appendArgumentsAdjuster(combineAdjusters( + getInsertArgumentAdjuster(Stage, ArgumentInsertPosition::BEGIN), + getClangSyntaxOnlyAdjuster())); + + Result = Tool.run(action.get()); + + Tool.clearArgumentsAdjusters(); + } + + LangOptions DefaultLangOptions; + IntrusiveRefCntPtr DiagOpts = new DiagnosticOptions(); + TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts); + DiagnosticsEngine Diagnostics( + IntrusiveRefCntPtr(new DiagnosticIDs()), + &*DiagOpts, &DiagnosticPrinter, false); + SourceManager Sources(Diagnostics, Tool.getFiles()); + + llvm::outs() << "Replacements collected by the tool:\n"; + for (auto &r : Tool.getReplacements()) { + llvm::outs() << r.toString() << "\n"; + } + + Rewriter Rewrite(Sources, DefaultLangOptions); + + if (!Tool.applyAllReplacements(Rewrite)) { + llvm::errs() << "Skipped some replacements.\n"; + } + + Result = Rewrite.overwriteChangedFiles(); + + + for (auto I : savedSources) + { + size_t pos = I.find(".cu"); + if (pos != std::string::npos) + { + rename(I.c_str(), I.substr(0, pos).c_str()); + } + } + return Result; +}