AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[HIPIFY][#1409] Fix for kernel launch macro expansion

Browse Source 
+ Add a corresponding test kernel_launch_01.cu
+ Add isBefore() check to avoid crash on Replacement with negative length

TODO:
+ Compatibility with former LLVM versions
+ More complicated kernel launch tests
This commit is contained in:
Evgeny Mankov
2019-11-05 14:00:13 +03:00
parent c48fdefee8
commit 976f8e8bf6
2 changed files with 64 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipify-clang/src/HipifyAction.cpp
+18 -10
View File
@@ -198,7 +198,7 @@ clang::SourceRange getWriteRange(clang::SourceManager &SM, const clang::SourceRa
// If the range is contained within a macro, update the macro definition.
// Otherwise, use the file location and hope for the best.
if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) {
return {SM.getFileLoc(begin), SM.getFileLoc(end)};
return {SM.getExpansionLoc(begin), SM.getExpansionLoc(end)};
}
return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)};
}
@@ -390,16 +390,24 @@ bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result)
OS << readSourceText(*SM, {argStart, argEnd});
}
OS << ")";
clang::SourceRange replacementRange = getWriteRange(*SM, {llcompat::getBeginLoc(launchKernel), llcompat::getEndLoc(launchKernel)});
clang::SourceLocation launchStart = replacementRange.getBegin();
clang::SourceLocation launchKernelExprLocBeg = launchKernel->getExprLoc();
clang::SourceLocation launchKernelExprLocEnd = launchKernelExprLocBeg.isMacroID() ? SM->getExpansionRange(launchKernelExprLocBeg).getEnd() : llcompat::getEndLoc(launchKernel);
clang::SourceLocation launchKernelEnd = llcompat::getEndLoc(launchKernel);
clang::BeforeThanCompare<clang::SourceLocation> isBefore(*SM);
launchKernelExprLocEnd = isBefore(launchKernelEnd, launchKernelExprLocEnd) ? launchKernelExprLocEnd : launchKernelEnd;
clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernelExprLocBeg, launchKernelExprLocEnd});
clang::SourceLocation launchBeg = replacementRange.getBegin();
clang::SourceLocation launchEnd = replacementRange.getEnd();
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart);
ct::Replacement Rep(*SM, launchStart, length, OS.str());
clang::FullSourceLoc fullSL(launchStart, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str());
return true;
if (isBefore(launchBeg, launchEnd)) {
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchBeg);
ct::Replacement Rep(*SM, launchBeg, length, OS.str());
clang::FullSourceLoc fullSL(launchBeg, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str());
return true;
}
return false;
}
bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) {