diff --git a/hipify-clang/src/CUDA2HIP_Perl.cpp b/hipify-clang/src/CUDA2HIP_Perl.cpp index f74d196cac..f0a7c093f8 100644 --- a/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -32,11 +32,12 @@ THE SOFTWARE. #include "LLVMCompat.h" #include "Statistics.h" -using namespace llvm; - namespace perl { - const std::string sCopyright = + using namespace std; + using namespace llvm; + + const string sCopyright = "##\n" "# Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.\n" "#\n" @@ -59,434 +60,410 @@ namespace perl { "# THE SOFTWARE.\n" "##\n"; - const std::string tab = " "; - const std::string double_tab = tab + tab; - const std::string triple_tab = double_tab + tab; - const std::string quad_tab = triple_tab + tab; - const std::string tab_5 = quad_tab + tab; - const std::string tab_6 = tab_5 + tab; - const std::string sSub = "sub"; - const std::string sReturn_0 = "return 0;\n"; - const std::string sReturn_k = "return $k;\n"; - const std::string sForeach = "foreach $func (\n"; - const std::string sMy = "my "; - const std::string sMy_k = sMy + "$k = 0;"; - const std::string sNoWarns = "no warnings qw/uninitialized/;"; + const string tab = " "; + const string tab_2 = tab + tab; + const string tab_3 = tab_2 + tab; + const string tab_4 = tab_3 + tab; + const string tab_5 = tab_4 + tab; + const string tab_6 = tab_5 + tab; + const string endl_2 = "\n\n"; + const string endl_tab = "\n" + tab; + const string endl_tab_2 = "\n" + tab_2; + const string endl_tab_3 = "\n" + tab_3; + const string endl_tab_4 = "\n" + tab_4; + const string endl_tab_5 = "\n" + tab_5; + const string sub = "sub "; + const string my = "my "; + const string my_k = my + "$k = 0;"; + const string return_0 = "return 0;\n"; + const string return_k = "return $k;\n"; + const string while_ = "while "; + const string unless_ = "unless "; + const string foreach = "foreach "; + const string foreach_func = foreach + "$func (\n"; + const string print = "print STDERR "; + const string printf = "printf STDERR "; + const string no_warns = "no warnings qw/uninitialized/;"; + const string hipify_perl = "hipify-perl"; - const std::string sCudaDevice = "cudaDevice"; - const std::string sCudaDeviceId = "cudaDeviceId"; - const std::string sCudaDevices = "cudaDevices"; - const std::string sCudaDevice_t = "cudaDevice_t"; - const std::string sCudaIDs = "cudaIDs"; - const std::string sCudaGridDim = "cudaGridDim"; - const std::string sCudaDimGrid = "cudaDimGrid"; - const std::string sCudaDimBlock = "cudaDimBlock"; - const std::string sCudaGradInput = "cudaGradInput"; - const std::string sCudaGradOutput = "cudaGradOutput"; - const std::string sCudaInput = "cudaInput"; - const std::string sCudaOutput = "cudaOutput"; - const std::string sCudaIndices = "cudaIndices"; - const std::string sCudaGaugeField = "cudaGaugeField"; - const std::string sCudaMom = "cudaMom"; - const std::string sCudaGauge = "cudaGauge"; - const std::string sCudaInGauge = "cudaInGauge"; - const std::string sCudaColorSpinorField = "cudaColorSpinorField"; - const std::string sCudaSiteLink = "cudaSiteLink"; - const std::string sCudaFatLink = "cudaFatLink"; - const std::string sCudaStaple = "cudaStaple"; - const std::string sCudaCloverField = "cudaCloverField"; - const std::string sCudaParam = "cudaParam"; + const string sCudaDevice = "cudaDevice"; + const string sCudaDeviceId = "cudaDeviceId"; + const string sCudaDevices = "cudaDevices"; + const string sCudaDevice_t = "cudaDevice_t"; + const string sCudaIDs = "cudaIDs"; + const string sCudaGridDim = "cudaGridDim"; + const string sCudaDimGrid = "cudaDimGrid"; + const string sCudaDimBlock = "cudaDimBlock"; + const string sCudaGradInput = "cudaGradInput"; + const string sCudaGradOutput = "cudaGradOutput"; + const string sCudaInput = "cudaInput"; + const string sCudaOutput = "cudaOutput"; + const string sCudaIndices = "cudaIndices"; + const string sCudaGaugeField = "cudaGaugeField"; + const string sCudaMom = "cudaMom"; + const string sCudaGauge = "cudaGauge"; + const string sCudaInGauge = "cudaInGauge"; + const string sCudaColorSpinorField = "cudaColorSpinorField"; + const string sCudaSiteLink = "cudaSiteLink"; + const string sCudaFatLink = "cudaFatLink"; + const string sCudaStaple = "cudaStaple"; + const string sCudaCloverField = "cudaCloverField"; + const string sCudaParam = "cudaParam"; - const std::set Whitelist{ + const set Whitelist{ {sCudaDevice}, {sCudaDevice_t}, {sCudaIDs}, {sCudaGridDim}, {sCudaDimGrid}, {sCudaDimBlock}, {sCudaDeviceId}, {sCudaDevices}, {sCudaGradInput}, {sCudaGradOutput}, {sCudaInput}, {sCudaOutput}, {sCudaIndices}, {sCudaGaugeField}, {sCudaMom}, {sCudaGauge}, {sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam} }; - void generateHeader(std::unique_ptr& streamPtr) { - *streamPtr.get() << "#!/usr/bin/perl -w" << std::endl << std::endl; - *streamPtr.get() << sCopyright << std::endl; - *streamPtr.get() << "#usage hipify-perl [OPTIONS] INPUT_FILE" << std::endl << std::endl; - *streamPtr.get() << "use Getopt::Long;" << std::endl; - *streamPtr.get() << sMy << "$whitelist = \"\";" << std::endl; - *streamPtr.get() << sMy << "$fileName = \"\";" << std::endl; - *streamPtr.get() << sMy << "%ft;" << std::endl; - *streamPtr.get() << sMy << "%Tkernels;" << std::endl << std::endl; - *streamPtr.get() << "GetOptions(" << std::endl; - *streamPtr.get() << tab << " \"examine\" => \\$examine # Combines -no-output and -print-stats options." << std::endl; - *streamPtr.get() << tab << ", \"inplace\" => \\$inplace # Modify input file inplace, replacing input with hipified output, save backup in .prehip file." << std::endl; - *streamPtr.get() << tab << ", \"no-output\" => \\$no_output # Don't write any translated output to stdout." << std::endl; - *streamPtr.get() << tab << ", \"print-stats\" => \\$print_stats # Print translation statistics." << std::endl; - *streamPtr.get() << tab << ", \"quiet-warnings\" => \\$quiet_warnings # Don't print warnings on unknown CUDA functions." << std::endl; - *streamPtr.get() << tab << ", \"whitelist=s\" => \\$whitelist # TODO: test it beforehand" << std::endl; - *streamPtr.get() << ");" << std::endl << std::endl; - *streamPtr.get() << "$print_stats = 1 if $examine;" << std::endl; - *streamPtr.get() << "$no_output = 1 if $examine;" << std::endl << std::endl; - *streamPtr.get() << "# Whitelist of cuda[A-Z] identifiers, which are commonly used in CUDA sources but don't map to any CUDA API:" << std::endl; + void generateHeader(unique_ptr& streamPtr) { + *streamPtr.get() << "#!/usr/bin/perl -w" << endl_2; + *streamPtr.get() << sCopyright << endl; + *streamPtr.get() << "#usage " << hipify_perl << " [OPTIONS] INPUT_FILE" << endl_2; + *streamPtr.get() << "use Getopt::Long;" << endl; + *streamPtr.get() << my << "$whitelist = \"\";" << endl; + *streamPtr.get() << my << "$fileName = \"\";" << endl; + *streamPtr.get() << my << "%ft;" << endl; + *streamPtr.get() << my << "%Tkernels;" << endl_2; + *streamPtr.get() << "GetOptions(" << endl; + *streamPtr.get() << tab << " \"examine\" => \\$examine # Combines -no-output and -print-stats options." << endl; + *streamPtr.get() << tab << ", \"inplace\" => \\$inplace # Modify input file inplace, replacing input with hipified output, save backup in .prehip file." << endl; + *streamPtr.get() << tab << ", \"no-output\" => \\$no_output # Don't write any translated output to stdout." << endl; + *streamPtr.get() << tab << ", \"print-stats\" => \\$print_stats # Print translation statistics." << endl; + *streamPtr.get() << tab << ", \"quiet-warnings\" => \\$quiet_warnings # Don't print warnings on unknown CUDA functions." << endl; + *streamPtr.get() << tab << ", \"whitelist=s\" => \\$whitelist # TODO: test it beforehand" << endl; + *streamPtr.get() << ");" << endl_2; + *streamPtr.get() << "$print_stats = 1 if $examine;" << endl; + *streamPtr.get() << "$no_output = 1 if $examine;" << endl_2; + *streamPtr.get() << "# Whitelist of cuda[A-Z] identifiers, which are commonly used in CUDA sources but don't map to any CUDA API:" << endl; *streamPtr.get() << "@whitelist = ("; unsigned int num = 0; - for (const std::string &m : Whitelist) { - *streamPtr.get() << std::endl << tab << (num ? ", " : " ") << "\"" << m << "\""; + for (const string &m : Whitelist) { + *streamPtr.get() << endl_tab << (num ? ", " : " ") << "\"" << m << "\""; ++num; } - *streamPtr.get() << std::endl << ");" << std::endl << std::endl; - *streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << std::endl << std::endl; + *streamPtr.get() << endl << ");" << endl_2; + *streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << endl_2; } - void generateStatFunctions(std::unique_ptr& streamPtr) { - *streamPtr.get() << std::endl << sSub << " totalStats" << " {" << std::endl; - *streamPtr.get() << tab << sMy << "%count = %{ shift() };" << std::endl; - *streamPtr.get() << tab << sMy << "$total = 0;" << std::endl; - *streamPtr.get() << tab << "foreach $key (keys %count) {" << std::endl; - *streamPtr.get() << double_tab << "$total += $count{$key};" << std::endl << tab << "}" << std::endl; - *streamPtr.get() << tab << "return $total;" << std::endl << "};" << std::endl; - *streamPtr.get() << std::endl << sSub << " printStats" << " {" << std::endl; - *streamPtr.get() << tab << sMy << "$label = shift();" << std::endl; - *streamPtr.get() << tab << sMy << "@statNames = @{ shift() };" << std::endl; - *streamPtr.get() << tab << sMy << "%counts = %{ shift() };" << std::endl; - *streamPtr.get() << tab << sMy << "$warnings = shift();" << std::endl; - *streamPtr.get() << tab << sMy << "$loc = shift();" << std::endl; - *streamPtr.get() << tab << sMy << "$total = totalStats(\\%counts);" << std::endl; - *streamPtr.get() << tab << "printf STDERR \"%s %d CUDA->HIP refs ( \", $label, $total;" << std::endl; - *streamPtr.get() << tab << "foreach $stat (@statNames) {" << std::endl; - *streamPtr.get() << double_tab << "printf STDERR \"%s:%d \", $stat, $counts{$stat};" << std::endl; - *streamPtr.get() << tab << "}" << std::endl; - *streamPtr.get() << tab << "printf STDERR \")\\n warn:%d LOC:%d\", $warnings, $loc;" << std::endl << "}" << std::endl; + void generateStatFunctions(unique_ptr& streamPtr) { + *streamPtr.get() << endl << sub << "totalStats" << " {" << endl; + *streamPtr.get() << tab << my << "%count = %{ shift() };" << endl; + *streamPtr.get() << tab << my << "$total = 0;" << endl; + *streamPtr.get() << tab << foreach << "$key (keys %count) {" << endl; + *streamPtr.get() << tab_2 << "$total += $count{$key};" << endl_tab << "}" << endl; + *streamPtr.get() << tab << "return $total;" << endl << "};" << endl; + *streamPtr.get() << endl << sub << "printStats" << " {" << endl; + *streamPtr.get() << tab << my << "$label = shift();" << endl; + *streamPtr.get() << tab << my << "@statNames = @{ shift() };" << endl; + *streamPtr.get() << tab << my << "%counts = %{ shift() };" << endl; + *streamPtr.get() << tab << my << "$warnings = shift();" << endl; + *streamPtr.get() << tab << my << "$loc = shift();" << endl; + *streamPtr.get() << tab << my << "$total = totalStats(\\%counts);" << endl; + *streamPtr.get() << tab << printf << "\"%s %d CUDA->HIP refs ( \", $label, $total;" << endl; + *streamPtr.get() << tab << foreach << "$stat (@statNames) {" << endl; + *streamPtr.get() << tab_2 << printf << "\"%s:%d \", $stat, $counts{$stat};" << endl_tab << "}" << endl; + *streamPtr.get() << tab << printf << "\")\\n warn:%d LOC:%d\", $warnings, $loc;" << endl << "}" << endl; for (int i = 0; i < 2; ++i) { - *streamPtr.get() << std::endl << sSub << " " << (i ? "clearStats" : "addStats") << " {" << std::endl; - *streamPtr.get() << tab << sMy << "$dest_ref = shift();" << std::endl; - *streamPtr.get() << tab << sMy << (i ? "@statNames = @{ shift() };" : "%adder = %{ shift() };") << std::endl; - *streamPtr.get() << tab << "foreach " << (i ? "$stat(@statNames)" : "$key (keys %adder)") << " {" << std::endl; - *streamPtr.get() << double_tab << "$dest_ref->" << (i ? "{$stat} = 0;" : "{$key} += $adder{$key};") << std::endl << tab << "}" << std::endl << "}" << std::endl; + *streamPtr.get() << endl << sub << (i ? "clearStats" : "addStats") << " {" << endl; + *streamPtr.get() << tab << my << "$dest_ref = shift();" << endl; + *streamPtr.get() << tab << my << (i ? "@statNames = @{ shift() };" : "%adder = %{ shift() };") << endl; + *streamPtr.get() << tab << foreach << (i ? "$stat(@statNames)" : "$key (keys %adder)") << " {" << endl; + *streamPtr.get() << tab_2 << "$dest_ref->" << (i ? "{$stat} = 0;" : "{$key} += $adder{$key};") << endl_tab << "}" << endl << "}" << endl; } } - void generateSimpleSubstitutions(std::unique_ptr& streamPtr) { - *streamPtr.get() << std::endl << sSub << " simpleSubstitutions" << " {" << std::endl; + void generateSimpleSubstitutions(unique_ptr& streamPtr) { + *streamPtr.get() << endl << sub << "simpleSubstitutions" << " {" << endl; for (int i = 0; i < NUM_CONV_TYPES; ++i) { if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) { for (auto& ma : CUDA_INCLUDE_MAP) { if (Statistics::isUnsupported(ma.second)) continue; if (i == ma.second.type) { - std::string sCUDA = ma.first.str(); - std::string sHIP = ma.second.hipName.str(); - sCUDA = std::regex_replace(sCUDA, std::regex("/"), "\\/"); - sHIP = std::regex_replace(sHIP, std::regex("/"), "\\/"); - *streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;" << std::endl; + string sCUDA = ma.first.str(); + string sHIP = ma.second.hipName.str(); + sCUDA = regex_replace(sCUDA, regex("/"), "\\/"); + sHIP = regex_replace(sHIP, regex("/"), "\\/"); + *streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;" << endl; } } } else { for (auto& ma : CUDA_RENAMES_MAP()) { if (Statistics::isUnsupported(ma.second)) continue; if (i == ma.second.type) { - *streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << std::endl; + *streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << endl; } } } } - *streamPtr.get() << "}" << std::endl; + *streamPtr.get() << "}" << endl; } - void generateExternShared(std::unique_ptr& 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 << 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 generateExternShared(unique_ptr& streamPtr) { + *streamPtr.get() << endl << "# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro" << endl; + *streamPtr.get() << sub << "transformExternShared" << " {" << endl; + *streamPtr.get() << tab << no_warns << endl; + *streamPtr.get() << tab << my_k << 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;" << endl; + *streamPtr.get() << tab << "$ft{'extern_shared'} += $k;" << endl << "}" << endl; } - void generateKernelLaunch(std::unique_ptr& streamPtr) { - *streamPtr.get() << std::endl << "# CUDA Kernel Launch Syntax" << std::endl; - *streamPtr.get() << sSub << " transformKernelLaunch" << " {" << std::endl; - *streamPtr.get() << tab << sNoWarns << std::endl; - *streamPtr.get() << tab << sMy_k << std::endl << std::endl; + void generateKernelLaunch(unique_ptr& streamPtr) { + *streamPtr.get() << endl << "# CUDA Kernel Launch Syntax" << endl << sub << "transformKernelLaunch" << " {" << endl; + *streamPtr.get() << tab << no_warns << endl; + *streamPtr.get() << tab << my_k << endl_2; - *streamPtr.get() << tab << "# Handle the kern<...><<>>() 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(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6)/g;" << std::endl; - *streamPtr.get() << tab << "# Handle the kern<<>>() 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<...><<>>() syntax with empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6)/g;" << endl; + *streamPtr.get() << tab << "# Handle the kern<<>>() syntax with empty args:" << 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;" << endl_2; - *streamPtr.get() << tab << "# Handle the kern<...><<>>(...) 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(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6, /g;" << std::endl; - *streamPtr.get() << tab << "# Handle the kern<<>>(...) 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<...><<>>(...) syntax with non-empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6, /g;" << endl; + *streamPtr.get() << tab << "# Handle the kern<<>>(...) syntax with non-empty args:" << 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;" << endl_2; - *streamPtr.get() << tab << "# Handle the kern<...><<>>() syntax with empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0)/g;" << std::endl; - *streamPtr.get() << tab << "# Handle the kern<<>>() 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<...><<>>() syntax with empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0)/g;" << endl; + *streamPtr.get() << tab << "# Handle the kern<<>>() syntax with empty args:" << 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;" << endl_2; - *streamPtr.get() << tab << "# Handle the kern<...><>>(...) syntax with non-empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0, /g;" << std::endl; - *streamPtr.get() << tab << "# Handle the kern<<>>(...) 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<...><>>(...) syntax with non-empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0, /g;" << endl; + *streamPtr.get() << tab << "# Handle the kern<<>>(...) syntax with non-empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g;" << endl_2; - *streamPtr.get() << tab << "# Handle the kern<...><<>>() syntax with empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0)/g;" << std::endl; - *streamPtr.get() << tab << "# Handle the kern<<>>() 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<...><<>>() syntax with empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0)/g;" << endl; + *streamPtr.get() << tab << "# Handle the kern<<>>() syntax with empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g;" << endl_2; - *streamPtr.get() << tab << "# Handle the kern<...><<>>(...) syntax with non-empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0, /g;" << std::endl; - *streamPtr.get() << tab << "# Handle the kern<<>>(...) 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 << "# Handle the kern<...><<>>(...) syntax with non-empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0, /g;" << endl; + *streamPtr.get() << tab << "# Handle the kern<<>>(...) syntax with non-empty args:" << endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g;" << endl_2; - *streamPtr.get() << tab << "if ($k) {" << std::endl; - *streamPtr.get() << double_tab << "$ft{'kernel_launch'} += $k;" << std::endl; - *streamPtr.get() << double_tab << "$Tkernels{$1}++;" << std::endl << tab << "}" << std::endl << "}" << std::endl; + *streamPtr.get() << tab << "if ($k) {" << endl; + *streamPtr.get() << tab_2 << "$ft{'kernel_launch'} += $k;" << endl; + *streamPtr.get() << tab_2 << "$Tkernels{$1}++;" << endl_tab << "}" << endl << "}" << endl; } - void generateHostFunctions(std::unique_ptr& streamPtr) { - *streamPtr.get() << std::endl << sSub << " transformHostFunctions" << " {" << std::endl << tab << sMy_k << std::endl; - std::set &funcSet = DeviceSymbolFunctions0; - const std::string s0 = "$k += s/(?& streamPtr) { + *streamPtr.get() << endl << sub << "transformHostFunctions" << " {" << endl_tab << my_k << endl; + set &funcSet = DeviceSymbolFunctions0; + const string s0 = "$k += s/(?second.hipName.str() << "\""; + *streamPtr.get() << (count ? ",\n" : "") << tab_2 << "\"" << found->second.hipName.str() << "\""; count++; } } - *streamPtr.get() << std::endl << tab << ")" << std::endl << tab << "{" << std::endl << double_tab; + *streamPtr.get() << endl_tab << ")" << endl_tab << "{" << endl_tab_2; switch (i) { case 0: - default: - *streamPtr.get() << s0 << sHIP_SYMBOL << "\\($2\\),/g" << std::endl; break; - case 1: - *streamPtr.get() << s1 << sHIP_SYMBOL << "\\($3\\)$4/g;" << std::endl; break; - case 2: - *streamPtr.get() << s0 << s_reinterpret_cast << "\\($2\\),/g" << std::endl; break; - case 3: - *streamPtr.get() << s1 << s_reinterpret_cast << "\\($3\\)$4/g;" << std::endl; break; + default: *streamPtr.get() << s0 << sHIP_SYMBOL << "\\($2\\),/g" << endl; break; + case 1: *streamPtr.get() << s1 << sHIP_SYMBOL << "\\($3\\)$4/g;" << endl; break; + case 2: *streamPtr.get() << s0 << s_reinterpret_cast << "\\($2\\),/g" << endl; break; + case 3: *streamPtr.get() << s1 << s_reinterpret_cast << "\\($3\\)$4/g;" << endl; break; } - *streamPtr.get() << tab << "}" << std::endl; + *streamPtr.get() << tab << "}" << endl; } - *streamPtr.get() << tab << sReturn_k << "}" << std::endl; + *streamPtr.get() << tab << return_k << "}" << endl; } - void generateDeviceFunctions(std::unique_ptr& streamPtr) { + void generateDeviceFunctions(unique_ptr& streamPtr) { unsigned int countUnsupported = 0; unsigned int countSupported = 0; - std::stringstream sSupported; - std::stringstream sUnsupported; + stringstream sSupported; + stringstream sUnsupported; for (auto& ma : CUDA_DEVICE_FUNC_MAP) { bool isUnsupported = Statistics::isUnsupported(ma.second); - (isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << double_tab << "\"" << ma.first.str() << "\""; + (isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << tab_2 << "\"" << ma.first.str() << "\""; if (isUnsupported) countUnsupported++; else countSupported++; } - std::stringstream subCountSupported; - std::stringstream subWarnUnsupported; - std::stringstream subCommon; - 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 + sMy + "$line_num = shift;\n" + sCommon : tab + sReturn_0); - if (countSupported) { - subCountSupported << sSupported.str() << std::endl << tab << ")" << std::endl; - } - if (countUnsupported) { - subWarnUnsupported << sUnsupported.str() << std::endl << tab << ")" << std::endl; - } + stringstream subCountSupported; + stringstream subWarnUnsupported; + stringstream subCommon; + string sCommon = tab + my_k + "\n" + tab + foreach_func; + subCountSupported << endl << sub << "countSupportedDeviceFunctions" << " {" << endl << (countSupported ? sCommon : tab + return_0); + subWarnUnsupported << endl << sub << "warnUnsupportedDeviceFunctions" << " {" << endl << (countUnsupported ? tab + my + "$line_num = shift;\n" + sCommon : tab + return_0); + if (countSupported) subCountSupported << sSupported.str() << endl_tab << ")" << endl; + if (countUnsupported) subWarnUnsupported << sUnsupported.str() << endl_tab << ")" << endl; if (countSupported || countUnsupported) { - subCommon << tab << "{" << std::endl; - subCommon << double_tab << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));" << std::endl; - subCommon << double_tab << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);" << std::endl; - subCommon << double_tab << sMy << "$mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl; - subCommon << double_tab << sMy << "$mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl; - subCommon << double_tab << "if ($mt && !$mt_namespace) {" << std::endl; - subCommon << triple_tab << "$k += $mt;" << std::endl; - } - if (countSupported) { - subCountSupported << subCommon.str(); + subCommon << tab << "{" << endl; + subCommon << tab_2 << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));" << endl; + subCommon << tab_2 << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);" << endl; + subCommon << tab_2 << my << "$mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;" << endl; + subCommon << tab_2 << my << "$mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;" << endl; + subCommon << tab_2 << "if ($mt && !$mt_namespace) {" << endl; + subCommon << tab_3 << "$k += $mt;" << endl; } + if (countSupported) subCountSupported << subCommon.str(); if (countUnsupported) { subWarnUnsupported << subCommon.str(); - 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_k; + subWarnUnsupported << tab_3 << print << "\" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";" << endl; } + if (countSupported || countUnsupported) sCommon = tab_2 + "}\n" + tab + "}\n" + tab + return_k; if (countSupported) subCountSupported << sCommon; if (countUnsupported) subWarnUnsupported << sCommon; - subCountSupported << "}" << std::endl; - subWarnUnsupported << "}" << std::endl; + subCountSupported << "}" << endl; + subWarnUnsupported << "}" << endl; *streamPtr.get() << subCountSupported.str(); *streamPtr.get() << subWarnUnsupported.str(); } bool generate(bool Generate) { if (!Generate) return true; - std::string dstHipifyPerl = "hipify-perl", dstHipifyPerlDir = OutputHipifyPerlDir; - std::error_code EC; + string dstHipifyPerl = hipify_perl, dstHipifyPerlDir = OutputHipifyPerlDir; + error_code EC; if (!dstHipifyPerlDir.empty()) { - std::string sOutputHipifyPerlDirAbsPath = getAbsoluteDirectoryPath(OutputHipifyPerlDir, EC, "output hipify-perl"); + string sOutputHipifyPerlDirAbsPath = getAbsoluteDirectoryPath(OutputHipifyPerlDir, EC, "output " + hipify_perl); if (EC) return false; dstHipifyPerl = sOutputHipifyPerlDirAbsPath + "/" + dstHipifyPerl; } SmallString<128> tmpFile; - StringRef ext = "hipify-perl-tmp"; - EC = sys::fs::createTemporaryFile(dstHipifyPerl, ext, tmpFile); + EC = sys::fs::createTemporaryFile(dstHipifyPerl, hipify_perl, tmpFile); if (EC) { llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n"; return false; } - std::unique_ptr streamPtr = std::unique_ptr(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc)); + unique_ptr streamPtr = unique_ptr(new ofstream(tmpFile.c_str(), ios_base::trunc)); generateHeader(streamPtr); - std::string sConv = sMy + "$apiCalls = "; + string sConv = my + "$apiCalls = "; unsigned int exclude[3] = { CONV_DEVICE_FUNC, CONV_EXTERN_SHARED, CONV_KERNEL_LAUNCH }; *streamPtr.get() << "@statNames = ("; for (unsigned int i = 0; i < NUM_CONV_TYPES - 1; ++i) { *streamPtr.get() << "\"" << counterNames[i] << "\", "; - if (std::any_of(exclude, exclude + 3, [&i](unsigned int x) { return x == i; })) continue; - sConv += "$ft{'" + std::string(counterNames[i]) + "'}" + (i < NUM_CONV_TYPES - 2 ? " + " : ";"); + if (any_of(exclude, exclude + 3, [&i](unsigned int x) { return x == i; })) continue; + sConv += "$ft{'" + string(counterNames[i]) + "'}" + (i < NUM_CONV_TYPES - 2 ? " + " : ";"); } - if (sConv.back() == ' ') { - sConv = sConv.substr(0, sConv.size() - 3) + ";"; - } - *streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << std::endl; + if (sConv.back() == ' ') sConv = sConv.substr(0, sConv.size() - 3) + ";"; + *streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << endl; generateStatFunctions(streamPtr); generateSimpleSubstitutions(streamPtr); generateExternShared(streamPtr); generateKernelLaunch(streamPtr); generateHostFunctions(streamPtr); generateDeviceFunctions(streamPtr); - *streamPtr.get() << std::endl << "# Count of transforms in all files" << std::endl; - *streamPtr.get() << sMy << "%tt;" << std::endl; - *streamPtr.get() << "clearStats(\\%tt, \\@statNames);" << std::endl; - *streamPtr.get() << "$Twarnings = 0;" << std::endl; - *streamPtr.get() << "$TlineCount = 0;" << std::endl; - *streamPtr.get() << sMy << "%TwarningTags;" << std::endl; - *streamPtr.get() << sMy << "$fileCount = @ARGV;" << std::endl << std::endl; - *streamPtr.get() << "while (@ARGV) {" << std::endl; - *streamPtr.get() << tab << "$fileName=shift (@ARGV);" << std::endl; - *streamPtr.get() << tab << "if ($inplace) {" << std::endl; - *streamPtr.get() << double_tab << sMy << "$file_prehip = \"$fileName\" . \".prehip\";" << std::endl; - *streamPtr.get() << double_tab << sMy << "$infile;" << std::endl; - *streamPtr.get() << double_tab << sMy << "$outfile;" << std::endl; - *streamPtr.get() << double_tab << "if (-e $file_prehip) {" << std::endl; - *streamPtr.get() << triple_tab << "$infile = $file_prehip;" << std::endl; - *streamPtr.get() << triple_tab << "$outfile = $fileName;" << std::endl; - *streamPtr.get() << double_tab << "} else {" << std::endl; - *streamPtr.get() << triple_tab << "system (\"cp $fileName $file_prehip\");" << std::endl; - *streamPtr.get() << triple_tab << "$infile = $file_prehip;" << std::endl; - *streamPtr.get() << triple_tab << "$outfile = $fileName;" << std::endl << double_tab << "}" << std::endl; - *streamPtr.get() << double_tab << "open(INFILE,\"<\", $infile) or die \"error: could not open $infile\";" << std::endl; - *streamPtr.get() << double_tab << "open(OUTFILE,\">\", $outfile) or die \"error: could not open $outfile\";" << std::endl; - *streamPtr.get() << double_tab << "$OUTFILE = OUTFILE;" << std::endl; - *streamPtr.get() << tab << "} else {" << std::endl; - *streamPtr.get() << double_tab << "open(INFILE,\"<\", $fileName) or die \"error: could not open $fileName\";" << std::endl; - *streamPtr.get() << double_tab << "$OUTFILE = STDOUT;" << std::endl << tab << "}" << std::endl; - *streamPtr.get() << tab << "# Count of transforms in this file" << std::endl; - *streamPtr.get() << tab << "clearStats(\\%ft, \\@statNames);" << std::endl; - *streamPtr.get() << tab << sMy << "$countIncludes = 0;" << std::endl; - *streamPtr.get() << tab << sMy << "$countKeywords = 0;" << std::endl; - *streamPtr.get() << tab << sMy << "$warnings = 0;" << std::endl; - *streamPtr.get() << tab << sMy << "%warningTags;" << std::endl; - *streamPtr.get() << tab << sMy << "$lineCount = 0;" << std::endl; - *streamPtr.get() << tab << "undef $/;" << std::endl; - *streamPtr.get() << tab << "# Read whole file at once, so we can match newlines" << std::endl; - *streamPtr.get() << tab << "while () {" << std::endl; - *streamPtr.get() << double_tab << "$countKeywords += m/__global__/;" << std::endl; - *streamPtr.get() << double_tab << "$countKeywords += m/__shared__/;" << std::endl; - *streamPtr.get() << double_tab << "simpleSubstitutions();" << std::endl; - *streamPtr.get() << double_tab << "transformExternShared();" << std::endl; - *streamPtr.get() << double_tab << "transformKernelLaunch();" << std::endl; - *streamPtr.get() << double_tab << "if ($print_stats) {" << std::endl; - *streamPtr.get() << triple_tab << "while (/(\\b(hip|HIP)([A-Z]|_)\\w+\\b)/g) {" << std::endl; - *streamPtr.get() << quad_tab << "$convertedTags{$1}++;" << std::endl; - *streamPtr.get() << triple_tab << "}" << std::endl; - *streamPtr.get() << double_tab << "}" << std::endl; - *streamPtr.get() << double_tab << sMy << "$hasDeviceCode = $countKeywords + $ft{'device_function'};" << std::endl; - *streamPtr.get() << double_tab << "unless ($quiet_warnings) {" << std::endl; - *streamPtr.get() << triple_tab << "# Copy into array of lines, process line-by-line to show warnings" << std::endl; - *streamPtr.get() << triple_tab << "if ($hasDeviceCode or (/\\bcu|CU/) or (/<<<.*>>>/)) {" << std::endl; - *streamPtr.get() << quad_tab << sMy << "@lines = split /\\n/, $_;" << std::endl; - *streamPtr.get() << quad_tab << "# Copy the whole file" << std::endl; - *streamPtr.get() << quad_tab << sMy << "$tmp = $_;" << std::endl; - *streamPtr.get() << quad_tab << sMy << "$line_num = 0;" << std::endl; - *streamPtr.get() << quad_tab << "foreach (@lines) {" << std::endl; - *streamPtr.get() << tab_5 << "$line_num++;" << std::endl; - *streamPtr.get() << tab_5 << "# Remove any whitelisted words" << std::endl; - *streamPtr.get() << tab_5 << "foreach $w (@whitelist) {" << std::endl; - *streamPtr.get() << tab_6 << "s/\\b$w\\b/ZAP/" << std::endl; - *streamPtr.get() << tab_5 << "}" << std::endl; - *streamPtr.get() << tab_5 << sMy << "$tag;" << std::endl; - *streamPtr.get() << tab_5 << "if ((/(\\bcuda[A-Z]\\w+)/) or (/<<<.*>>>/)) {" << std::endl; - *streamPtr.get() << tab_6 << "# Flag any remaining code that look like cuda API calls: may want to add these to hipify" << std::endl; - *streamPtr.get() << tab_6 << "$tag = (defined $1) ? $1 : \"Launch\";" << std::endl; - *streamPtr.get() << tab_5 << "}" << std::endl; - *streamPtr.get() << tab_5 << "if (defined $tag) {" << std::endl; - *streamPtr.get() << tab_6 << "$warnings++;" << std::endl; - *streamPtr.get() << tab_6 << "$warningTags{$tag}++;" << std::endl; - *streamPtr.get() << tab_6 << "print STDERR \" warning: $fileName:#$line_num : $_\\n\";" << std::endl; - *streamPtr.get() << tab_5 << "}" << std::endl; - *streamPtr.get() << tab_5 << "$s = warnUnsupportedDeviceFunctions($line_num);" << std::endl; - *streamPtr.get() << tab_5 << "$warnings += $s;" << std::endl; - *streamPtr.get() << quad_tab << "}" << std::endl; - *streamPtr.get() << quad_tab << "$_ = $tmp;" << std::endl; - *streamPtr.get() << triple_tab << "}" << std::endl; - *streamPtr.get() << double_tab << "}" << std::endl; - *streamPtr.get() << double_tab << "if ($hasDeviceCode > 0) {" << std::endl; - *streamPtr.get() << triple_tab << "$ft{'device_function'} += countSupportedDeviceFunctions();" << std::endl; - *streamPtr.get() << double_tab << "}" << std::endl; - *streamPtr.get() << double_tab << "transformHostFunctions();" << std::endl; - *streamPtr.get() << double_tab << "# TODO: would like to move this code outside loop but it uses $_ which contains the whole file" << std::endl; - *streamPtr.get() << double_tab << "unless ($no_output) {" << std::endl; - *streamPtr.get() << triple_tab << sConv << std::endl; - *streamPtr.get() << triple_tab << sMy << "$kernStuff = $hasDeviceCode + $ft{'" << counterNames[CONV_KERNEL_LAUNCH] << "'} + $ft{'" << counterNames[CONV_DEVICE_FUNC] << "'};" << std::endl; - *streamPtr.get() << triple_tab << sMy << "$totalCalls = $apiCalls + $kernStuff;" << std::endl; - *streamPtr.get() << triple_tab << "$is_dos = m/\\r\\n$/;" << std::endl; - *streamPtr.get() << triple_tab << "if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) {" << std::endl; - *streamPtr.get() << quad_tab << "# TODO: implement hipify-clang's logic with header files AMAP" << std::endl; - *streamPtr.get() << quad_tab << "print $OUTFILE '#include \"hip/hip_runtime.h\"' . ($is_dos ? \"\\r\\n\" : \"\\n\");" << std::endl; - *streamPtr.get() << triple_tab << "}" << std::endl; - *streamPtr.get() << triple_tab << "print $OUTFILE \"$_\";" << std::endl; - *streamPtr.get() << double_tab << "}" << std::endl; - *streamPtr.get() << double_tab << "$lineCount = $_ =~ tr/\\n//;" << std::endl; - *streamPtr.get() << tab << "}" << std::endl; - *streamPtr.get() << tab << sMy << "$totalConverted = totalStats(\\%ft);" << std::endl; - *streamPtr.get() << tab << "if (($totalConverted+$warnings) and $print_stats) {" << std::endl; - *streamPtr.get() << double_tab << "printStats(\" info: converted\", \\@statNames, \\%ft, $warnings, $lineCount);" << std::endl; - *streamPtr.get() << double_tab << "print STDERR \" in '$fileName'\\n\";" << std::endl; - *streamPtr.get() << tab << "}" << std::endl; - *streamPtr.get() << tab << "# Update totals for all files" << std::endl; - *streamPtr.get() << tab << "addStats(\\%tt, \\%ft);" << std::endl; - *streamPtr.get() << tab << "$Twarnings += $warnings;" << std::endl; - *streamPtr.get() << tab << "$TlineCount += $lineCount;" << std::endl; - *streamPtr.get() << tab << "foreach $key (keys %warningTags) {" << std::endl; - *streamPtr.get() << double_tab << "$TwarningTags{$key} += $warningTags{$key};" << std::endl; - *streamPtr.get() << tab << "}" << std::endl; - *streamPtr.get() << "}" << std::endl; - *streamPtr.get() << "# Print total stats for all files processed:" << std::endl; - *streamPtr.get() << "if ($print_stats and ($fileCount > 1)) {" << std::endl; - *streamPtr.get() << tab << "print STDERR \"\\n\";" << std::endl; - *streamPtr.get() << tab << "printStats(\" info: TOTAL-converted\", \\@statNames, \\%tt, $Twarnings, $TlineCount);" << std::endl; - *streamPtr.get() << tab << "print STDERR \"\\n\";" << std::endl; - *streamPtr.get() << tab << "foreach my $key (sort { $TwarningTags{$b} <=> $TwarningTags{$a} } keys %TwarningTags) {" << std::endl; - *streamPtr.get() << double_tab << "printf STDERR \" warning: unconverted %s : %d\\n\", $key, $TwarningTags{$key};" << std::endl; - *streamPtr.get() << tab << "}" << std::endl; - *streamPtr.get() << tab << sMy << "$kernelCnt = keys %Tkernels;" << std::endl; - *streamPtr.get() << tab << "printf STDERR \" kernels (%d total) : \", $kernelCnt;" << std::endl; - *streamPtr.get() << tab << "foreach my $key (sort { $Tkernels{$b} <=> $Tkernels{$a} } keys %Tkernels) {" << std::endl; - *streamPtr.get() << double_tab << "printf STDERR \" %s(%d)\", $key, $Tkernels{$key};" << std::endl; - *streamPtr.get() << tab << "}" << std::endl; - *streamPtr.get() << tab << "print STDERR \"\\n\\n\";" << std::endl; - *streamPtr.get() << "}" << std::endl; - *streamPtr.get() << "if ($print_stats) {" << std::endl; - *streamPtr.get() << tab << "foreach my $key (sort { $convertedTags{$b} <=> $convertedTags{$a} } keys %convertedTags) {" << std::endl; - *streamPtr.get() << double_tab << "printf STDERR \" %s %d\\n\", $key, $convertedTags{$key};" << std::endl; - *streamPtr.get() << tab << "}" << std::endl << "}" << std::endl; + *streamPtr.get() << endl << "# Count of transforms in all files" << endl; + *streamPtr.get() << my << "%tt;" << endl; + *streamPtr.get() << "clearStats(\\%tt, \\@statNames);" << endl; + *streamPtr.get() << "$Twarnings = 0;" << endl; + *streamPtr.get() << "$TlineCount = 0;" << endl; + *streamPtr.get() << my << "%TwarningTags;" << endl; + *streamPtr.get() << my << "$fileCount = @ARGV;" << endl_2; + *streamPtr.get() << while_ << "(@ARGV) {" << endl; + *streamPtr.get() << tab << "$fileName=shift (@ARGV);" << endl; + *streamPtr.get() << tab << "if ($inplace) {" << endl; + *streamPtr.get() << tab_2 << my << "$file_prehip = \"$fileName\" . \".prehip\";" << endl; + *streamPtr.get() << tab_2 << my << "$infile;" << endl; + *streamPtr.get() << tab_2 << my << "$outfile;" << endl; + *streamPtr.get() << tab_2 << "if (-e $file_prehip) {" << endl; + *streamPtr.get() << tab_3 << "$infile = $file_prehip;" << endl; + *streamPtr.get() << tab_3 << "$outfile = $fileName;" << endl; + *streamPtr.get() << tab_2 << "} else {" << endl; + *streamPtr.get() << tab_3 << "system (\"cp $fileName $file_prehip\");" << endl; + *streamPtr.get() << tab_3 << "$infile = $file_prehip;" << endl; + *streamPtr.get() << tab_3 << "$outfile = $fileName;" << endl_tab_2 << "}" << endl; + *streamPtr.get() << tab_2 << "open(INFILE,\"<\", $infile) or die \"error: could not open $infile\";" << endl; + *streamPtr.get() << tab_2 << "open(OUTFILE,\">\", $outfile) or die \"error: could not open $outfile\";" << endl; + *streamPtr.get() << tab_2 << "$OUTFILE = OUTFILE;" << endl; + *streamPtr.get() << tab << "} else {" << endl; + *streamPtr.get() << tab_2 << "open(INFILE,\"<\", $fileName) or die \"error: could not open $fileName\";" << endl; + *streamPtr.get() << tab_2 << "$OUTFILE = STDOUT;" << endl_tab << "}" << endl; + *streamPtr.get() << tab << "# Count of transforms in this file" << endl; + *streamPtr.get() << tab << "clearStats(\\%ft, \\@statNames);" << endl; + *streamPtr.get() << tab << my << "$countIncludes = 0;" << endl; + *streamPtr.get() << tab << my << "$countKeywords = 0;" << endl; + *streamPtr.get() << tab << my << "$warnings = 0;" << endl; + *streamPtr.get() << tab << my << "%warningTags;" << endl; + *streamPtr.get() << tab << my << "$lineCount = 0;" << endl; + *streamPtr.get() << tab << "undef $/;" << endl; + *streamPtr.get() << tab << "# Read whole file at once, so we can match newlines" << endl; + *streamPtr.get() << tab << while_ << "() {" << endl; + *streamPtr.get() << tab_2 << "$countKeywords += m/__global__/;" << endl; + *streamPtr.get() << tab_2 << "$countKeywords += m/__shared__/;" << endl; + *streamPtr.get() << tab_2 << "simpleSubstitutions();" << endl; + *streamPtr.get() << tab_2 << "transformExternShared();" << endl; + *streamPtr.get() << tab_2 << "transformKernelLaunch();" << endl; + *streamPtr.get() << tab_2 << "if ($print_stats) {" << endl; + *streamPtr.get() << tab_3 << while_ << "(/(\\b(hip|HIP)([A-Z]|_)\\w+\\b)/g) {" << endl; + *streamPtr.get() << tab_4 << "$convertedTags{$1}++;" << endl_tab_3 << "}" << endl_tab_2 << "}" << endl; + *streamPtr.get() << tab_2 << my << "$hasDeviceCode = $countKeywords + $ft{'device_function'};" << endl; + *streamPtr.get() << tab_2 << unless_ << "($quiet_warnings) {" << endl; + *streamPtr.get() << tab_3 << "# Copy into array of lines, process line-by-line to show warnings" << endl; + *streamPtr.get() << tab_3 << "if ($hasDeviceCode or (/\\bcu|CU/) or (/<<<.*>>>/)) {" << endl; + *streamPtr.get() << tab_4 << my << "@lines = split /\\n/, $_;" << endl; + *streamPtr.get() << tab_4 << "# Copy the whole file" << endl; + *streamPtr.get() << tab_4 << my << "$tmp = $_;" << endl; + *streamPtr.get() << tab_4 << my << "$line_num = 0;" << endl; + *streamPtr.get() << tab_4 << foreach << "(@lines) {" << endl; + *streamPtr.get() << tab_5 << "$line_num++;" << endl; + *streamPtr.get() << tab_5 << "# Remove any whitelisted words" << endl; + *streamPtr.get() << tab_5 << foreach << "$w (@whitelist) {" << endl; + *streamPtr.get() << tab_6 << "s/\\b$w\\b/ZAP/" << endl_tab_5 << "}" << endl; + *streamPtr.get() << tab_5 << my << "$tag;" << endl; + *streamPtr.get() << tab_5 << "if ((/(\\bcuda[A-Z]\\w+)/) or (/<<<.*>>>/)) {" << endl; + *streamPtr.get() << tab_6 << "# Flag any remaining code that look like cuda API calls: may want to add these to hipify" << endl; + *streamPtr.get() << tab_6 << "$tag = (defined $1) ? $1 : \"Launch\";" << endl_tab_5 << "}" << endl; + *streamPtr.get() << tab_5 << "if (defined $tag) {" << endl; + *streamPtr.get() << tab_6 << "$warnings++;" << endl; + *streamPtr.get() << tab_6 << "$warningTags{$tag}++;" << endl; + *streamPtr.get() << tab_6 << print << "\" warning: $fileName:#$line_num : $_\\n\";" << endl_tab_5 << "}" << endl; + *streamPtr.get() << tab_5 << "$s = warnUnsupportedDeviceFunctions($line_num);" << endl; + *streamPtr.get() << tab_5 << "$warnings += $s;" << endl_tab_4 << "}" << endl; + *streamPtr.get() << tab_4 << "$_ = $tmp;" << endl_tab_3 << "}" << endl_tab_2 << "}" << endl; + *streamPtr.get() << tab_2 << "if ($hasDeviceCode > 0) {" << endl; + *streamPtr.get() << tab_3 << "$ft{'device_function'} += countSupportedDeviceFunctions();" << endl_tab_2 << "}" << endl; + *streamPtr.get() << tab_2 << "transformHostFunctions();" << endl; + *streamPtr.get() << tab_2 << "# TODO: would like to move this code outside loop but it uses $_ which contains the whole file" << endl; + *streamPtr.get() << tab_2 << unless_ << "($no_output) {" << endl; + *streamPtr.get() << tab_3 << sConv << endl; + *streamPtr.get() << tab_3 << my << "$kernStuff = $hasDeviceCode + $ft{'" << counterNames[CONV_KERNEL_LAUNCH] << "'} + $ft{'" << counterNames[CONV_DEVICE_FUNC] << "'};" << endl; + *streamPtr.get() << tab_3 << my << "$totalCalls = $apiCalls + $kernStuff;" << endl; + *streamPtr.get() << tab_3 << "$is_dos = m/\\r\\n$/;" << endl; + *streamPtr.get() << tab_3 << "if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) {" << endl; + *streamPtr.get() << tab_4 << "# TODO: implement hipify-clang's logic with header files AMAP" << endl; + *streamPtr.get() << tab_4 << "print $OUTFILE '#include \"hip/hip_runtime.h\"' . ($is_dos ? \"\\r\\n\" : \"\\n\");" << endl_tab_3 << "}" << endl; + *streamPtr.get() << tab_3 << "print $OUTFILE \"$_\";" << endl_tab_2 << "}" << endl; + *streamPtr.get() << tab_2 << "$lineCount = $_ =~ tr/\\n//;" << endl_tab << "}" << endl; + *streamPtr.get() << tab << my << "$totalConverted = totalStats(\\%ft);" << endl; + *streamPtr.get() << tab << "if (($totalConverted+$warnings) and $print_stats) {" << endl; + *streamPtr.get() << tab_2 << "printStats(\" info: converted\", \\@statNames, \\%ft, $warnings, $lineCount);" << endl; + *streamPtr.get() << tab_2 << print << "\" in '$fileName'\\n\";" << endl_tab << "}" << endl; + *streamPtr.get() << tab << "# Update totals for all files" << endl; + *streamPtr.get() << tab << "addStats(\\%tt, \\%ft);" << endl; + *streamPtr.get() << tab << "$Twarnings += $warnings;" << endl; + *streamPtr.get() << tab << "$TlineCount += $lineCount;" << endl; + *streamPtr.get() << tab << foreach << "$key (keys %warningTags) {" << endl; + *streamPtr.get() << tab_2 << "$TwarningTags{$key} += $warningTags{$key};" << endl_tab << "}" << endl << "}" << endl; + *streamPtr.get() << "# Print total stats for all files processed:" << endl; + *streamPtr.get() << "if ($print_stats and ($fileCount > 1)) {" << endl; + *streamPtr.get() << tab << print << "\"\\n\";" << endl; + *streamPtr.get() << tab << "printStats(\" info: TOTAL-converted\", \\@statNames, \\%tt, $Twarnings, $TlineCount);" << endl; + *streamPtr.get() << tab << print << "\"\\n\";" << endl; + *streamPtr.get() << tab << foreach << "my $key (sort { $TwarningTags{$b} <=> $TwarningTags{$a} } keys %TwarningTags) {" << endl; + *streamPtr.get() << tab_2 << printf << "\" warning: unconverted %s : %d\\n\", $key, $TwarningTags{$key};" << endl_tab << "}" << endl; + *streamPtr.get() << tab << my << "$kernelCnt = keys %Tkernels;" << endl; + *streamPtr.get() << tab << printf << "\" kernels (%d total) : \", $kernelCnt;" << endl; + *streamPtr.get() << tab << foreach << "my $key (sort { $Tkernels{$b} <=> $Tkernels{$a} } keys %Tkernels) {" << endl; + *streamPtr.get() << tab_2 << printf << "\" %s(%d)\", $key, $Tkernels{$key};" << endl_tab << "}" << endl; + *streamPtr.get() << tab << print << "\"\\n\\n\";" << endl << "}" << endl; + *streamPtr.get() << "if ($print_stats) {" << endl; + *streamPtr.get() << tab << foreach << "my $key (sort { $convertedTags{$b} <=> $convertedTags{$a} } keys %convertedTags) {" << endl; + *streamPtr.get() << tab_2 << printf << "\" %s %d\\n\", $key, $convertedTags{$key};" << endl_tab << "}" << endl << "}" << endl; streamPtr.get()->flush(); bool ret = true; EC = sys::fs::copy_file(tmpFile, dstHipifyPerl);