From 403f5c71bd7ec1ff908ea53f8bc3ed977aa3eb53 Mon Sep 17 00:00:00 2001 From: dfukalov Date: Thu, 24 Mar 2016 14:54:14 +0300 Subject: [PATCH] source reformatted to LLVM style, minor cleanups --- hipamd/src/Cuda2Hip.cpp | 905 +++++++++++++++++++++------------------- 1 file changed, 479 insertions(+), 426 deletions(-) diff --git a/hipamd/src/Cuda2Hip.cpp b/hipamd/src/Cuda2Hip.cpp index d80e685c23..db8527835f 100644 --- a/hipamd/src/Cuda2Hip.cpp +++ b/hipamd/src/Cuda2Hip.cpp @@ -24,28 +24,28 @@ THE SOFTWARE. * * This file is compiled and linked into clang based hipify tool. */ -#include "clang/ASTMatchers/ASTMatchers.h" #include "clang/ASTMatchers/ASTMatchFinder.h" +#include "clang/ASTMatchers/ASTMatchers.h" #include "clang/Basic/SourceManager.h" +#include "clang/Frontend/CompilerInstance.h" #include "clang/Frontend/FrontendActions.h" +#include "clang/Frontend/TextDiagnosticPrinter.h" #include "clang/Lex/Lexer.h" +#include "clang/Lex/MacroArgs.h" +#include "clang/Lex/MacroInfo.h" +#include "clang/Lex/PPCallbacks.h" +#include "clang/Lex/Preprocessor.h" +#include "clang/Rewrite/Core/Rewriter.h" #include "clang/Tooling/CommonOptionsParser.h" #include "clang/Tooling/Refactoring.h" #include "clang/Tooling/Tooling.h" #include "llvm/Support/CommandLine.h" +#include "llvm/Support/Debug.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 "clang/Lex/MacroArgs.h" -#include #include +#include using namespace clang; using namespace clang::ast_matchers; @@ -55,428 +55,436 @@ using namespace llvm; #define DEBUG_TYPE "cuda2hip" namespace { - struct hipName { +struct hipName { hipName() { - // defines - cuda2hipRename["__CUDACC__"] = "__HIPCC__"; + // defines + cuda2hipRename["__CUDACC__"] = "__HIPCC__"; - // includes - cuda2hipRename["cuda_runtime.h"] = "hip_runtime.h"; - cuda2hipRename["cuda_runtime_api.h"] = "hip_runtime_api.h"; + // 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"; - cuda2hipRename["cudaSuccess"] = "hipSuccess"; + // Error codes and return types: + cuda2hipRename["cudaError_t"] = "hipError_t"; + cuda2hipRename["cudaError"] = "hipError"; + 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"; + 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"; + // 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"; + // 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"] = "hipHostAlloc"; - cuda2hipRename["cudaFree"] = "hipFree"; - cuda2hipRename["cudaFreeHost"] = "hipHostFree"; + // Memory management : + cuda2hipRename["cudaMalloc"] = "hipMalloc"; + cuda2hipRename["cudaMallocHost"] = "hipHostAlloc"; + cuda2hipRename["cudaFree"] = "hipFree"; + cuda2hipRename["cudaFreeHost"] = "hipHostFree"; - // Coordinate Indexing and Dimensions: - cuda2hipRename["threadIdx.x"] = "hipThreadIdx_x"; - cuda2hipRename["threadIdx.y"] = "hipThreadIdx_y"; - cuda2hipRename["threadIdx.z"] = "hipThreadIdx_z"; + // 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["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["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["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["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["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["gridDim.x"] = "hipGridDim_x"; + cuda2hipRename["gridDim.y"] = "hipGridDim_y"; + cuda2hipRename["gridDim.z"] = "hipGridDim_z"; + cuda2hipRename["warpSize"] = "hipWarpSize"; - 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"; - // 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"; - // 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"; - // 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"; - // 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"; - // 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"; + cuda2hipRename["cudaRuntimeGetVersion"] = "hipRuntimeGetVersion"; - cuda2hipRename["cudaDriverGetVersion"] = "hipDriverGetVersion"; - cuda2hipRename["cudaRuntimeGetVersion"] = "hipRuntimeGetVersion"; + // Peer2Peer + cuda2hipRename["cudaDeviceCanAccessPeer"] = "hipDeviceCanAccessPeer"; + cuda2hipRename["cudaDeviceDisablePeerAccess"] = + "hipDeviceDisablePeerAccess"; + cuda2hipRename["cudaDeviceEnablePeerAccess"] = "hipDeviceEnablePeerAccess"; + cuda2hipRename["cudaMemcpyPeerAsync"] = "hipMemcpyPeerAsync"; + cuda2hipRename["cudaMemcpyPeer"] = "hipMemcpyPeer"; - // 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"; - // 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"; - cuda2hipRename["cudaGetDeviceCount"] = "hipGetDeviceCount"; + // Profiler + // cuda2hipRename["cudaProfilerInitialize"] = "hipProfilerInitialize"; // + // see if these are called anywhere. + cuda2hipRename["cudaProfilerStart"] = "hipProfilerStart"; + cuda2hipRename["cudaProfilerStop"] = "hipProfilerStop"; - // 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["cudaChannelFormatDesc"] = "hipChannelFormatDesc"; - cuda2hipRename["cudaFilterModePoint"] = "hipFilterModePoint"; - cuda2hipRename["cudaReadModeElementType"] = "hipReadModeElementType"; + cuda2hipRename["cudaCreateChannelDesc"] = "hipCreateChannelDesc"; + cuda2hipRename["cudaBindTexture"] = "hipBindTexture"; + cuda2hipRename["cudaUnbindTexture"] = "hipUnbindTexture"; + } + DenseMap cuda2hipRename; +}; - cuda2hipRename["cudaCreateChannelDesc"] = "hipCreateChannelDesc"; - cuda2hipRename["cudaBindTexture"] = "hipBindTexture"; - cuda2hipRename["cudaUnbindTexture"] = "hipUnbindTexture"; +StringRef unquoteStr(StringRef s) { + if (s.size() > 1 && s.front() == '"' && s.back() == '"') + return s.substr(1, s.size() - 2); + return s; +} + +void processString(StringRef s, struct hipName &map, Replacements *Replace, + SourceManager &SM, SourceLocation start) { + size_t begin = 0; + while ((begin = s.find("cuda", begin)) != StringRef::npos) { + const size_t end = s.find_first_of(" ", begin + 4); + StringRef name = s.slice(begin, end); + StringRef repName = map.cuda2hipRename[name]; + if (!repName.empty()) { + SourceLocation sl = start.getLocWithOffset(begin + 1); + Replacement Rep(SM, sl, name.size(), repName); + Replace->insert(Rep); } - DenseMap cuda2hipRename; - }; + if (end == StringRef::npos) + break; + begin = end + 1; + } +} - StringRef unquoteStr(StringRef s) { - if (s.size() > 1 && s.front() == '"' && s.back() == '"') - return s.substr(1, s.size()-2); - return s; +struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { + HipifyPPCallbacks(Replacements *R) + : SeenEnd(false), _sm(nullptr), _pp(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(); + setPreprocessor(&PP); + return true; } - void processString(StringRef s, struct hipName & map, - Replacements * Replace, SourceManager & SM, SourceLocation start) - { - size_t begin = 0; - while ((begin = s.find("cuda", begin)) != StringRef::npos) { - const size_t end = s.find_first_of(" ", begin + 4); - StringRef name = s.slice(begin, end); - StringRef repName = map.cuda2hipRename[name]; - if (!repName.empty()) { - SourceLocation sl = start.getLocWithOffset(begin + 1); - Replacement Rep(SM, sl, name.size(), repName); - Replace->insert(Rep); + 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 clang::Module *imported) override { + if (_sm->isWrittenInMainFile(hash_loc)) { + if (is_angled) { + if (N.cuda2hipRename.count(file_name)) { + StringRef repName = N.cuda2hipRename[file_name]; + DEBUG(dbgs() << "Include file found: " << file_name << "\n" + << "SourceLocation:" + << filename_range.getBegin().printToString(*_sm) << "\n" + << "Will 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); + SmallString<128> tmpData; + Replacement Rep(*_sm, sl, E - B, + Twine("<" + repName + ">").toStringRef(tmpData)); + Replace->insert(Rep); + } } - if (end == StringRef::npos) break; - begin = end + 1; } } - - struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { - HipifyPPCallbacks(Replacements * R) - : SeenEnd(false), _sm(nullptr), _pp(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(); - setPreprocessor(&PP); - 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 clang::Module *imported) override - { - if (_sm->isWrittenInMainFile(hash_loc)) { - if (is_angled) { - if (N.cuda2hipRename.count(file_name)) { - StringRef repName = N.cuda2hipRename[file_name]; - DEBUG(dbgs() << "Include file found: " << file_name << "\n"); - DEBUG(dbgs() << "SourceLocation:" - << filename_range.getBegin().printToString(*_sm) << "\n"); - DEBUG(dbgs() << "Will 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); - SmallString<128> tmpData; - Replacement Rep(*_sm, sl, E - B, Twine("<" + repName + ">").toStringRef(tmpData)); + 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]; + SourceLocation sl = T.getLocation(); + DEBUG(dbgs() << "Identifier " << name + << " found in definition of macro " + << MacroNameTok.getIdentifierInfo()->getName() << "\n" + << "will be replaced with: " << repName << "\n" + << "SourceLocation: " << sl.printToString(*_sm) + << "\n"); + Replacement Rep(*_sm, sl, name.size(), 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(); + virtual void MacroExpands(const Token &MacroNameTok, + const MacroDefinition &MD, SourceRange Range, + const MacroArgs *Args) override { + if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { + for (unsigned int i = 0; Args && i < MD.getMacroInfo()->getNumArgs(); + i++) { + StringRef macroName = MacroNameTok.getIdentifierInfo()->getName(); + std::vector toks; + // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument' + // to workaround the 'const' MacroArgs passed into this hook. + const Token *start = Args->getUnexpArgument(i); + size_t len = Args->getArgLength(start) + 1; +#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) + _pp->EnterTokenStream(ArrayRef(start, len), false); +#else + _pp->EnterTokenStream(start, len, false, false); +#endif + do { + toks.push_back(Token()); + Token &tk = toks.back(); + _pp->Lex(tk); + } while (toks.back().isNot(tok::eof)); + _pp->RemoveTopOfLexerStack(); + // end of stolen code + for (auto tok : toks) { + if (tok.isAnyIdentifier()) { + StringRef name = tok.getIdentifierInfo()->getName(); if (N.cuda2hipRename.count(name)) { StringRef repName = N.cuda2hipRename[name]; DEBUG(dbgs() << "Identifier " << name - << " found in definition of macro " - << MacroNameTok.getIdentifierInfo()->getName() << "\n"); - DEBUG(dbgs() << "will be replaced with: " << repName << "\n"); - SourceLocation sl = T.getLocation(); - DEBUG(dbgs() << "SourceLocation: " << sl.printToString(*_sm) << "\n"); + << " found as an actual argument in expansion of macro " + << macroName << "\n" + << "will be replaced with: " << repName << "\n"); + SourceLocation sl = tok.getLocation(); Replacement Rep(*_sm, sl, name.size(), repName); Replace->insert(Rep); } } - } - } - } - - virtual void MacroExpands(const Token &MacroNameTok, - const MacroDefinition &MD, SourceRange Range, - const MacroArgs *Args) override - { - if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) - { - for (unsigned int i = 0; Args && i < MD.getMacroInfo()->getNumArgs(); i++) - { - StringRef macroName = MacroNameTok.getIdentifierInfo()->getName(); - std::vector toks; - // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument' - // to workaround the 'const' MacroArgs passed into this hook. - const Token * start = Args->getUnexpArgument(i); - size_t len = Args->getArgLength(start) + 1; -#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) - _pp->EnterTokenStream(ArrayRef(start,len), false); -#else - _pp->EnterTokenStream(start, len, false, false); -#endif - do { - toks.push_back(Token()); - Token & tk = toks.back(); - _pp->Lex(tk); - } while (toks.back().isNot(tok::eof)); - _pp->RemoveTopOfLexerStack(); - // end of stolen code - for (auto tok : toks) { - if (tok.isAnyIdentifier()) - { - StringRef name = tok.getIdentifierInfo()->getName(); - if (N.cuda2hipRename.count(name)) { - StringRef repName = N.cuda2hipRename[name]; - DEBUG(dbgs() << "Identifier " << name - << " found as an actual argument in expansion of macro " - << macroName << "\n"); - DEBUG(dbgs() << "will be replaced with: " << repName << "\n"); - SourceLocation sl = tok.getLocation(); - Replacement Rep(*_sm, sl, name.size(), repName); - Replace->insert(Rep); - } - } - if (tok.is(tok::string_literal)) - { - StringRef s(tok.getLiteralData(), tok.getLength()); - processString(unquoteStr(s), N, Replace, *_sm, tok.getLocation()); - } + if (tok.is(tok::string_literal)) { + StringRef s(tok.getLiteralData(), tok.getLength()); + processString(unquoteStr(s), N, Replace, *_sm, tok.getLocation()); } } } } + } - void EndOfMainFile() override - { + void EndOfMainFile() override {} - } + bool SeenEnd; + void setSourceManager(SourceManager *sm) { _sm = sm; } + void setPreprocessor(Preprocessor *pp) { _pp = pp; } - bool SeenEnd; - void setSourceManager(SourceManager * sm) { _sm = sm; } - void setPreprocessor (Preprocessor * pp) { _pp = pp; } +private: + SourceManager *_sm; + Preprocessor *_pp; - private: - - SourceManager * _sm; - Preprocessor * _pp; - - Replacements * Replace; - struct hipName N; - }; + Replacements *Replace; + struct hipName N; +}; class Cuda2HipCallback : public MatchFinder::MatchCallback { - public: - Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent) - : Replace(Replace), owner(parent) {} +public: + Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent) + : Replace(Replace), owner(parent) {} - void convertKernelDecl(const FunctionDecl * kernelDecl, const MatchFinder::MatchResult &Result) - { - SourceManager * SM = Result.SourceManager; - LangOptions DefaultLangOptions; + void convertKernelDecl(const FunctionDecl *kernelDecl, + const MatchFinder::MatchResult &Result) { + SourceManager *SM = Result.SourceManager; + LangOptions DefaultLangOptions; - SmallString<40> XStr; - raw_svector_ostream OS(XStr); - StringRef initialParamList; - OS << "hipLaunchParm lp"; - size_t replacementLength = OS.str().size(); - SourceLocation sl = kernelDecl->getNameInfo().getEndLoc(); - SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(sl, clang::tok::l_paren, *SM, DefaultLangOptions, true); - DEBUG(dbgs() << kernelArgListStart.printToString(*SM)); - if (kernelDecl->getNumParams() > 0) { - 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); - initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), replacementLength); - OS << ", " << initialParamList; - } - DEBUG(dbgs() << "initial paramlist: " << initialParamList << "\n"); - DEBUG(dbgs() << "new paramlist: " << OS.str() << "\n"); - Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, OS.str()); - Replace->insert(Rep0); + SmallString<40> XStr; + raw_svector_ostream OS(XStr); + StringRef initialParamList; + OS << "hipLaunchParm lp"; + size_t replacementLength = OS.str().size(); + SourceLocation sl = kernelDecl->getNameInfo().getEndLoc(); + SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken( + sl, clang::tok::l_paren, *SM, DefaultLangOptions, true); + DEBUG(dbgs() << kernelArgListStart.printToString(*SM)); + if (kernelDecl->getNumParams() > 0) { + 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); + initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), + replacementLength); + OS << ", " << initialParamList; + } + DEBUG(dbgs() << "initial paramlist: " << initialParamList << "\n" + << "new paramlist: " << OS.str() << "\n"); + Replacement Rep0(*(Result.SourceManager), kernelArgListStart, + replacementLength, OS.str()); + Replace->insert(Rep0); } void run(const MatchFinder::MatchResult &Result) override { - - SourceManager * SM = Result.SourceManager; + SourceManager *SM = Result.SourceManager; LangOptions DefaultLangOptions; - if (const CallExpr * call = Result.Nodes.getNodeAs("cudaCall")) - { - const FunctionDecl * funcDcl = call->getDirectCallee(); + if (const CallExpr *call = + Result.Nodes.getNodeAs("cudaCall")) { + const FunctionDecl *funcDcl = call->getDirectCallee(); StringRef name = funcDcl->getDeclName().getAsString(); if (N.cuda2hipRename.count(name)) { StringRef repName = N.cuda2hipRename[name]; SourceLocation sl = call->getLocStart(); - Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? - SM->getImmediateSpellingLoc(sl) : sl, name.size(), repName); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) + ? SM->getImmediateSpellingLoc(sl) + : sl, + name.size(), repName); Replace->insert(Rep); } } - if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs("cudaLaunchKernel")) - { + if (const CUDAKernelCallExpr *launchKernel = + Result.Nodes.getNodeAs( + "cudaLaunchKernel")) { SmallString<40> XStr; raw_svector_ostream OS(XStr); StringRef calleeName; - const FunctionDecl * kernelDecl = launchKernel->getDirectCallee(); + const FunctionDecl *kernelDecl = launchKernel->getDirectCallee(); if (kernelDecl) { calleeName = kernelDecl->getName(); convertKernelDecl(kernelDecl, Result); - } - else { - const Expr * e = launchKernel->getCallee(); - if (const UnresolvedLookupExpr * ule = dyn_cast(e)) { + } else { + const Expr *e = launchKernel->getCallee(); + if (const UnresolvedLookupExpr *ule = + dyn_cast(e)) { calleeName = ule->getName().getAsIdentifierInfo()->getName(); - owner->addMatcher(functionTemplateDecl(hasName(calleeName)).bind("unresolvedTemplateName"), this); + owner->addMatcher(functionTemplateDecl(hasName(calleeName)) + .bind("unresolvedTemplateName"), + this); } } - XStr.clear(); - OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << calleeName << "), "; + OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << calleeName << "),"; - const CallExpr * config = launchKernel->getConfig(); + const CallExpr *config = launchKernel->getConfig(); DEBUG(dbgs() << "Kernel config arguments:" << "\n"); - for (unsigned argno = 0; argno < config->getNumArgs(); argno++) - { - const Expr * arg = config->getArg(argno); + for (unsigned argno = 0; argno < config->getNumArgs(); argno++) { + const Expr *arg = config->getArg(argno); if (!isa(arg)) { - const ParmVarDecl * pvd = config->getDirectCallee()->getParamDecl(argno); + const ParmVarDecl *pvd = + config->getDirectCallee()->getParamDecl(argno); SourceLocation sl(arg->getLocStart()); SourceLocation el(arg->getLocEnd()); - SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); - StringRef outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl)); + SourceLocation stop = + clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); + StringRef outs(SM->getCharacterData(sl), + SM->getCharacterData(stop) - SM->getCharacterData(sl)); DEBUG(dbgs() << "args[ " << argno << "]" << outs << " <" - << pvd->getType().getAsString() << ">" << "\n"); + << pvd->getType().getAsString() << ">" << "\n"); if (pvd->getType().getAsString().compare("dim3") == 0) OS << " dim3(" << outs << "),"; else @@ -485,50 +493,56 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { OS << " 0,"; } - for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++) - { - const Expr * arg = launchKernel->getArg(argno); + 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)); + SourceLocation stop = + clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); + std::string outs(SM->getCharacterData(sl), + SM->getCharacterData(stop) - SM->getCharacterData(sl)); DEBUG(dbgs() << outs << "\n"); OS << " " << outs << ","; } XStr.pop_back(); OS << ")"; - size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) - - SM->getCharacterData(launchKernel->getLocStart()); + size_t length = + SM->getCharacterData(clang::Lexer::getLocForEndOfToken( + launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) - + SM->getCharacterData(launchKernel->getLocStart()); Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str()); Replace->insert(Rep); - } + } - if (const FunctionTemplateDecl * templateDecl = Result.Nodes.getNodeAs("unresolvedTemplateName")) - { - FunctionDecl * kernelDecl = templateDecl->getTemplatedDecl(); + if (const FunctionTemplateDecl *templateDecl = + Result.Nodes.getNodeAs( + "unresolvedTemplateName")) { + FunctionDecl *kernelDecl = templateDecl->getTemplatedDecl(); convertKernelDecl(kernelDecl, Result); } - if (const MemberExpr * threadIdx = Result.Nodes.getNodeAs("cudaBuiltin")) - { - if (const OpaqueValueExpr * refBase = dyn_cast(threadIdx->getBase())) { - if (const DeclRefExpr * declRef = dyn_cast(refBase->getSourceExpr())) { - StringRef name = declRef->getDecl()->getName(); - StringRef memberName = threadIdx->getMemberDecl()->getName(); - size_t pos = memberName.find_first_not_of("__fetch_builtin_"); - memberName = memberName.slice(pos, memberName.size()); - SmallString<128> tmpData; - name = Twine(name+"."+memberName).toStringRef(tmpData); - StringRef repName = N.cuda2hipRename[name]; - SourceLocation sl = threadIdx->getLocStart(); - Replacement Rep(*SM, sl, name.size(), 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())) { + StringRef name = declRef->getDecl()->getName(); + StringRef memberName = threadIdx->getMemberDecl()->getName(); + size_t pos = memberName.find_first_not_of("__fetch_builtin_"); + memberName = memberName.slice(pos, memberName.size()); + SmallString<128> tmpData; + name = Twine(name + "." + memberName).toStringRef(tmpData); + StringRef repName = N.cuda2hipRename[name]; + SourceLocation sl = threadIdx->getLocStart(); + Replacement Rep(*SM, sl, name.size(), repName); + Replace->insert(Rep); } - } - } + } + } - if (const DeclRefExpr * cudaEnumConstantRef = Result.Nodes.getNodeAs("cudaEnumConstantRef")) - { + if (const DeclRefExpr *cudaEnumConstantRef = + Result.Nodes.getNodeAs("cudaEnumConstantRef")) { StringRef name = cudaEnumConstantRef->getDecl()->getNameAsString(); StringRef repName = N.cuda2hipRename[name]; SourceLocation sl = cudaEnumConstantRef->getLocStart(); @@ -536,18 +550,22 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { Replace->insert(Rep); } - if (const VarDecl * cudaEnumConstantDecl = Result.Nodes.getNodeAs("cudaEnumConstantDecl")) - { - StringRef name = cudaEnumConstantDecl->getType()->getAsTagDecl()->getNameAsString(); + if (const VarDecl *cudaEnumConstantDecl = + Result.Nodes.getNodeAs("cudaEnumConstantDecl")) { + StringRef name = + cudaEnumConstantDecl->getType()->getAsTagDecl()->getNameAsString(); StringRef repName = N.cuda2hipRename[name]; SourceLocation sl = cudaEnumConstantDecl->getLocStart(); Replacement Rep(*SM, sl, name.size(), repName); Replace->insert(Rep); } - if (const VarDecl * cudaStructVar = Result.Nodes.getNodeAs("cudaStructVar")) - { - StringRef name = cudaStructVar->getType()->getAsStructureType()->getDecl()->getNameAsString(); + if (const VarDecl *cudaStructVar = + Result.Nodes.getNodeAs("cudaStructVar")) { + StringRef name = cudaStructVar->getType() + ->getAsStructureType() + ->getDecl() + ->getNameAsString(); StringRef repName = N.cuda2hipRename[name]; TypeLoc TL = cudaStructVar->getTypeSourceInfo()->getTypeLoc(); SourceLocation sl = TL.getUnqualifiedLoc().getLocStart(); @@ -555,9 +573,9 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { Replace->insert(Rep); } - if (const VarDecl * cudaStructVarPtr = Result.Nodes.getNodeAs("cudaStructVarPtr")) - { - const Type * t = cudaStructVarPtr->getType().getTypePtrOrNull(); + if (const VarDecl *cudaStructVarPtr = + Result.Nodes.getNodeAs("cudaStructVarPtr")) { + const Type *t = cudaStructVarPtr->getType().getTypePtrOrNull(); if (t) { StringRef name = t->getPointeeCXXRecordDecl()->getName(); StringRef repName = N.cuda2hipRename[name]; @@ -568,12 +586,11 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { } } - - if (const ParmVarDecl * cudaParamDecl = Result.Nodes.getNodeAs("cudaParamDecl")) - { + if (const ParmVarDecl *cudaParamDecl = + Result.Nodes.getNodeAs("cudaParamDecl")) { QualType QT = cudaParamDecl->getOriginalType().getUnqualifiedType(); StringRef name = QT.getAsString(); - const Type * t = QT.getTypePtr(); + const Type *t = QT.getTypePtr(); if (t->isStructureOrClassType()) { name = t->getAsCXXRecordDecl()->getName(); } @@ -584,14 +601,15 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { Replace->insert(Rep); } - if (const ParmVarDecl * cudaParamDeclPtr = Result.Nodes.getNodeAs("cudaParamDeclPtr")) - { - const Type * pt = cudaParamDeclPtr->getType().getTypePtrOrNull(); + if (const ParmVarDecl *cudaParamDeclPtr = + Result.Nodes.getNodeAs("cudaParamDeclPtr")) { + const Type *pt = cudaParamDeclPtr->getType().getTypePtrOrNull(); if (pt) { QualType QT = pt->getPointeeType(); - const Type * t = QT.getTypePtr(); - StringRef name = t->isStructureOrClassType()? - t->getAsCXXRecordDecl()->getName() : StringRef(QT.getAsString()); + const Type *t = QT.getTypePtr(); + StringRef name = t->isStructureOrClassType() + ? t->getAsCXXRecordDecl()->getName() + : StringRef(QT.getAsString()); StringRef repName = N.cuda2hipRename[name]; TypeLoc TL = cudaParamDeclPtr->getTypeSourceInfo()->getTypeLoc(); SourceLocation sl = TL.getUnqualifiedLoc().getLocStart(); @@ -600,20 +618,20 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { } } - - if (const StringLiteral * stringLiteral = Result.Nodes.getNodeAs("stringLiteral")) - { + if (const StringLiteral *stringLiteral = + Result.Nodes.getNodeAs("stringLiteral")) { if (stringLiteral->getCharByteWidth() == 1) { StringRef s = stringLiteral->getString(); processString(s, N, Replace, *SM, stringLiteral->getLocStart()); } } - if (const UnaryExprOrTypeTraitExpr * expr = Result.Nodes.getNodeAs("cudaStructSizeOf")) - { - TypeSourceInfo * typeInfo = expr->getArgumentTypeInfo(); + if (const UnaryExprOrTypeTraitExpr *expr = + Result.Nodes.getNodeAs( + "cudaStructSizeOf")) { + TypeSourceInfo *typeInfo = expr->getArgumentTypeInfo(); QualType QT = typeInfo->getType().getUnqualifiedType(); - const Type * type = QT.getTypePtr(); + const Type *type = QT.getTypePtr(); StringRef name = type->getAsCXXRecordDecl()->getName(); StringRef repName = N.cuda2hipRename[name]; TypeLoc TL = typeInfo->getTypeLoc(); @@ -623,27 +641,29 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { } } - private: +private: Replacements *Replace; - ast_matchers::MatchFinder * owner; + ast_matchers::MatchFinder *owner; struct hipName N; }; } // end anonymous namespace // Set up the command line options -static cl::OptionCategory ToolTemplateCategory("CUDA to HIP source translator options"); -static cl::extrahelp MoreHelp( " specify the path of source file\n\n" ); +static cl::OptionCategory + ToolTemplateCategory("CUDA to HIP source translator options"); +static cl::extrahelp MoreHelp(" specify the path of source file\n\n"); -static cl::opt -OutputFilename("o", cl::desc("Output filename"), cl::value_desc("filename"), cl::cat(ToolTemplateCategory)); - -//static cl::opt -//Debug("debug", cl::desc("Enable debug output"), cl::Hidden, cl::location(llvm::DebugFlag)); +static cl::opt OutputFilename("o", cl::desc("Output filename"), + cl::value_desc("filename"), + cl::cat(ToolTemplateCategory)); static cl::opt -Inplace("inplace", cl::desc("Modify input file inplace, replacing input with hipified output, save backup in .prehip file. " - "If .prehip file exists, use that as input to hip."), cl::value_desc("inplace"), cl::cat(ToolTemplateCategory)); + Inplace("inplace", + cl::desc("Modify input file inplace, replacing input with hipified " + "output, save backup in .prehip file. " + "If .prehip file exists, use that as input to hip."), + cl::value_desc("inplace"), cl::cat(ToolTemplateCategory)); int main(int argc, const char **argv) { @@ -651,7 +671,8 @@ int main(int argc, const char **argv) { int Result; - CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::Required); + CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, + llvm::cl::Required); std::string dst = OutputFilename; std::vector fileSources = OptionsParser.getSourcePathList(); if (dst.empty()) { @@ -673,7 +694,7 @@ int main(int argc, const char **argv) { } std::ifstream source(fileSources[0], std::ios::binary); - std::ofstream dest(Inplace ? dst+".prehip" : dst, std::ios::binary); + std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary); dest << source.rdbuf(); source.close(); dest.close(); @@ -682,17 +703,49 @@ int main(int argc, const char **argv) { ast_matchers::MatchFinder Finder; Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder); HipifyPPCallbacks PPCallbacks(&Tool.getReplacements()); - Finder.addMatcher(callExpr(isExpansionInMainFile(), callee(functionDecl(matchesName("cuda.*")))).bind("cudaCall"), &Callback); + 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("cudaEnumConstantRef"), &Callback); - Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(enumDecl(matchesName("cuda.*")))).bind("cudaEnumConstantDecl"), &Callback); - Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(cxxRecordDecl(matchesName("cuda.*")))).bind("cudaStructVar"), &Callback); - Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(pointsTo(cxxRecordDecl(matchesName("cuda.*"))))).bind("cudaStructVarPtr"), &Callback); - Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(namedDecl(matchesName("cuda.*")))).bind("cudaParamDecl"), &Callback); - Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(pointsTo(namedDecl(matchesName("cuda.*"))))).bind("cudaParamDeclPtr"), &Callback); - Finder.addMatcher(expr(isExpansionInMainFile(), sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(cxxRecordDecl(matchesName("cuda.*"))))))).bind("cudaStructSizeOf"), &Callback); - Finder.addMatcher(stringLiteral(isExpansionInMainFile()).bind("stringLiteral"), &Callback); + Finder.addMatcher(memberExpr(isExpansionInMainFile(), + hasObjectExpression(hasType(cxxRecordDecl( + matchesName("__cuda_builtin_"))))) + .bind("cudaBuiltin"), + &Callback); + Finder.addMatcher(declRefExpr(isExpansionInMainFile(), + to(enumConstantDecl(matchesName("cuda.*")))) + .bind("cudaEnumConstantRef"), + &Callback); + Finder.addMatcher( + varDecl(isExpansionInMainFile(), hasType(enumDecl(matchesName("cuda.*")))) + .bind("cudaEnumConstantDecl"), + &Callback); + Finder.addMatcher(varDecl(isExpansionInMainFile(), + hasType(cxxRecordDecl(matchesName("cuda.*")))) + .bind("cudaStructVar"), + &Callback); + Finder.addMatcher( + varDecl(isExpansionInMainFile(), + hasType(pointsTo(cxxRecordDecl(matchesName("cuda.*"))))) + .bind("cudaStructVarPtr"), + &Callback); + Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), + hasType(namedDecl(matchesName("cuda.*")))) + .bind("cudaParamDecl"), + &Callback); + Finder.addMatcher( + parmVarDecl(isExpansionInMainFile(), + hasType(pointsTo(namedDecl(matchesName("cuda.*"))))) + .bind("cudaParamDeclPtr"), + &Callback); + Finder.addMatcher(expr(isExpansionInMainFile(), + sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration( + cxxRecordDecl(matchesName("cuda.*"))))))) + .bind("cudaStructSizeOf"), + &Callback); + Finder.addMatcher( + stringLiteral(isExpansionInMainFile()).bind("stringLiteral"), &Callback); auto action = newFrontendActionFactory(&Finder, &PPCallbacks); @@ -700,12 +753,13 @@ int main(int argc, const char **argv) { compilationStages.push_back("--cuda-host-only"); compilationStages.push_back("--cuda-device-only"); - for (auto Stage : compilationStages) - { - Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster(Stage, ArgumentInsertPosition::BEGIN)); + for (auto Stage : compilationStages) { + Tool.appendArgumentsAdjuster( + getInsertArgumentAdjuster(Stage, ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11")); #if defined(HIPIFY_CLANG_RES) - Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); + Tool.appendArgumentsAdjuster( + getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); #endif // defined(HIPIFY_CLANG_HEADERS) Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster()); Result = Tool.run(action.get()); @@ -717,8 +771,8 @@ int main(int argc, const char **argv) { IntrusiveRefCntPtr DiagOpts = new DiagnosticOptions(); TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts); DiagnosticsEngine Diagnostics( - IntrusiveRefCntPtr(new DiagnosticIDs()), - &*DiagOpts, &DiagnosticPrinter, false); + IntrusiveRefCntPtr(new DiagnosticIDs()), &*DiagOpts, + &DiagnosticPrinter, false); SourceManager Sources(Diagnostics, Tool.getFiles()); DEBUG(dbgs() << "Replacements collected by the tool:\n"); @@ -736,8 +790,7 @@ int main(int argc, const char **argv) { if (!Inplace) { size_t pos = dst.rfind(".cu"); - if (pos != std::string::npos) - { + if (pos != std::string::npos) { rename(dst.c_str(), dst.substr(0, pos).c_str()); } }