AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge branch 'master' of https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-hipify

Просмотр исходного кода 
Conflicts:
	src/Cuda2Hip.cpp


[ROCm/clr commit: 1b926d03b5]
Этот коммит содержится в:
atimofee
2016-03-01 21:51:01 +03:00
родитель d4a54ef91c 7990aba877
Коммит f43e55503f
5 изменённых файлов: 128 добавлений и 126 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/clr/hipamd/CMakeLists.txt
+1 -2
Просмотреть файл
@@ -14,9 +14,8 @@ if(${HIPIFY_STANDLONE})
include_directories(${LLVM_INCLUDE_DIRS})
link_directories(${LLVM_LIBRARY_DIRS})
add_definitions(${LLVM_DEFINITIONS})
add_llvm_executable( hipify src/Cuda2Hip.cpp )
add_llvm_executable(hipify src/Cuda2Hip.cpp )
find_program(LIT_COMMAND lit)
else()
set(LLVM_LINK_COMPONENTS
Option
projects/clr/hipamd/src/Cuda2Hip.cpp
+123 -121
Просмотреть файл
@@ -158,7 +158,7 @@ namespace {
cuda2hipRename["cudaStreamDefault"] = "hipStreamDefault";
cuda2hipRename["cudaStreamNonBlocking"] = "hipStreamNonBlocking";
// Other synchronization
// Other synchronization
cuda2hipRename["cudaDeviceSynchronize"] = "hipDeviceSynchronize";
cuda2hipRename["cudaThreadSynchronize"] = "hipDeviceSynchronize"; // translate deprecated cudaThreadSynchronize
cuda2hipRename["cudaDeviceReset"] = "hipDeviceReset";
@@ -166,7 +166,7 @@ namespace {
cuda2hipRename["cudaSetDevice"] = "hipSetDevice";
cuda2hipRename["cudaGetDevice"] = "hipGetDevice";
// Device
// Device
cuda2hipRename["cudaDeviceProp"] = "hipDeviceProp_t";
cuda2hipRename["cudaGetDeviceProperties"] = "hipDeviceGetProperties";
@@ -357,11 +357,11 @@ namespace {
}
}
void EndOfMainFile() override
{
void EndOfMainFile() override
{
}
bool SeenEnd;
void setSourceManager(SourceManager * sm) { _sm = sm; }
void setPreprocessor (Preprocessor * pp) { _pp = pp; }
@@ -376,102 +376,103 @@ namespace {
};
class Cuda2HipCallback : public MatchFinder::MatchCallback {
public:
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, 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);
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();
@@ -558,23 +559,24 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback {
if (const StringLiteral * stringLiteral = Result.Nodes.getNodeAs<clang::StringLiteral>("stringLiteral"))
{
StringRef s = stringLiteral->getString();
size_t begin = 0;
while ((begin = s.find("cuda", begin)) != StringRef::npos) {
const size_t end = s.find_first_of(" ", begin + 4);
StringRef name = s.slice(begin, end);
llvm::outs() << "\nToken: <" << name << "> found in string literal " << s << "\n";
StringRef repName = N.cuda2hipRename[name];
if (!repName.empty()) {
llvm::outs() << "\nWill be replaced with: <" << repName << "\n";
SourceLocation sl = stringLiteral->getLocationOfByte(begin, *SM,
Result.Context->getLangOpts(), Result.Context->getTargetInfo());
Replacement Rep(*SM, sl, name.size(), repName);
Replace->insert(Rep);
}
if (end == StringRef::npos) break;
begin = end + 1;
}
StringRef s = stringLiteral->getString();
size_t begin = 0;
while ((begin = s.find("cuda", begin)) != StringRef::npos) {
const size_t end = s.find_first_of(" ", begin + 4);
StringRef name = s.slice(begin, end);
llvm::outs() << "\nToken: <" << name << "> found in string literal " << s << "\n";
StringRef repName = N.cuda2hipRename[name];
if (!repName.empty()) {
llvm::outs() << "\nWill be replaced with: <" << repName << "\n";
SourceLocation sl = stringLiteral->getLocationOfByte(begin, *SM,
Result.Context->getLangOpts(), Result.Context->getTargetInfo());
Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ?
SM->getImmediateSpellingLoc(sl) : sl, name.size(), repName);
Replace->insert(Rep);
}
if (end == StringRef::npos) break;
begin = end + 1;
}
}
if (const UnaryExprOrTypeTraitExpr * expr = Result.Nodes.getNodeAs<clang::UnaryExprOrTypeTraitExpr>("cudaStructSizeOf"))
@@ -609,7 +611,7 @@ int main(int argc, const char **argv) {
llvm::sys::PrintStackTraceOnErrorSignal();
int Result;
CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::Required);
std::string dst = OutputFilename;
std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
@@ -642,9 +644,9 @@ int main(int argc, const char **argv) {
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(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(expr(sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(cxxRecordDecl(matchesName("cuda.*"))))))).bind("cudaStructSizeOf"), &Callback);
Finder.addMatcher(stringLiteral().bind("stringLiteral"), &Callback);
auto action = newFrontendActionFactory(&Finder, &PPCallbacks);
@@ -663,7 +665,7 @@ int main(int argc, const char **argv) {
Tool.clearArgumentsAdjusters();
}
LangOptions DefaultLangOptions;
IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts = new DiagnosticOptions();
TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
projects/clr/hipamd/test/CMakeLists.txt
+2 -2
Просмотреть файл
@@ -6,7 +6,7 @@ if( NOT PYTHONINTERP_FOUND )
"Please install Python or specify the PYTHON_EXECUTABLE CMake variable.")
endif()
set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}/.. )
configure_file(
${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.in
@@ -16,7 +16,7 @@ configure_file(
add_lit_testsuite(check-hipify "Running HIPify regression tests"
${CMAKE_CURRENT_SOURCE_DIR}
PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg
DEPENDS hipify
DEPENDS hipify FileCheck lit
)
add_custom_target(check)
projects/clr/hipamd/test/lit.cfg
+1
Просмотреть файл
@@ -51,4 +51,5 @@ if not tool_path:
lit_config.note('Did not find ' + tool_name + ' in ' + llvm_tools_dir)
tool_path = llvm_tools_dir + '/' + tool_name
config.substitutions.append((tool_name, tool_path))
config.substitutions.append(("hipify", obj_root+"/hipify"))
projects/clr/hipamd/test/lit.site.cfg.in
+1 -1
Просмотреть файл
@@ -1,7 +1,7 @@
import sys
config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
config.obj_root = "@BINARY_DIR@"
# Support substitution of the tools and libs dirs with user parameters. This is
# used when we can't determine the tool dir at configuration time.