diff --git a/hipamd/hipify-clang/src/HipifyAction.cpp b/hipamd/hipify-clang/src/HipifyAction.cpp index 4acbb7e78e..01c2a7f092 100644 --- a/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/hipamd/hipify-clang/src/HipifyAction.cpp @@ -42,6 +42,8 @@ const std::string sCudaGetSymbolSize = "cudaGetSymbolSize"; const std::string sCudaGetSymbolAddress = "cudaGetSymbolAddress"; const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol"; const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync"; +const std::string sCudaFuncSetCacheConfig = "cudaFuncSetCacheConfig"; +const std::string sCudaFuncGetAttributes = "cudaFuncGetAttributes"; std::set DeviceSymbolFunctions0 { {sCudaMemcpyToSymbol}, @@ -413,29 +415,31 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match insertReplacement(Rep, fullSL); hipCounter counter = {"HIP_DYNAMIC_SHARED", "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME}; Statistics::current().incrementCounter(counter, refName.str()); + return true; } - return true; + return false; } bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { if (const clang::CallExpr *call = Result.Nodes.getNodeAs("cudaDeviceFuncCall")) { const clang::FunctionDecl *funcDcl = call->getDirectCallee(); if (!funcDcl) { - return true; + return false; } FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false); + return true; } - return true; + return false; } bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { if (const clang::CallExpr * call = Result.Nodes.getNodeAs("cudaSymbolFuncCall")) { if (!call->getNumArgs()) { - return true; + return false; } const clang::FunctionDecl* funcDcl = call->getDirectCallee(); if (!funcDcl) { - return true; + return false; } std::string sName = funcDcl->getDeclName().getAsString(); unsigned int argNum = 0; @@ -444,7 +448,7 @@ bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::Ma } else if (call->getNumArgs() > 1 && DeviceSymbolFunctions1.find(sName) != DeviceSymbolFunctions1.end()) { argNum = 1; } else { - return true; + return false; } clang::SmallString<40> XStr; llvm::raw_svector_ostream OS(XStr); @@ -460,8 +464,46 @@ bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::Ma ct::Replacement Rep(*SM, s, length, OS.str()); clang::FullSourceLoc fullSL(s, *SM); insertReplacement(Rep, fullSL); + return true; } - return true; + return false; +} + +bool HipifyAction::cudaReinterpretCastArgFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) { + if (const clang::CallExpr * call = Result.Nodes.getNodeAs("cudaReinterpretCastArgFuncCall")) { + if (!call->getNumArgs()) { + return false; + } + const clang::FunctionDecl* funcDcl = call->getDirectCallee(); + if (!funcDcl) { + return false; + } + std::string sName = funcDcl->getDeclName().getAsString(); + unsigned int argNum = 0; + if (sCudaFuncSetCacheConfig == sName) { + argNum = 0; + } else if (call->getNumArgs() > 1 && sCudaFuncGetAttributes == sName) { + argNum = 1; + } else { + return false; + } + clang::SmallString<40> XStr; + llvm::raw_svector_ostream OS(XStr); + clang::SourceRange sr = call->getArg(argNum)->getSourceRange(); + clang::SourceManager* SM = Result.SourceManager; + const std::string sCast = "reinterpret_cast"; + OS << sCast << "(" << readSourceText(*SM, sr) << ")"; + clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() }); + clang::SourceLocation s = replacementRange.getBegin(); + clang::SourceLocation e = replacementRange.getEnd(); + clang::LangOptions DefaultLangOptions; + size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(e, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(s); + ct::Replacement Rep(*SM, s, length, OS.str()); + clang::FullSourceLoc fullSL(s, *SM); + insertReplacement(Rep, fullSL); + return true; + } + return false; } void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) { @@ -505,6 +547,20 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi ).bind("cudaSymbolFuncCall"), this ); + Finder->addMatcher( + mat::callExpr( + mat::isExpansionInMainFile(), + mat::callee( + mat::functionDecl( + mat::hasAnyName( + sCudaFuncSetCacheConfig, + sCudaFuncGetAttributes + ) + ) + ) + ).bind("cudaReinterpretCastArgFuncCall"), + this + ); Finder->addMatcher( mat::callExpr( mat::isExpansionInMainFile(), @@ -643,5 +699,6 @@ void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Resu if (cudaLaunchKernel(Result)) return; if (cudaSharedIncompleteArrayVar(Result)) return; if (cudaSymbolFuncCall(Result)) return; + if (cudaReinterpretCastArgFuncCall(Result)) return; if (cudaDeviceFuncCall(Result)) return; } diff --git a/hipamd/hipify-clang/src/HipifyAction.h b/hipamd/hipify-clang/src/HipifyAction.h index 6c04e2c0cc..1c87738854 100644 --- a/hipamd/hipify-clang/src/HipifyAction.h +++ b/hipamd/hipify-clang/src/HipifyAction.h @@ -72,6 +72,8 @@ public: bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result); bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); bool cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); + bool cudaReinterpretCastArgFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result); + // Called by the preprocessor for each include directive during the non-raw lexing pass. void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token &include_token, diff --git a/hipamd/tests/hipify-clang/unit_tests/samples/reinterpret_cast.cu b/hipamd/tests/hipify-clang/unit_tests/samples/reinterpret_cast.cu new file mode 100644 index 0000000000..fe67629ec1 --- /dev/null +++ b/hipamd/tests/hipify-clang/unit_tests/samples/reinterpret_cast.cu @@ -0,0 +1,53 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +/* +Copyright (c) 2015-present 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. +*/ + +#include +// CHECK: #include +#include + +__global__ +void fn(float* px, float* py) { + bool a[42]; + __shared__ double b[69]; + for (auto&& x : b) x = *py++; + for (auto&& x : a) x = *px++ > 0.0; + for (auto&& x : a) if (x)* --py = *--px; +} + +int main() { + // CHECK: hipFuncCache_t cacheConfig; + cudaFuncCache cacheConfig; + void* func; + // CHECK: hipFuncSetCacheConfig(reinterpret_cast(func), cacheConfig); + cudaFuncSetCacheConfig(func, cacheConfig); + // CHECK: hipFuncAttributes attr{}; + cudaFuncAttributes attr{}; + // CHECK: auto r = hipFuncGetAttributes(&attr, reinterpret_cast(&fn)); + auto r = cudaFuncGetAttributes(&attr, &fn); + // CHECK: if (r != hipSuccess || attr.maxThreadsPerBlock == 0) { + if (r != cudaSuccess || attr.maxThreadsPerBlock == 0) { + return 1; + } + return 0; +}