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

[HIPIFY][perl] Generate transformKernelLaunch

Browse Source 
+ Update hipify-perl accordingly
+ Minor refactoring


[ROCm/hip commit: c8f73ea491]
This commit is contained in:
Evgeny Mankov
2019-10-01 11:17:27 +03:00
parent ef49c4f789
commit 34e20b8568
2 changed files with 229 additions and 183 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp
+56 -12
View File
@@ -64,10 +64,11 @@ namespace perl {
const std::string triple_tab = double_tab + tab;
const std::string sSub = "sub";
const std::string sReturn_0 = "return 0;\n";
const std::string sReturn_m = "return $m;\n";
const std::string sReturn_k = "return $k;\n";
const std::string sForeach = "foreach $func (\n";
const std::string sMy = "my $m = 0;\n";
const std::string sMy_k = "my $k = 0;";
const std::string sNoWarns = "no warnings qw/uninitialized/;";
const std::string sCudaDevice = "cudaDevice";
const std::string sCudaDeviceId = "cudaDeviceId";
const std::string sCudaDevices = "cudaDevices";
@@ -182,17 +183,59 @@ namespace perl {
void generateExternShared(std::unique_ptr<std::ostream>& streamPtr) {
*streamPtr.get() << std::endl << "# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro" << std::endl;
*streamPtr.get() << sSub << " transformExternShared" << " {" << std::endl;
*streamPtr.get() << tab << "no warnings qw/uninitialized/;" << std::endl;
*streamPtr.get() << tab << "my $k = 0;" << std::endl;
*streamPtr.get() << tab << sNoWarns << std::endl;
*streamPtr.get() << tab << sMy_k << std::endl;
*streamPtr.get() << tab << "$k += s/extern\\s+([\\w\\(\\)]+)?\\s*__shared__\\s+([\\w:<>\\s]+)\\s+(\\w+)\\s*\\[\\s*\\]\\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g;" << std::endl;
*streamPtr.get() << tab << "$ft{ 'extern_shared' } += $k;" << std::endl << "}" << std::endl;
}
void generateKernelLaunch(std::unique_ptr<std::ostream>& streamPtr) {
*streamPtr.get() << std::endl << "# CUDA Kernel Launch Syntax" << std::endl;
*streamPtr.get() << sSub << " transformKernelLaunch" << " {" << std::endl;
*streamPtr.get() << tab << "my $TkernRef = @_;" << std::endl;
*streamPtr.get() << tab << sNoWarns << std::endl;
*streamPtr.get() << tab << sMy_k << std::endl << std::endl;
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db, Ns, S>>>() syntax with empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, $6)/g;" << std::endl;
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns, S>>>() syntax with empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5)/g;" << std::endl << std::endl;
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db, Ns, S>>>(...) syntax with non-empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, $6, /g;" << std::endl;
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns, S>>>(...) syntax with non-empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5, /g;" << std::endl << std::endl;
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db, Ns>>>() syntax with empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, 0)/g;" << std::endl;
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns>>>() syntax with empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0)/g;" << std::endl << std::endl;
*streamPtr.get() << tab << "# Handle the kern<...><<Dg, Db, Ns>>>(...) syntax with non-empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, 0, /g;" << std::endl;
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns>>>(...) syntax with non-empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g;" << std::endl << std::endl;
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db>>>() syntax with empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), 0, 0)/g;" << std::endl;
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db>>>() syntax with empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g;" << std::endl << std::endl;
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db>>>(...) syntax with non-empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), 0, 0, /g;" << std::endl;
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db>>>(...) syntax with non-empty args:" << std::endl;
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g;" << std::endl << std::endl;
*streamPtr.get() << tab << "if ($k) {" << std::endl;
*streamPtr.get() << double_tab << "$ft{'kernel_launch'} += $k;" << std::endl;
*streamPtr.get() << double_tab << "@$TkernRef{$1} ++;" << std::endl << tab << "}" << std::endl << "}" << std::endl;
}
void generateHostFunctions(std::unique_ptr<std::ostream>& streamPtr) {
*streamPtr.get() << std::endl << sSub << " transformHostFunctions" << " {" << std::endl << tab << sMy;
*streamPtr.get() << std::endl << sSub << " transformHostFunctions" << " {" << std::endl << tab << sMy_k << std::endl;
std::set<std::string> &funcSet = DeviceSymbolFunctions0;
const std::string s0 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
const std::string s1 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, ";
const std::string s0 = "$k += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
const std::string s1 = "$k += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, ";
for (int i = 0; i < 4; ++i) {
*streamPtr.get() << tab + sForeach;
switch (i) {
@@ -223,7 +266,7 @@ namespace perl {
}
*streamPtr.get() << tab << "}" << std::endl;
}
*streamPtr.get() << tab << sReturn_m << "}" << std::endl;
*streamPtr.get() << tab << sReturn_k << "}" << std::endl;
}
void generateDeviceFunctions(std::unique_ptr<std::ostream>& streamPtr) {
@@ -240,7 +283,7 @@ namespace perl {
std::stringstream subCountSupported;
std::stringstream subWarnUnsupported;
std::stringstream subCommon;
std::string sCommon = tab + sMy + tab + sForeach;
std::string sCommon = tab + sMy_k + "\n" + tab + sForeach;
subCountSupported << std::endl << sSub << " countSupportedDeviceFunctions" << " {" << std::endl << (countSupported ? sCommon : tab + sReturn_0);
subWarnUnsupported << std::endl << sSub << " warnUnsupportedDeviceFunctions" << " {" << std::endl << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0);
if (countSupported) {
@@ -256,7 +299,7 @@ namespace perl {
subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl;
subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl;
subCommon << double_tab << "if ($mt && !$mt_namespace) {" << std::endl;
subCommon << triple_tab << "$m += $mt;" << std::endl;
subCommon << triple_tab << "$k += $mt;" << std::endl;
}
if (countSupported) {
subCountSupported << subCommon.str();
@@ -266,7 +309,7 @@ namespace perl {
subWarnUnsupported << triple_tab << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";" << std::endl;
}
if (countSupported || countUnsupported) {
sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_m;
sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_k;
}
if (countSupported) subCountSupported << sCommon;
if (countUnsupported) subWarnUnsupported << sCommon;
@@ -305,6 +348,7 @@ namespace perl {
generateStatFunctions(streamPtr);
generateSimpleSubstitutions(streamPtr);
generateExternShared(streamPtr);
generateKernelLaunch(streamPtr);
generateHostFunctions(streamPtr);
generateDeviceFunctions(streamPtr);
streamPtr.get()->flush();