Merge branch 'master' of https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-hipify
[ROCm/hip commit: d5861b078c]
This commit is contained in:
dfukalov
@@ -183,6 +183,7 @@ namespace {
|
||||
cuda2hipRename["cudaFuncSetCacheConfig"] = "hipFuncSetCacheConfig";
|
||||
|
||||
cuda2hipRename["cudaDriverGetVersion"] = "hipDriverGetVersion";
|
||||
cuda2hipRename["cudaRuntimeGetVersion"] = "hipRuntimeGetVersion";
|
||||
|
||||
// Peer2Peer
|
||||
cuda2hipRename["cudaDeviceCanAccessPeer"] = "hipDeviceCanAccessPeer";
|
||||
@@ -307,98 +308,103 @@ 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<clang::CallExpr>("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<clang::CUDAKernelCallExpr>("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<CXXDefaultArgExpr>(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<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);
|
||||
public:
|
||||
Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {}
|
||||
|
||||
void run(const MatchFinder::MatchResult &Result) override {
|
||||
|
||||
SourceManager * SM = Result.SourceManager;
|
||||
|
||||
if (const CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("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<clang::CUDAKernelCallExpr>("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 * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1);
|
||||
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<CXXDefaultArgExpr>(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<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;
|
||||
std::string repName = N.cuda2hipRename[name];
|
||||
SourceLocation sl = threadIdx->getLocStart();
|
||||
@@ -434,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<clang::VarDecl>("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<clang::ParmVarDecl>("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<clang::ParmVarDecl>("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<clang::StringLiteral>("stringLiteral"))
|
||||
{
|
||||
StringRef s = stringLiteral->getString();
|
||||
@@ -471,6 +517,20 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback {
|
||||
begin = end + 1;
|
||||
}
|
||||
}
|
||||
|
||||
if (const UnaryExprOrTypeTraitExpr * expr = Result.Nodes.getNodeAs<clang::UnaryExprOrTypeTraitExpr>("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:
|
||||
@@ -523,7 +583,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);
|
||||
|
||||
Reference in New Issue
Block a user