AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Fixed tool crash on kernels with empty parameter list

Przeglądaj źródła
This commit is contained in:
atimofee
2016-02-25 22:54:58 +03:00
rodzic 36833c78f3
commit 63df6f915f
1 zmienionych plików z 99 dodań i 92 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
hipamd/src/Cuda2Hip.cpp
+99 -92
Wyświetl plik
@@ -307,98 +307,105 @@ namespace {
}; };
class Cuda2HipCallback : public MatchFinder::MatchCallback { class Cuda2HipCallback : public MatchFinder::MatchCallback {
public: public:
Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {} Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {}
void run(const MatchFinder::MatchResult &Result) override { void run(const MatchFinder::MatchResult &Result) override {
SourceManager * SM = Result.SourceManager; SourceManager * SM = Result.SourceManager;
if (const CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaCall")) if (const CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaCall"))
{ {
const FunctionDecl * funcDcl = call->getDirectCallee(); const FunctionDecl * funcDcl = call->getDirectCallee();
std::string name = funcDcl->getDeclName().getAsString(); std::string name = funcDcl->getDeclName().getAsString();
if (N.cuda2hipRename.count(name)) { if (N.cuda2hipRename.count(name)) {
std::string repName = N.cuda2hipRename[name]; std::string repName = N.cuda2hipRename[name];
SourceLocation sl = call->getLocStart(); SourceLocation sl = call->getLocStart();
Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ?
SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName);
Replace->insert(Rep); Replace->insert(Rep);
} }
} }
if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs<clang::CUDAKernelCallExpr>("cudaLaunchKernel")) if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs<clang::CUDAKernelCallExpr>("cudaLaunchKernel"))
{ {
LangOptions DefaultLangOptions; LangOptions DefaultLangOptions;
StringRef initialParamList;
const FunctionDecl * kernelDecl = launchKernel->getDirectCallee(); SmallString<40> XStr;
raw_svector_ostream OS(XStr);
const ParmVarDecl * pvdFirst = kernelDecl->getParamDecl(0); OS << "hipLaunchParm lp";
const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1); const FunctionDecl * kernelDecl = launchKernel->getDirectCallee();
SourceLocation kernelArgListStart(pvdFirst->getLocStart()); SourceLocation l1 = kernelDecl->getNameInfo().getLocStart();
SourceLocation kernelArgListEnd(pvdLast->getLocEnd()); l1.dump(*SM);llvm::outs() << "\n";
SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(l1, clang::tok::l_paren, *SM, DefaultLangOptions, true);
size_t replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); kernelArgListStart.dump(*SM);llvm::outs() << "\n";
std::string outs(SM->getCharacterData(kernelArgListStart), replacementLength); size_t replacementLength = 0;
llvm::outs() << "initial paramlist: " << outs.c_str() << "\n"; if (kernelDecl->getNumParams() > 0) {
outs = "hipLaunchParm lp, " + outs; //const ParmVarDecl * pvdFirst = kernelDecl->getParamDecl(0);
llvm::outs() << "new paramlist: " << outs.c_str() << "\n"; const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1);
Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, outs); //kernelArgListStart = SourceLocation(pvdFirst->getLocStart());
Replace->insert(Rep0); SourceLocation kernelArgListEnd = SourceLocation(pvdLast->getLocEnd());
SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions);
replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart);
std::string name = kernelDecl->getDeclName().getAsString(); initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), replacementLength);
std::string repName = "hipLaunchKernel(HIP_KERNEL_NAME(" + name + "), "; OS << ", " << initialParamList;
}
const CallExpr * config = launchKernel->getConfig(); llvm::outs() << "initial paramlist: " << initialParamList << "\n";
llvm::outs() << "\nKernel config arguments:\n"; llvm::outs() << "new paramlist: " << OS.str() << "\n";
for (unsigned argno = 0; argno < config->getNumArgs(); argno++) Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, OS.str());
{ Replace->insert(Rep0);
const Expr * arg = config->getArg(argno);
if (!isa<CXXDefaultArgExpr>(arg)) { XStr.clear();
std::string typeCtor = ""; OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << kernelDecl->getName() << "), ";
const ParmVarDecl * pvd = config->getDirectCallee()->getParamDecl(argno);
const CallExpr * config = launchKernel->getConfig();
SourceLocation sl(arg->getLocStart()); llvm::outs() << "\nKernel config arguments:\n";
SourceLocation el(arg->getLocEnd()); for (unsigned argno = 0; argno < config->getNumArgs(); argno++)
SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions); {
std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl)); const Expr * arg = config->getArg(argno);
llvm::outs() << "args[ " << argno << "]" << outs.c_str() << " <" << pvd->getType().getAsString() << ">\n"; if (!isa<CXXDefaultArgExpr>(arg)) {
if (pvd->getType().getAsString().compare("dim3") == 0) const ParmVarDecl * pvd = config->getDirectCallee()->getParamDecl(argno);
repName += " dim3(" + outs + "),";
else SourceLocation sl(arg->getLocStart());
repName += " " + outs + ","; SourceLocation el(arg->getLocEnd());
} else SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
repName += " 0,"; StringRef outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
} llvm::outs() << "args[ " << argno << "]" << outs << " <" << pvd->getType().getAsString() << ">\n";
if (pvd->getType().getAsString().compare("dim3") == 0)
for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++) OS << " dim3(" << outs << "),";
{ else
const Expr * arg = launchKernel->getArg(argno); OS << " " << outs << ",";
SourceLocation sl(arg->getLocStart()); } else
SourceLocation el(arg->getLocEnd()); OS << " 0,";
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"; for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++)
repName += " " + outs + ","; {
} const Expr * arg = launchKernel->getArg(argno);
repName.pop_back(); SourceLocation sl(arg->getLocStart());
repName += ")"; SourceLocation el(arg->getLocEnd());
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) - SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
SM->getCharacterData(launchKernel->getLocStart()); std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
Replacement Rep(*SM, launchKernel->getLocStart(), length, repName); llvm::outs() << outs.c_str() << "\n";
Replace->insert(Rep); OS << " " << outs << ",";
} }
XStr.pop_back();
if (const MemberExpr * threadIdx = Result.Nodes.getNodeAs<clang::MemberExpr>("cudaBuiltin")) OS << ")";
{ size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) -
if (const OpaqueValueExpr * refBase = dyn_cast<OpaqueValueExpr>(threadIdx->getBase())) { SM->getCharacterData(launchKernel->getLocStart());
if (const DeclRefExpr * declRef = dyn_cast<DeclRefExpr>(refBase->getSourceExpr())) { Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str());
std::string name = declRef->getDecl()->getNameAsString(); Replace->insert(Rep);
std::string memberName = threadIdx->getMemberDecl()->getNameAsString(); }
size_t pos = memberName.find_first_not_of("__fetch_builtin_");
memberName = memberName.substr(pos, memberName.length() - pos); if (const MemberExpr * threadIdx = Result.Nodes.getNodeAs<clang::MemberExpr>("cudaBuiltin"))
{
if (const OpaqueValueExpr * refBase = dyn_cast<OpaqueValueExpr>(threadIdx->getBase())) {
if (const DeclRefExpr * declRef = dyn_cast<DeclRefExpr>(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; name += "." + memberName;
std::string repName = N.cuda2hipRename[name]; std::string repName = N.cuda2hipRename[name];
SourceLocation sl = threadIdx->getLocStart(); SourceLocation sl = threadIdx->getLocStart();