From b2ceadc19755df532c580a1952f4b57cf6ce53ce Mon Sep 17 00:00:00 2001 From: atimofee Date: Thu, 25 Feb 2016 22:54:58 +0300 Subject: [PATCH 1/3] Fixed tool crash on kernels with empty parameter list [ROCm/hip commit: e2379adae30c0bdf5a0cd8010e458ae8065de297] --- projects/hip/src/Cuda2Hip.cpp | 191 ++++++++++++++++++---------------- 1 file changed, 99 insertions(+), 92 deletions(-) diff --git a/projects/hip/src/Cuda2Hip.cpp b/projects/hip/src/Cuda2Hip.cpp index 5a1432a185..57df437e77 100644 --- a/projects/hip/src/Cuda2Hip.cpp +++ b/projects/hip/src/Cuda2Hip.cpp @@ -307,98 +307,105 @@ namespace { }; 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); + 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; + StringRef initialParamList; + SmallString<40> XStr; + raw_svector_ostream OS(XStr); + OS << "hipLaunchParm lp"; + const FunctionDecl * kernelDecl = launchKernel->getDirectCallee(); + SourceLocation l1 = kernelDecl->getNameInfo().getLocStart(); + l1.dump(*SM);llvm::outs() << "\n"; + SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(l1, clang::tok::l_paren, *SM, DefaultLangOptions, true); + kernelArgListStart.dump(*SM);llvm::outs() << "\n"; + size_t replacementLength = 0; + if (kernelDecl->getNumParams() > 0) { + //const ParmVarDecl * pvdFirst = kernelDecl->getParamDecl(0); + const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1); + //kernelArgListStart = SourceLocation(pvdFirst->getLocStart()); + SourceLocation kernelArgListEnd = SourceLocation(pvdLast->getLocEnd()); + SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); + replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); + initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), replacementLength); + OS << ", " << initialParamList; + } + + llvm::outs() << "initial paramlist: " << initialParamList << "\n"; + llvm::outs() << "new paramlist: " << OS.str() << "\n"; + Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, OS.str()); + Replace->insert(Rep0); + + XStr.clear(); + OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << kernelDecl->getName() << "), "; + + 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)) { + 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)); + llvm::outs() << "args[ " << argno << "]" << outs << " <" << pvd->getType().getAsString() << ">\n"; + if (pvd->getType().getAsString().compare("dim3") == 0) + OS << " dim3(" << outs << "),"; + else + OS << " " << outs << ","; + } else + OS << " 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"; + OS << " " << outs << ","; + } + XStr.pop_back(); + OS << ")"; + 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 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(); From 5304ddfa1deb2305227f0e30ee954a05665c54e8 Mon Sep 17 00:00:00 2001 From: atimofee Date: Fri, 26 Feb 2016 23:35:27 +0300 Subject: [PATCH 2/3] =?UTF-8?q?1.CUDA=20structure=20types=20(and=20pointer?= =?UTF-8?q?s=20to=20the=20types)=20in=20function=20parameter=20declaration?= =?UTF-8?q?=20conversion=20=20=20=E2=80=93=20DONE.=202.Wrong=20source=20lo?= =?UTF-8?q?cations=20in=20qualified=20types=20declaration:=20=20=E2=80=93?= =?UTF-8?q?=20FIXED=203.cudaRuntimeGetVersion=20added=20to=20the=20names?= =?UTF-8?q?=20map=204.sizeof=20expression=20matcher=20added,=20expression?= =?UTF-8?q?=20handler=20is=20not=20yet=20ready=20-=20upcoming=20soon?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit [ROCm/hip commit: 5e3ab927ad2dbc2a23a627053023a1352e071833] --- projects/hip/src/Cuda2Hip.cpp | 58 ++++++++++++++++++++++++++++++----- 1 file changed, 50 insertions(+), 8 deletions(-) diff --git a/projects/hip/src/Cuda2Hip.cpp b/projects/hip/src/Cuda2Hip.cpp index 57df437e77..bb02357869 100644 --- a/projects/hip/src/Cuda2Hip.cpp +++ b/projects/hip/src/Cuda2Hip.cpp @@ -183,6 +183,7 @@ namespace { cuda2hipRename["cudaFuncSetCacheConfig"] = "hipFuncSetCacheConfig"; cuda2hipRename["cudaDriverGetVersion"] = "hipDriverGetVersion"; + cuda2hipRename["cudaRuntimeGetVersion"] = "hipRuntimeGetVersion"; // Peer2Peer cuda2hipRename["cudaDeviceCanAccessPeer"] = "hipDeviceCanAccessPeer"; @@ -341,9 +342,7 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { kernelArgListStart.dump(*SM);llvm::outs() << "\n"; size_t replacementLength = 0; if (kernelDecl->getNumParams() > 0) { - //const ParmVarDecl * pvdFirst = kernelDecl->getParamDecl(0); const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1); - //kernelArgListStart = SourceLocation(pvdFirst->getLocStart()); SourceLocation kernelArgListEnd = SourceLocation(pvdLast->getLocEnd()); SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions); replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart); @@ -441,22 +440,62 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { std::string name = cudaStructVar->getType()->getAsStructureType()->getDecl()->getNameAsString(); std::string repName = N.cuda2hipRename[name]; - SourceLocation sl = cudaStructVar->getLocStart(); + TypeLoc TL = cudaStructVar->getTypeSourceInfo()->getTypeLoc(); + SourceLocation sl = TL.getUnqualifiedLoc().getLocStart(); Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); Replace->insert(Rep); } + 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]; + TypeLoc TL = cudaStructVarPtr->getTypeSourceInfo()->getTypeLoc(); + SourceLocation sl = TL.getUnqualifiedLoc().getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.size(), repName); + Replace->insert(Rep); + } + } + + if (const ParmVarDecl * cudaParamDecl = Result.Nodes.getNodeAs("cudaParamDecl")) { - QualType QT = cudaParamDecl->getOriginalType(); - std::string name = QT.getAsString(); - std::string repName = N.cuda2hipRename[name]; - SourceLocation sl = cudaParamDecl->getLocStart(); + QualType QT = cudaParamDecl->getOriginalType().getUnqualifiedType(); + StringRef name = QT.getAsString(); + const Type * t = QT.getTypePtr(); + if (t->isStructureOrClassType()) { + name = t->getAsCXXRecordDecl()->getName(); + } + StringRef repName = N.cuda2hipRename[name]; + TypeLoc TL = cudaParamDecl->getTypeSourceInfo()->getTypeLoc(); + SourceLocation sl = TL.getUnqualifiedLoc().getLocStart(); Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? - SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName); + SM->getImmediateSpellingLoc(sl) : sl, name.size(), repName); Replace->insert(Rep); } + + 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()); + StringRef repName = N.cuda2hipRename[name]; + TypeLoc TL = cudaParamDeclPtr->getTypeSourceInfo()->getTypeLoc(); + SourceLocation sl = TL.getUnqualifiedLoc().getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.size(), repName); + Replace->insert(Rep); + } + } + + if (const StringLiteral * stringLiteral = Result.Nodes.getNodeAs("stringLiteral")) { StringRef s = stringLiteral->getString(); @@ -530,7 +569,10 @@ int main(int argc, const char **argv) { 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(sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(cxxRecordDecl(matchesName("cuda.*"))))))).bind("cudaStructSizeOf"), &Callback); Finder.addMatcher(stringLiteral().bind("stringLiteral"), &Callback); auto action = newFrontendActionFactory(&Finder, &PPCallbacks); From 2e75136197cc53fe73be1f7d2f3fe4a88f970318 Mon Sep 17 00:00:00 2001 From: atimofee Date: Mon, 29 Feb 2016 20:27:38 +0300 Subject: [PATCH 3/3] CUDA type names in sizeof() expression conversion added [ROCm/hip commit: 432500056a8b8af01910001d884569d4660bb950] --- projects/hip/src/Cuda2Hip.cpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/projects/hip/src/Cuda2Hip.cpp b/projects/hip/src/Cuda2Hip.cpp index bb02357869..c86a80bfb7 100644 --- a/projects/hip/src/Cuda2Hip.cpp +++ b/projects/hip/src/Cuda2Hip.cpp @@ -517,6 +517,20 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { begin = end + 1; } } + + if (const UnaryExprOrTypeTraitExpr * expr = Result.Nodes.getNodeAs("cudaStructSizeOf")) + { + TypeSourceInfo * typeInfo = expr->getArgumentTypeInfo(); + QualType QT = typeInfo->getType().getUnqualifiedType(); + const Type * type = QT.getTypePtr(); + StringRef name = type->getAsCXXRecordDecl()->getName(); + StringRef repName = N.cuda2hipRename[name]; + TypeLoc TL = typeInfo->getTypeLoc(); + SourceLocation sl = TL.getUnqualifiedLoc().getLocStart(); + Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? + SM->getImmediateSpellingLoc(sl) : sl, name.size(), repName); + Replace->insert(Rep); + } } private: