From de6c26040d9e7e63caba64bcd34bb3ab6668ce1a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 3 Oct 2019 16:00:52 +0300 Subject: [PATCH 01/17] [HIPIFY] Finishing with generation of hipify-perl + hipify-perl script is entirely generated by hipify-clang under an option -perl now + hipify-perl still has correctness gaps comparing to hipify-clang: they will be eliminated AMAP further [ROCm/hip commit: 6badebd963d92f37b9c84997a4d1bcd642ec65bf] --- projects/hip/bin/hipify-perl | 34 ++-- .../hip/hipify-clang/src/CUDA2HIP_Perl.cpp | 179 ++++++++++++++++-- 2 files changed, 171 insertions(+), 42 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 609b0fc3f4..225bb08887 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -27,6 +27,8 @@ use Getopt::Long; my $whitelist = ""; my $fileName = ""; +my %ft; +my %Tkernels; GetOptions( "examine" => \$examine # Combines -no-output and -print-stats options. @@ -1600,12 +1602,11 @@ sub transformExternShared { no warnings qw/uninitialized/; my $k = 0; $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g; - $ft{ 'extern_shared' } += $k; + $ft{'extern_shared'} += $k; } # CUDA Kernel Launch Syntax sub transformKernelLaunch { - my $TkernRef = @_; no warnings qw/uninitialized/; my $k = 0; @@ -1641,7 +1642,7 @@ sub transformKernelLaunch { if ($k) { $ft{'kernel_launch'} += $k; - @$TkernRef{$1} ++; + $Tkernels{$1}++; } } @@ -2299,7 +2300,6 @@ clearStats(\%tt, \@statNames); $Twarnings = 0; $TlineCount = 0; my %TwarningTags; -my %Tkernels; my $fileCount = @ARGV; while (@ARGV) { @@ -2323,42 +2323,36 @@ while (@ARGV) { open(INFILE,"<", $fileName) or die "error: could not open $fileName"; $OUTFILE = STDOUT; } - # Note : \b is used in perl to indicate the start of a word # Count of transforms in this file - my %ft; clearStats(\%ft, \@statNames); my $countIncludes = 0; my $countKeywords = 0; - my $warnings = 0; + my $warnings = 0; my %warningTags; my $lineCount = 0; undef $/; # Read whole file at once, so we can match newlines - while () - { - simpleSubstitutions(); - + while () { $countKeywords += m/__global__/; $countKeywords += m/__shared__/; - + simpleSubstitutions(); transformExternShared(); - transformKernelLaunch(\%Tkernels); - + transformKernelLaunch(); if ($print_stats) { - while (/(\bhip[A-Z]\w+\b)/g) { + while (/(\b(hip|HIP)([A-Z]|_)\w+\b)/g) { $convertedTags{$1}++; } } my $hasDeviceCode = $countKeywords + $ft{'device_function'}; unless ($quiet_warnings) { # Copy into array of lines, process line-by-line to show warnings - if ($hasDeviceCode or (/\bcu/) or (/\bCU_/) or (/\bCUDA_/) or (/<<<.*>>>/)) { + if ($hasDeviceCode or (/\bcu|CU/) or (/<<<.*>>>/)) { my @lines = split /\n/, $_; # Copy the whole file my $tmp = $_; my $line_num = 0; foreach (@lines) { - $line_num ++; + $line_num++; # Remove any whitelisted words foreach $w (@whitelist) { s/\b$w\b/ZAP/ @@ -2371,8 +2365,7 @@ while (@ARGV) { if (defined $tag) { $warnings++; $warningTags{$tag}++; - print STDERR " warning: $fileName:#$line_num : $_"; - print STDERR "\n"; + print STDERR " warning: $fileName:#$line_num : $_\n"; } $s = warnUnsupportedDeviceFunctions($line_num); $warnings += $s; @@ -2424,8 +2417,7 @@ if ($print_stats and ($fileCount > 1)) { foreach my $key (sort { $Tkernels{$b} <=> $Tkernels{$a} } keys %Tkernels) { printf STDERR " %s(%d)", $key, $Tkernels{$key}; } - print STDERR "\n"; - print STDERR "\n"; + print STDERR "\n\n"; } if ($print_stats) { foreach my $key (sort { $convertedTags{$b} <=> $convertedTags{$a} } keys %convertedTags) { diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index f101f9b8cb..675826824a 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -62,11 +62,15 @@ namespace perl { 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_k = "my $k = 0;"; + const std::string sMy = "my "; + const std::string sMy_k = sMy + "$k = 0;"; const std::string sNoWarns = "no warnings qw/uninitialized/;"; const std::string sCudaDevice = "cudaDevice"; @@ -104,7 +108,10 @@ namespace perl { *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() << "my $whitelist = \"\";" << std::endl << 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; @@ -128,18 +135,18 @@ namespace perl { void generateStatFunctions(std::unique_ptr& streamPtr) { *streamPtr.get() << std::endl << sSub << " totalStats" << " {" << std::endl; - *streamPtr.get() << tab << "my %count = %{ shift() };" << std::endl; - *streamPtr.get() << tab << "my $total = 0;" << 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 << "my $label = shift();" << std::endl; - *streamPtr.get() << tab << "my @statNames = @{ shift() };" << std::endl; - *streamPtr.get() << tab << "my %counts = %{ shift() };" << std::endl; - *streamPtr.get() << tab << "my $warnings = shift();" << std::endl; - *streamPtr.get() << tab << "my $loc = shift();" << std::endl; - *streamPtr.get() << tab << "my $total = totalStats(\\%counts);" << 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; @@ -147,8 +154,8 @@ namespace perl { *streamPtr.get() << tab << "printf STDERR \")\\n warn:%d LOC:%d\", $warnings, $loc;" << std::endl << "}" << std::endl; for (int i = 0; i < 2; ++i) { *streamPtr.get() << std::endl << sSub << " " << (i ? "clearStats" : "addStats") << " {" << std::endl; - *streamPtr.get() << tab << "my $dest_ref = shift();" << std::endl; - *streamPtr.get() << tab << (i ? "my @statNames = @{ shift() };" : "my %adder = %{ shift() };") << 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; } @@ -186,13 +193,12 @@ namespace perl { *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; + *streamPtr.get() << tab << "$ft{'extern_shared'} += $k;" << std::endl << "}" << std::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 << "my $TkernRef = @_;" << std::endl; *streamPtr.get() << tab << sNoWarns << std::endl; *streamPtr.get() << tab << sMy_k << std::endl << std::endl; @@ -228,7 +234,7 @@ namespace perl { *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; + *streamPtr.get() << double_tab << "$Tkernels{$1}++;" << std::endl << tab << "}" << std::endl << "}" << std::endl; } void generateHostFunctions(std::unique_ptr& streamPtr) { @@ -285,7 +291,7 @@ namespace perl { 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 + "my $line_num = shift;\n" + 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; } @@ -296,8 +302,8 @@ namespace perl { 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 << "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 << 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; } @@ -338,11 +344,16 @@ namespace perl { } std::unique_ptr streamPtr = std::unique_ptr(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc)); generateHeader(streamPtr); - std::string sConv = "my $conversions = "; + std::string sConv = sMy + "$apiCalls = "; + unsigned int exclude[3] = { CONV_DEVICE_FUNC, CONV_EXTERN_SHARED, CONV_KERNEL_LAUNCH }; *streamPtr.get() << "@statNames = ("; - for (int i = 0; i < NUM_CONV_TYPES - 1; ++i) { + for (unsigned int i = 0; i < NUM_CONV_TYPES - 1; ++i) { *streamPtr.get() << "\"" << counterNames[i] << "\", "; - sConv += "$ft{'" + std::string(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 (sConv.back() == ' ') { + sConv = sConv.substr(0, sConv.size() - 3) + ";"; } *streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << std::endl; generateStatFunctions(streamPtr); @@ -351,6 +362,132 @@ namespace perl { 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()->flush(); bool ret = true; EC = sys::fs::copy_file(tmpFile, dstPerlMap); From a7edb7fb157ed820c683ff1b4be11d2a73f815ba Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 3 Oct 2019 16:30:48 +0300 Subject: [PATCH 02/17] [HIPIFY][perl] Start to generate hipify-perl under an option -perl + Get rid of --o-perl-map option for Perl output filename: the output name now is always hipify-perl + Option --o-perl-map-dir renamed to --o-hipify-perl-dir Cmd example for generating hipify-perl script to the particular directory: hipify-clang -perl --o-hipify-perl-dir=../../bin [ROCm/hip commit: b3cf982ee7a12e3050f9ae93c42f8c1a3d340124] --- projects/hip/hipify-clang/src/ArgParse.cpp | 9 ++------- projects/hip/hipify-clang/src/ArgParse.h | 3 +-- projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp | 17 ++++++++--------- 3 files changed, 11 insertions(+), 18 deletions(-) diff --git a/projects/hip/hipify-clang/src/ArgParse.cpp b/projects/hip/hipify-clang/src/ArgParse.cpp index 751de0d028..5ce92dfdca 100644 --- a/projects/hip/hipify-clang/src/ArgParse.cpp +++ b/projects/hip/hipify-clang/src/ArgParse.cpp @@ -44,13 +44,8 @@ cl::opt GeneratePython("python", cl::value_desc("python"), cl::cat(ToolTemplateCategory)); -cl::opt OutputPerlMapFilename("o-perl-map", - cl::desc("Output filename for Perl map"), - cl::value_desc("filename"), - cl::cat(ToolTemplateCategory)); - -cl::opt OutputPerlMapDir("o-perl-map-dir", - cl::desc("Output direcory for Perl map"), +cl::opt OutputHipifyPerlDir("o-hipify-perl-dir", + cl::desc("Output direcory for hipify-perl script"), cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); diff --git a/projects/hip/hipify-clang/src/ArgParse.h b/projects/hip/hipify-clang/src/ArgParse.h index 64fb2e9678..886a658d78 100644 --- a/projects/hip/hipify-clang/src/ArgParse.h +++ b/projects/hip/hipify-clang/src/ArgParse.h @@ -30,8 +30,7 @@ namespace ct = clang::tooling; extern cl::OptionCategory ToolTemplateCategory; extern cl::opt OutputFilename; -extern cl::opt OutputPerlMapFilename; -extern cl::opt OutputPerlMapDir; +extern cl::opt OutputHipifyPerlDir; extern cl::opt OutputPythonMapDir; extern cl::opt OutputDir; extern cl::opt TemporaryDir; diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index 675826824a..f74d196cac 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -327,17 +327,16 @@ namespace perl { bool generate(bool Generate) { if (!Generate) return true; - std::string dstPerlMap = OutputPerlMapFilename, dstPerlMapDir = OutputPerlMapDir; - if (dstPerlMap.empty()) dstPerlMap = "hipify-perl-map"; + std::string dstHipifyPerl = "hipify-perl", dstHipifyPerlDir = OutputHipifyPerlDir; std::error_code EC; - if (!dstPerlMapDir.empty()) { - std::string sOutputPerlMapDirAbsPath = getAbsoluteDirectoryPath(OutputPerlMapDir, EC, "output hipify-perl map"); + if (!dstHipifyPerlDir.empty()) { + std::string sOutputHipifyPerlDirAbsPath = getAbsoluteDirectoryPath(OutputHipifyPerlDir, EC, "output hipify-perl"); if (EC) return false; - dstPerlMap = sOutputPerlMapDirAbsPath + "/" + dstPerlMap; + dstHipifyPerl = sOutputHipifyPerlDirAbsPath + "/" + dstHipifyPerl; } SmallString<128> tmpFile; - StringRef ext = "hipify-tmp"; - EC = sys::fs::createTemporaryFile(dstPerlMap, ext, tmpFile); + StringRef ext = "hipify-perl-tmp"; + EC = sys::fs::createTemporaryFile(dstHipifyPerl, ext, tmpFile); if (EC) { llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n"; return false; @@ -490,9 +489,9 @@ namespace perl { *streamPtr.get() << tab << "}" << std::endl << "}" << std::endl; streamPtr.get()->flush(); bool ret = true; - EC = sys::fs::copy_file(tmpFile, dstPerlMap); + EC = sys::fs::copy_file(tmpFile, dstHipifyPerl); if (EC) { - llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstPerlMap << "\n"; + llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstHipifyPerl << "\n"; ret = false; } if (!SaveTemps) sys::fs::remove(tmpFile); From 0f9074b5685d9297f90b962eedec9021d17380a6 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Tue, 1 Oct 2019 17:13:07 -0400 Subject: [PATCH 03/17] fix kernel descriptor bug with code object v3 Change-Id: I9306b2baf36d338e36c5ab1226f74373a61a5ae0 [ROCm/hip commit: dcf70ff9a27c2245bc761aea182f8ac0fb3dcf0b] --- projects/hip/src/hip_module.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 5fe6ce1996..d98b98a378 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -932,10 +932,10 @@ hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t if (!*func) return hipErrorInvalidValue; - auto kernel = find_kernel_by_name(hmod->executable, name, agent); + std::string name_str(name); + auto kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent); if (kernel.handle == 0u) { - std::string name_str(name); name_str.append(".kd"); kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent); } @@ -945,7 +945,7 @@ hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t // TODO: refactor the whole ihipThisThat, which is a mess and yields the // below, due to hipFunction_t being a pointer to ihipModuleSymbol_t. func[0][0] = *static_cast( - Kernel_descriptor{kernel_object(kernel), name, hmod->kernargs[name]}); + Kernel_descriptor{kernel_object(kernel), name_str, hmod->kernargs[name_str]}); return hipSuccess; } From 7dbc574a4b9336afefe019aa976a7c4edad698bc Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 4 Oct 2019 13:35:21 +0530 Subject: [PATCH 04/17] hipcc defaults to code object v3 (#1298) Change-Id: I272ec61090fa504d631fb54371a73dd6fd17a167 [ROCm/hip commit: d39a2a0749e4b3deffd2717dc8209d668d5f9871] --- projects/hip/bin/hipcc | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index c759802e2e..81f394102a 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -337,6 +337,7 @@ my $runCmd = 1; my $buildDeps = 0; my $linkType = 1; my $setLinkType = 0; +my $coFormatv3 = 1; my @options = (); my @inputs = (); @@ -456,6 +457,22 @@ foreach $arg (@ARGV) $swallowArg = 1; } + # code object format parsing + if ($trimarg eq '-mcode-object-v3') { + $coFormatv3 = 1; + # hip-clang already recognizes -mcode-object-v3, so we just pass it on + if ($HIP_PLATFORM eq 'hcc') { + $swallowArg = 1; + } + } + if ($trimarg eq '-mno-code-object-v3') { + $coFormatv3 = 0; + # hip-clang already recognizes -mno-code-object-v3, so we just pass it on + if ($HIP_PLATFORM eq 'hcc') { + $swallowArg = 1; + } + } + if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) { $arg = "--cuda-device-only"; } @@ -831,6 +848,12 @@ if($HIP_PLATFORM eq "hcc" or $HIP_PLATFORM eq "clang"){ } } +# hcc defaults to v2, so we need to convert to the appropriate flag +# hip-clang defaults to v3, so we don't need to do anything +if ($coFormatv3 and $HIP_PLATFORM eq 'hcc') { + $HIPLDFLAGS .= " -Wl,-hcc-cov3 "; +} + if ($hasC and $HIP_PLATFORM eq 'nvcc') { $HIPCXXFLAGS .= " -x cu"; } From 66e94b1b1fa0c29d9c3d1e32c1c119a07a32260e Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Fri, 4 Oct 2019 13:36:31 +0530 Subject: [PATCH 05/17] Added new Memory API's (#1399) Added new memory API's hipMemAllocPitch, hipMemAllocHost, hipMemsetD16, hipMemsetD16Async, hipMemsetD8Async Modified to support all scenarios hipMemcpyParam2DAsync, hipMemcpyParam2D. [ROCm/hip commit: ba9c6e13e438cae8242ccfe47018ed450fec2955] --- projects/hip/bin/hipify-perl | 8 ++ ...A_Driver_API_functions_supported_by_HIP.md | 10 +- .../src/CUDA2HIP_Driver_API_functions.cpp | 16 +-- .../include/hip/hcc_detail/hip_runtime_api.h | 84 +++++++++++++ .../include/hip/nvcc_detail/hip_runtime_api.h | 24 ++++ projects/hip/src/hip_memory.cpp | 116 +++++++++++++++--- .../hip/tests/src/deviceLib/hipTestHalf.cpp | 2 +- .../tests/src/runtimeApi/memory/hipArray.cpp | 47 +++++-- .../tests/src/runtimeApi/memory/hipMemset.cpp | 116 +++++++++++++++++- .../src/runtimeApi/memory/hipMemset2D.cpp | 13 +- projects/hip/tests/src/test_common.cpp | 14 +++ projects/hip/tests/src/test_common.h | 2 + 12 files changed, 404 insertions(+), 48 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 225bb08887..5bf6a4fe9a 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -198,6 +198,8 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcuIpcOpenMemHandle\b/hipIpcOpenMemHandle/g; $ft{'memory'} += s/\bcuMemAlloc\b/hipMalloc/g; $ft{'memory'} += s/\bcuMemAllocManaged\b/hipMemAllocManaged/g; + $ft{'memory'} += s/\bcuMemAllocPitch\b/hipMemAllocPitch/g; + $ft{'memory'} += s/\bcuMemAllocPitch_v2\b/hipMemAllocPitch/g; $ft{'memory'} += s/\bcuMemAlloc_v2\b/hipMalloc/g; $ft{'memory'} += s/\bcuMemFree\b/hipFree/g; $ft{'memory'} += s/\bcuMemFreeHost\b/hipHostFree/g; @@ -207,6 +209,8 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcuMemGetInfo\b/hipMemGetInfo/g; $ft{'memory'} += s/\bcuMemGetInfo_v2\b/hipMemGetInfo/g; $ft{'memory'} += s/\bcuMemHostAlloc\b/hipHostMalloc/g; + $ft{'memory'} += s/\bcuMemAllocHost\b/hipMemAllocHost/g; + $ft{'memory'} += s/\bcuMemAllocHost_v2\b/hipMemAllocHost/g; $ft{'memory'} += s/\bcuMemHostGetDevicePointer\b/hipHostGetDevicePointer/g; $ft{'memory'} += s/\bcuMemHostGetDevicePointer_v2\b/hipHostGetDevicePointer/g; $ft{'memory'} += s/\bcuMemHostGetFlags\b/hipMemHostGetFlags/g; @@ -238,6 +242,10 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcuMemsetD32_v2\b/hipMemsetD32/g; $ft{'memory'} += s/\bcuMemsetD8\b/hipMemsetD8/g; $ft{'memory'} += s/\bcuMemsetD8_v2\b/hipMemsetD8/g; + $ft{'memory'} += s/\bcuMemsetD8Async\b/hipMemsetD8Async/g; + $ft{'memory'} += s/\bcuMemsetD16\b/hipMemsetD16/g; + $ft{'memory'} += s/\bcuMemsetD16_v2\b/hipMemsetD16/g; + $ft{'memory'} += s/\bcuMemsetD16Async\b/hipMemsetD16Async/g; $ft{'memory'} += s/\bcudaFree\b/hipFree/g; $ft{'memory'} += s/\bcudaFreeArray\b/hipFreeArray/g; $ft{'memory'} += s/\bcudaFreeHost\b/hipHostFree/g; diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index cbd4ccb818..5a6ca9fc1e 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -857,9 +857,9 @@ | `cuIpcOpenEventHandle` | | | `cuIpcOpenMemHandle` | `hipIpcOpenMemHandle` | | `cuMemAlloc` | `hipMalloc` | -| `cuMemAllocHost` | | +| `cuMemAllocHost` | `hipMemAllocHost` | | `cuMemAllocManaged` | `hipMemAllocManaged` | -| `cuMemAllocPitch` | | +| `cuMemAllocPitch` | `hipMemAllocPitch` | | `cuMemcpy` | | | `cuMemcpy2D` | `hipMemcpyParam2D` | | `cuMemcpy2DAsync` | `hipMemcpyParam2DAsync` | @@ -893,8 +893,8 @@ | `cuMemHostGetFlags` | `hipHostGetFlags` | | `cuMemHostRegister` | `hipHostRegister` | | `cuMemHostUnregister` | `hipHostUnregister` | -| `cuMemsetD16` | | -| `cuMemsetD16Async` | | +| `cuMemsetD16` | `hipMemsetD16` | +| `cuMemsetD16Async` | `hipMemsetD16Async` | | `cuMemsetD2D16` | | | `cuMemsetD2D16Async` | | | `cuMemsetD2D32` | | @@ -904,7 +904,7 @@ | `cuMemsetD32` | `hipMemsetD32` | | `cuMemsetD32Async` | `hipMemsetD32Async` | | `cuMemsetD8` | `hipMemsetD8` | -| `cuMemsetD8Async` | | +| `cuMemsetD8Async` | `hipMemsetD8Async` | | `cuMipmappedArrayCreate` | | | `cuMipmappedArrayDestroy` | | | `cuMipmappedArrayGetLevel` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 72534a96b6..792d34b3ce 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -162,14 +162,14 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuMemAlloc", {"hipMalloc", "", CONV_MEMORY, API_DRIVER}}, {"cuMemAlloc_v2", {"hipMalloc", "", CONV_MEMORY, API_DRIVER}}, // cudaHostAlloc - {"cuMemAllocHost", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuMemAllocHost_v2", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemAllocHost", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemAllocHost_v2", {"hipMemAllocHost", "", CONV_MEMORY, API_DRIVER}}, // cudaMallocManaged {"cuMemAllocManaged", {"hipMemAllocManaged", "", CONV_MEMORY, API_DRIVER}}, // no analogue // NOTE: Not equal to cudaMallocPitch due to different signatures - {"cuMemAllocPitch", {"hipMemAllocPitch", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuMemAllocPitch_v2", {"hipMemAllocPitch", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemAllocPitch", {"hipMemAllocPitch", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemAllocPitch_v2", {"hipMemAllocPitch", "", CONV_MEMORY, API_DRIVER}}, // no analogue // NOTE: Not equal to cudaMemcpy due to different signatures {"cuMemcpy", {"hipMemcpy_", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, @@ -271,10 +271,10 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // cudaHostUnregister {"cuMemHostUnregister", {"hipHostUnregister", "", CONV_MEMORY, API_DRIVER}}, // no analogue - {"cuMemsetD16", {"hipMemsetD16", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuMemsetD16_v2", {"hipMemsetD16", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemsetD16", {"hipMemsetD16", "", CONV_MEMORY, API_DRIVER}}, + {"cuMemsetD16_v2", {"hipMemsetD16", "", CONV_MEMORY, API_DRIVER}}, // no analogue - {"cuMemsetD16Async", {"hipMemsetD16Async", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemsetD16Async", {"hipMemsetD16Async", "", CONV_MEMORY, API_DRIVER}}, // no analogue {"cuMemsetD2D16", {"hipMemsetD2D16", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, {"cuMemsetD2D16_v2", {"hipMemsetD2D16", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, @@ -299,7 +299,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuMemsetD8", {"hipMemsetD8", "", CONV_MEMORY, API_DRIVER}}, {"cuMemsetD8_v2", {"hipMemsetD8", "", CONV_MEMORY, API_DRIVER}}, // no analogue - {"cuMemsetD8Async", {"hipMemsetD8Async", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemsetD8Async", {"hipMemsetD8Async", "", CONV_MEMORY, API_DRIVER}}, // no analogue // NOTE: Not equal to cudaMallocMipmappedArray due to different signatures {"cuMipmappedArrayCreate", {"hipMipmappedArrayCreate", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index edd63b96f0..1ccc87d216 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1115,6 +1115,21 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag DEPRECATED("use hipHostMalloc instead") hipError_t hipMallocHost(void** ptr, size_t size); +/** + * @brief Allocate pinned host memory [Deprecated] + * + * @param[out] ptr Pointer to the allocated host pinned memory + * @param[in] size Requested memory size + * + * If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned. + * + * @return #hipSuccess, #hipErrorMemoryAllocation + * + * @deprecated use hipHostMalloc() instead + */ +DEPRECATED("use hipHostMalloc instead") +hipError_t hipMemAllocHost(void** ptr, size_t size); + /** * @brief Allocate device accessible page locked host memory * @@ -1250,6 +1265,30 @@ hipError_t hipHostUnregister(void* hostPtr); hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height); +/** + * Allocates at least width (in bytes) * height bytes of linear memory + * Padding may occur to ensure alighnment requirements are met for the given row + * The change in width size due to padding will be returned in *pitch. + * Currently the alignment is set to 128 bytes + * + * @param[out] dptr Pointer to the allocated device memory + * @param[out] pitch Pitch for allocation (in bytes) + * @param[in] width Requested pitched allocation width (in bytes) + * @param[in] height Requested pitched allocation height + * + * If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned. + * The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array. + * Given the row and column of an array element of type T, the address is computed as: + * T* pElement = (T*)((char*)BaseAddress + Row * Pitch) + Column; + * + * @return Error code + * + * @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hipHostFree, hipMalloc3D, + * hipMalloc3DArray, hipHostMalloc + */ + +hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, size_t height, unsigned int elementSizeBytes); + /** * @brief Free memory allocated by the hcc hip memory allocation API. * This API performs an implicit hipDeviceSynchronize() call. @@ -1708,6 +1747,51 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes); */ hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes); +/** + * @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant + * byte value value. + * + * hipMemsetD8Async() is asynchronous with respect to the host, so the call may return before the + * memset is complete. The operation can optionally be associated to a stream by passing a non-zero + * stream argument. If stream is non-zero, the operation may overlap with operations in other + * streams. + * + * @param[out] dst Data ptr to be filled + * @param[in] constant value to be set + * @param[in] sizeBytes Data size in bytes + * @param[in] stream - Stream identifier + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized + */ +hipError_t hipMemsetD8Async(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes, hipStream_t stream __dparm(0)); + +/** + * @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant + * short value value. + * + * @param[out] dst Data ptr to be filled + * @param[in] constant value to be set + * @param[in] sizeBytes Data size in bytes + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized + */ +hipError_t hipMemsetD16(hipDeviceptr_t dest, unsigned short value, size_t sizeBytes); + +/** + * @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant + * short value value. + * + * hipMemsetD16Async() is asynchronous with respect to the host, so the call may return before the + * memset is complete. The operation can optionally be associated to a stream by passing a non-zero + * stream argument. If stream is non-zero, the operation may overlap with operations in other + * streams. + * + * @param[out] dst Data ptr to be filled + * @param[in] constant value to be set + * @param[in] sizeBytes Data size in bytes + * @param[in] stream - Stream identifier + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized + */ +hipError_t hipMemsetD16Async(hipDeviceptr_t dest, unsigned short value, size_t sizeBytes, hipStream_t stream __dparm(0)); + /** * @brief Fills the memory area pointed to by dest with the constant integer * value for specified number of times. diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 20d2fa8842..b788aa5cdc 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -459,6 +459,10 @@ inline static hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, return hipCUDAErrorTohipError(cudaMallocPitch(ptr, pitch, width, height)); } +inline static hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr,size_t* pitch,size_t widthInBytes,size_t height,unsigned int elementSizeBytes){ + return hipCUResultTohipError(cuMemAllocPitch(dptr,pitch,widthInBytes,height,elementSizeBytes)); +} + inline static hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { return hipCUDAErrorTohipError(cudaMalloc3D(pitchedDevPtr, extent)); } @@ -471,6 +475,12 @@ inline static hipError_t hipMallocHost(void** ptr, size_t size) { return hipCUDAErrorTohipError(cudaMallocHost(ptr, size)); } +inline static hipError_t hipMemAllocHost(void** ptr, size_t size) + __attribute__((deprecated("use hipHostMalloc instead"))); +inline static hipError_t hipMemAllocHost(void** ptr, size_t size) { + return hipCUResultTohipError(cuMemAllocHost(ptr, size)); +} + inline static hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute__((deprecated("use hipHostMalloc instead"))); inline static hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) { @@ -761,6 +771,20 @@ inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, s return hipCUResultTohipError(cuMemsetD8(dest, value, sizeBytes)); } +inline static hipError_t hipMemsetD8Async(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes, + hipStream_t stream __dparm(0)) { + return hipCUResultTohipError(cuMemsetD8Async(dest, value, sizeBytes, stream)); +} + +inline static hipError_t hipMemsetD16(hipDeviceptr_t dest, unsigned short value, size_t sizeBytes) { + return hipCUResultTohipError(cuMemsetD16(dest, value, sizeBytes)); +} + +inline static hipError_t hipMemsetD16Async(hipDeviceptr_t dest, unsigned short value, size_t sizeBytes, + hipStream_t stream __dparm(0)) { + return hipCUResultTohipError(cuMemsetD16Async(dest, value, sizeBytes, stream)); +} + inline static hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) { return hipCUDAErrorTohipError(cudaMemset2D(dst, pitch, value, width, height)); } diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 43eeac739d..ed1422fcda 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -391,6 +391,8 @@ hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) { // Deprecated function: hipError_t hipMallocHost(void** ptr, size_t sizeBytes) { return hipHostMalloc(ptr, sizeBytes, 0); } +// Deprecated function: +hipError_t hipMemAllocHost(void** ptr, size_t sizeBytes) { return hipHostMalloc(ptr, sizeBytes, 0); } // Deprecated function: hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { @@ -465,6 +467,15 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height return ihipLogStatus(hip_status); } +hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, size_t height, unsigned int elementSizeBytes){ + HIP_INIT_SPECIAL_API(hipMemAllocPitch, (TRACE_MEM), dptr, pitch, widthInBytes, height,elementSizeBytes); + HIP_SET_DEVICE(); + + if (widthInBytes == 0 || height == 0) return ihipLogStatus(hipErrorInvalidValue); + + return ihipLogStatus(ihipMallocPitch(tls, dptr, pitch, widthInBytes, height, 0)); +} + hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { HIP_INIT_API(hipMalloc3D, pitchedDevPtr, &extent); HIP_SET_DEVICE(); @@ -1786,28 +1797,65 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp return ihipLogStatus(e); } +hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, hipStream_t stream, bool isAsync) { + if (pCopy == nullptr) { + return hipErrorInvalidValue; + } + void* dst; const void* src; + size_t spitch = pCopy->srcPitch; + size_t dpitch = pCopy->dstPitch; + switch(pCopy->srcMemoryType){ + case hipMemoryTypeHost: + src = pCopy->srcHost; + break; + case hipMemoryTypeArray: + src = pCopy->srcArray->data; + spitch = pCopy->WidthInBytes; + break; + case hipMemoryTypeUnified: + case hipMemoryTypeDevice: + src = pCopy->srcDevice; + break; + default: + return hipErrorInvalidValue; + } + switch(pCopy->dstMemoryType){ + case hipMemoryTypeHost: + dst = pCopy->dstHost; + break; + case hipMemoryTypeArray: + dst = pCopy->dstArray->data; + dpitch = pCopy->WidthInBytes; + break; + case hipMemoryTypeUnified: + case hipMemoryTypeDevice: + dst = pCopy->dstDevice; + break; + default: + return hipErrorInvalidValue; + } + if(pCopy->srcPitch < pCopy->WidthInBytes + pCopy->srcXInBytes || pCopy->srcY >= pCopy->Height){ + return hipErrorInvalidValue; + } else if(pCopy->dstPitch < pCopy->WidthInBytes + pCopy->dstXInBytes || pCopy->dstY >= pCopy->Height){ + return hipErrorInvalidValue; + } + src = (void*)((char*)src+pCopy->srcY*pCopy->srcPitch + pCopy->srcXInBytes); + dst = (void*)((char*)dst+pCopy->dstY*pCopy->dstPitch + pCopy->dstXInBytes); + if(isAsync){ + return ihipMemcpy2DAsync(dst, dpitch, src, spitch, pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault, stream); + } else{ + return ihipMemcpy2D(dst, dpitch, src, spitch, pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault); + } +} + hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { HIP_INIT_SPECIAL_API(hipMemcpyParam2D, (TRACE_MCMD), pCopy); - hipError_t e = hipSuccess; - if (pCopy == nullptr) { - e = hipErrorInvalidValue; - } else { - e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch, - pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault); - } - return ihipLogStatus(e); + return ihipLogStatus(ihipMemcpyParam2D(pCopy, hipStreamNull, false)); } hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, hipStream_t stream) { HIP_INIT_SPECIAL_API(hipMemcpyParam2DAsync, (TRACE_MCMD), pCopy, stream); - hipError_t e = hipSuccess; - if (pCopy == nullptr) { - e = hipErrorInvalidValue; - } else { - e = ihipMemcpy2DAsync(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch, - pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault, stream); - } - return ihipLogStatus(e); + return ihipLogStatus(ihipMemcpyParam2D(pCopy, stream, true)); } // TODO-sync: function is async unless target is pinned host memory - then these are fully sync. @@ -1903,6 +1951,42 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes return ihipLogStatus(e); } +hipError_t hipMemsetD8Async(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes , hipStream_t stream ) { + HIP_INIT_SPECIAL_API(hipMemsetD8Async, (TRACE_MCMD), dst, value, sizeBytes, stream); + + stream = ihipSyncAndResolveStream(stream); + if (stream) { + return ihipLogStatus(ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeChar)); + } else { + return ihipLogStatus(hipErrorInvalidValue); + } +} + +hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t sizeBytes){ + HIP_INIT_SPECIAL_API(hipMemsetD16, (TRACE_MCMD), dst, value, sizeBytes); + hipError_t e = hipSuccess; + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + if (stream) { + e = ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeShort); + if(hipSuccess == e) + stream->locked_wait(); + } else { + e = hipErrorInvalidValue; + } + return ihipLogStatus(e); +} + +hipError_t hipMemsetD16Async(hipDeviceptr_t dst, unsigned short value, size_t sizeBytes, hipStream_t stream ){ + HIP_INIT_SPECIAL_API(hipMemsetD16Async, (TRACE_MCMD), dst, value, sizeBytes, stream); + + stream = ihipSyncAndResolveStream(stream); + if (stream) { + return ihipLogStatus(ihipMemset(dst, value, sizeBytes, stream, ihipMemsetDataTypeShort)); + } else { + return ihipLogStatus(hipErrorInvalidValue); + } +} + hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) { HIP_INIT_SPECIAL_API(hipMemsetD32, (TRACE_MCMD), dst, value, count); diff --git a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp index 2056cffa91..b78e1d8c63 100644 --- a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp @@ -234,7 +234,7 @@ void checkFunctional() { int main() { bool* result{nullptr}; - hipHostMalloc(&result, sizeof(result)); + hipMemAllocHost((void**)&result, sizeof(result)); result[0] = false; hipLaunchKernelGGL( diff --git a/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp b/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp index c022a40079..9b3b18521e 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ @@ -34,6 +34,21 @@ void printSep() { "======================================================================================\n"); } +inline void initMemCpyParam2D(hip_Memcpy2D &ins, const size_t dpitch, + const size_t spitch, const size_t width, + const size_t height, hipMemoryType dstType, + enum hipMemoryType srcType) { + ins.srcXInBytes=0; + ins.srcY=0; + ins.srcPitch=spitch; + ins.dstXInBytes=0; + ins.dstY=0; + ins.dstPitch=dpitch; + ins.WidthInBytes=width; + ins.Height=height; + ins.dstMemoryType= dstType; + ins.srcMemoryType= srcType; +} //--- // Test copies of a matrix numW by numH @@ -65,7 +80,11 @@ void memcpy2Dtest(size_t numW, size_t numH, bool usePinnedHost) { unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numW * numH); HIPCHECK(hipMemcpy2D(A_d, pitch_A, A_h, width, width, numH, hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpy2D(B_d, pitch_B, B_h, width, width, numH, hipMemcpyHostToDevice)); + hip_Memcpy2D ins; + initMemCpyParam2D(ins,pitch_B,width,width,numH,hipMemoryTypeDevice,hipMemoryTypeHost); + ins.dstDevice = (hipDeviceptr_t)B_d; + ins.srcHost = B_h; + HIPCHECK(hipMemcpyParam2D(&ins)); hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, C_d, (pitch_C / sizeof(T)) * numH); @@ -115,7 +134,11 @@ void memcpyArraytest(size_t numW, size_t numH, bool usePinnedHost, bool usePitch unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numW * numH); HIPCHECK(hipMemcpyToArray(A_d, 0, 0, (void*)A_h, width, hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpyToArray(B_d, 0, 0, (void*)B_h, width, hipMemcpyHostToDevice)); + hip_Memcpy2D ins; + initMemCpyParam2D(ins,width,width,width,numH,hipMemoryTypeArray,hipMemoryTypeHost); + ins.dstArray = B_d; + ins.srcHost = B_h; + HIPCHECK(hipMemcpyParam2D(&ins)); hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, (T*)A_d->data, (T*)B_d->data, (T*)C_d->data, numW); @@ -152,15 +175,21 @@ void memcpyArraytest(size_t numW, size_t numH, bool usePinnedHost, bool usePitch } else { HIPCHECK(hipMemcpy2DToArray(A_d, 0, 0, (void*)A_h, width, width, numH, hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpy2DToArray(B_d, 0, 0, (void*)B_h, width, width, numH, - hipMemcpyHostToDevice)); + hip_Memcpy2D ins; + initMemCpyParam2D(ins,width,width,width,numH,hipMemoryTypeArray,hipMemoryTypeHost); + ins.dstArray = B_d; + ins.srcHost = B_h; + HIPCHECK(hipMemcpyParam2D(&ins)); } hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, (T*)A_d->data, (T*)B_d->data, (T*)C_d->data, numW * numH); - - HIPCHECK(hipMemcpy2D((void*)C_h, width, (void*)C_d->data, width, width, numH, - hipMemcpyDeviceToHost)); + printf("memcpy srcArray to dstHost\n"); + hip_Memcpy2D ins; + initMemCpyParam2D(ins,width,width,width,numH,hipMemoryTypeHost,hipMemoryTypeArray); + ins.srcArray = C_d; + ins.dstHost = C_h; + HIPCHECK(hipMemcpyParam2D(&ins)); HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, numW * numH); @@ -227,8 +256,6 @@ int main(int argc, char* argv[]) { printf("\n\n=== tests&1 (types)\n"); printSep(); HIPCHECK(hipDeviceReset()); - size_t width = N / 6; - size_t height = N / 6; memcpy2Dtest(321, 211, 0); memcpy2Dtest(322, 211, 0); memcpy2Dtest(320, 211, 0); diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp index 7edfe16b3c..fac83ec003 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp @@ -26,11 +26,11 @@ THE SOFTWARE. * BUILD: %t %s ../../test_common.cpp * TEST: %t * //Small copy - * TEST: %t -N 10 --memsetval 0x42 --memsetD32val 0x101 + * TEST: %t -N 10 --memsetval 0x42 --memsetD32val 0x101 --memsetD16val 0x10 --memsetD8val 0x1 * // Oddball size - * TEST: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF + * TEST: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF --memsetD16val 0xDEAD --memsetD8val 0xDE * // Big copy - * TEST: %t -N 256M --memsetval 0xa6 --memsetD32val 0xCAFEBABE + * TEST: %t -N 256M --memsetval 0xa6 --memsetD32val 0xCAFEBABE --memsetD16val 0xCAFE --memsetD8val 0xCA * HIT_END */ @@ -86,6 +86,54 @@ bool testhipMemsetD32(int memsetD32val,int p_gpuDevice) return testResult; } +bool testhipMemsetD16(short memsetD16val,int p_gpuDevice) +{ + size_t Nbytes = N*sizeof(int); + printf ("testhipMemsetD16 N=%zu memsetD16val=%4x device=%d\n", N, memsetD16val, p_gpuDevice); + short *A_d; + short *A_h; + bool testResult = true; + + HIPCHECK ( hipMalloc(&A_d, Nbytes) ); + A_h = (short*)malloc(Nbytes); + HIPCHECK ( hipMemsetD16((hipDeviceptr_t)A_d, memsetD16val, N) ); + HIPCHECK ( hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + + for (int i=0; i= argc || !HipTest::parseInt(argv[i], &ex)) { + failed("Bad memsetD16val argument"); + } + memsetD16val = ex; + } else if (!strcmp(arg, "--memsetD8val")) { + int ex; + if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) { + failed("Bad memsetD8val argument"); + } + memsetD8val = ex; } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { if (++i >= argc || !HipTest::parseInt(argv[i], &iterations)) { failed("Bad iterations argument"); diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index 016cc34d52..73a952b0d1 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/tests/src/test_common.h @@ -109,6 +109,8 @@ THE SOFTWARE. extern size_t N; extern char memsetval; extern int memsetD32val; +extern short memsetD16val; +extern char memsetD8val; extern int iterations; extern unsigned blocksPerCU; extern unsigned threadsPerBlock; From d8a79dda8c2f1751220a9f67539910addb8e42c6 Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Fri, 4 Oct 2019 13:38:02 +0530 Subject: [PATCH 06/17] Removed definition of abs(), real() & imag() from hip_complex.h (#1448) Addresses SWDEV-201461. [ROCm/hip commit: 58a476abc255f1ead80dff8cd862989087489f0e] --- .../hip/include/hip/hcc_detail/hip_complex.h | 18 +----------------- 1 file changed, 1 insertion(+), 17 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 75930c469e..4acee9e235 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -301,20 +301,4 @@ __device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, h return make_hipDoubleComplex(real, imag); } -// Complex functions returning real numbers. -#define __DEFINE_HIP_COMPLEX_REAL_FUN(func, hipFun) \ -__device__ __host__ inline float func(const hipFloatComplex& z) { return hipFun##f(z); } \ -__device__ __host__ inline double func(const hipDoubleComplex& z) { return hipFun(z); } - -__DEFINE_HIP_COMPLEX_REAL_FUN(abs, hipCabs) -__DEFINE_HIP_COMPLEX_REAL_FUN(real, hipCreal) -__DEFINE_HIP_COMPLEX_REAL_FUN(imag, hipCimag) - -// Complex functions returning complex numbers. -#define __DEFINE_HIP_COMPLEX_FUN(func, hipFun) \ -__device__ __host__ inline hipFloatComplex func(const hipFloatComplex& z) { return hipFun##f(z); } \ -__device__ __host__ inline hipDoubleComplex func(const hipDoubleComplex& z) { return hipFun(z); } - -__DEFINE_HIP_COMPLEX_FUN(conj, hipConj) - -#endif +#endif //HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H From 3b9b5fca8de33a8572e3297ea21a57e2b7919d17 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 4 Oct 2019 01:08:45 -0700 Subject: [PATCH 07/17] Add texref get APIs support (#1471) Added support for - hipTexRefGetArray hipTexRefGetAddressMode hipTexRefGetAddress [ROCm/hip commit: bec725dec24b0db3ca2e66dab00715e203c3fece] --- .../include/hip/hcc_detail/hip_runtime_api.h | 6 +++ projects/hip/src/hip_texture.cpp | 48 +++++++++++++++++++ projects/hip/src/hip_texture.h | 1 + 3 files changed, 55 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 1ccc87d216..c07d2ad9f1 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -3375,8 +3375,12 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, hipTextureObject_t textureObject); hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags); +hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex); + hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am); +hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* am, textureReference tex, int dim); + hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm); hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags); @@ -3386,6 +3390,8 @@ hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int Nu hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr, size_t size); +hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex); + hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc, hipDeviceptr_t devPtr, size_t pitch); diff --git a/projects/hip/src/hip_texture.cpp b/projects/hip/src/hip_texture.cpp index fe97895381..6d4ef24673 100644 --- a/projects/hip/src/hip_texture.cpp +++ b/projects/hip/src/hip_texture.cpp @@ -437,6 +437,7 @@ hipError_t ihipBindTextureImpl(TlsData *tls_, int dim, enum hipTextureReadMode r return hipErrorRuntimeOther; } getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler); + pTexture->devPtr = (void*) devPtr; textureHash[textureObject] = pTexture; } @@ -509,6 +510,7 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode return hipErrorRuntimeOther; } getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler); + pTexture->devPtr = (void*) devPtr; textureHash[textureObject] = pTexture; } @@ -607,6 +609,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData *tls_, int dim, enum hipTextureRea return hipErrorRuntimeOther; } getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler); + pTexture->devPtr = (void*) array; textureHash[textureObject] = pTexture; } @@ -724,6 +727,17 @@ hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAdd return ihipLogStatus(hip_status); } +hipError_t hipTexRefGetAddressMode(hipTextureAddressMode* am, textureReference tex, int dim) { + HIP_INIT_API(hipTexRefGetAddressMode,am, &tex, dim); + + if ((am == nullptr) || (dim >= 3)) + return ihipLogStatus(hipErrorInvalidValue); + + *am = tex.addressMode[dim]; + + return ihipLogStatus(hipSuccess); +} + hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags) { HIP_INIT_API(hipTexRefSetArray, tex, array, flags); hipError_t hip_status = hipSuccess; @@ -733,6 +747,23 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi return ihipLogStatus(hip_status); } +hipError_t hipTexRefGetArray(hipArray_t* array, textureReference tex) { + HIP_INIT_API(hipTexRefGetArray, array, &tex); + + if (array == nullptr) + return ihipLogStatus(hipErrorInvalidValue); + + hipTexture* pTexture = textureHash[tex.textureObject]; + if((pTexture == nullptr) || (hipResourceTypeArray != pTexture->resDesc.resType)) + return ihipLogStatus(hipErrorInvalidImage); + + if (pTexture->devPtr == nullptr) + return ihipLogStatus(hipErrorUnknown); + + *array = reinterpret_cast(pTexture->devPtr); + + return ihipLogStatus(hipSuccess); +} hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr, size_t size) { @@ -744,6 +775,23 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep return ihipLogStatus(hip_status); } +hipError_t hipTexRefGetAddress(hipDeviceptr_t* dev_ptr, textureReference tex) { + HIP_INIT_API(hipTexRefGetAddress,dev_ptr, &tex); + + if (dev_ptr == nullptr) + return ihipLogStatus(hipErrorInvalidValue); + + hipTexture* pTexture = textureHash[tex.textureObject]; + if (pTexture == nullptr) + return ihipLogStatus(hipErrorInvalidImage); + + if (pTexture->devPtr == nullptr) + return ihipLogStatus(hipErrorUnknown); + + *dev_ptr = reinterpret_cast(pTexture->devPtr); + return ihipLogStatus(hipSuccess); +} + hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc, hipDeviceptr_t devPtr, size_t pitch) { HIP_INIT_API(hipTexRefSetAddress2D, tex, desc, devPtr, pitch); diff --git a/projects/hip/src/hip_texture.h b/projects/hip/src/hip_texture.h index 365a02e103..1affc48b85 100644 --- a/projects/hip/src/hip_texture.h +++ b/projects/hip/src/hip_texture.h @@ -31,6 +31,7 @@ struct hipTexture { hipResourceViewDesc resViewDesc; hsa_ext_image_t image; hsa_ext_sampler_t sampler; + void* devPtr; }; #endif From aeba964a7e4c0815bf5a2713330cc2610f52f8f3 Mon Sep 17 00:00:00 2001 From: Philip Salzmann Date: Fri, 4 Oct 2019 10:09:19 +0200 Subject: [PATCH 08/17] Fix uninitialized var in hipDeviceGetAttribute (#1497) This fixes the usage of an uninitialized cdattr variable in hipDeviceGetAttribute for the CUDA backend when taking the switch default, as detailed in #1317. Note that the directed_tests/runtimeApi/device/hipGetDeviceAttribute.tst test fails for me, but it already did before applying this patch. Let's see what CI says! [ROCm/hip commit: 9ababa4276e1950b6966cf2752ce825ad6f2b299] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index b788aa5cdc..97551d15d4 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -992,8 +992,7 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att cdattr = cudaDevAttrEccEnabled; break; default: - cerror = cudaErrorInvalidValue; - break; + return hipCUDAErrorTohipError(cudaErrorInvalidValue); } cerror = cudaDeviceGetAttribute(pi, cdattr, device); From 4b05ebc19479d95b93ec672ae7d6360e3afd62b2 Mon Sep 17 00:00:00 2001 From: mhbliao <47895780+mhbliao@users.noreply.github.com> Date: Fri, 4 Oct 2019 04:09:29 -0400 Subject: [PATCH 09/17] [hip] Add the missing header . (#1498) [ROCm/hip commit: 4d632166bb4b552bc58ef0efeb52cb1d3f47a9e7] --- projects/hip/src/trace_helper.h | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip/src/trace_helper.h b/projects/hip/src/trace_helper.h index 0f09b98e08..84f218a438 100644 --- a/projects/hip/src/trace_helper.h +++ b/projects/hip/src/trace_helper.h @@ -27,6 +27,7 @@ THE SOFTWARE. #include #include +#include #include //--- // Helper functions to convert HIP function arguments into strings. From f05b0fcb6046c36380937ee0c1ceab29d140b4ef Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 4 Oct 2019 14:31:15 +0300 Subject: [PATCH 10/17] [HIPIFY] CUDA2HIP_Perl code clean-up and formatting + Generated hipify-perl is unchanged [ROCm/hip commit: 6d6258d15e9c6fe836c0968dc238084252675659] --- .../hip/hipify-clang/src/CUDA2HIP_Perl.cpp | 633 +++++++++--------- 1 file changed, 305 insertions(+), 328 deletions(-) diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index f74d196cac..f0a7c093f8 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/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); From ffee68f4f9ec085f1a944329e8fc9c7d0de43909 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Sat, 5 Oct 2019 15:53:04 +0300 Subject: [PATCH 11/17] [HIPIFY][SPARSE] sync with hipSPARSE + Update doc and hipify-perl accordingly [ROCm/hip commit: 410f6c4a953b57ffc8ed0d127e865d88262caa68] --- projects/hip/bin/hipify-perl | 26 ++++++++++++++----- .../markdown/CUSPARSE_API_supported_by_HIP.md | 26 ++++++++++--------- .../src/CUDA2HIP_SPARSE_API_functions.cpp | 24 ++++++++--------- .../src/CUDA2HIP_SPARSE_API_types.cpp | 3 +++ 4 files changed, 49 insertions(+), 30 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 5bf6a4fe9a..f13b865be4 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -197,6 +197,8 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcuIpcGetMemHandle\b/hipIpcGetMemHandle/g; $ft{'memory'} += s/\bcuIpcOpenMemHandle\b/hipIpcOpenMemHandle/g; $ft{'memory'} += s/\bcuMemAlloc\b/hipMalloc/g; + $ft{'memory'} += s/\bcuMemAllocHost\b/hipMemAllocHost/g; + $ft{'memory'} += s/\bcuMemAllocHost_v2\b/hipMemAllocHost/g; $ft{'memory'} += s/\bcuMemAllocManaged\b/hipMemAllocManaged/g; $ft{'memory'} += s/\bcuMemAllocPitch\b/hipMemAllocPitch/g; $ft{'memory'} += s/\bcuMemAllocPitch_v2\b/hipMemAllocPitch/g; @@ -209,8 +211,6 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcuMemGetInfo\b/hipMemGetInfo/g; $ft{'memory'} += s/\bcuMemGetInfo_v2\b/hipMemGetInfo/g; $ft{'memory'} += s/\bcuMemHostAlloc\b/hipHostMalloc/g; - $ft{'memory'} += s/\bcuMemAllocHost\b/hipMemAllocHost/g; - $ft{'memory'} += s/\bcuMemAllocHost_v2\b/hipMemAllocHost/g; $ft{'memory'} += s/\bcuMemHostGetDevicePointer\b/hipHostGetDevicePointer/g; $ft{'memory'} += s/\bcuMemHostGetDevicePointer_v2\b/hipHostGetDevicePointer/g; $ft{'memory'} += s/\bcuMemHostGetFlags\b/hipMemHostGetFlags/g; @@ -237,15 +237,15 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcuMemcpyHtoDAsync\b/hipMemcpyHtoDAsync/g; $ft{'memory'} += s/\bcuMemcpyHtoDAsync_v2\b/hipMemcpyHtoDAsync/g; $ft{'memory'} += s/\bcuMemcpyHtoD_v2\b/hipMemcpyHtoD/g; + $ft{'memory'} += s/\bcuMemsetD16\b/hipMemsetD16/g; + $ft{'memory'} += s/\bcuMemsetD16Async\b/hipMemsetD16Async/g; + $ft{'memory'} += s/\bcuMemsetD16_v2\b/hipMemsetD16/g; $ft{'memory'} += s/\bcuMemsetD32\b/hipMemsetD32/g; $ft{'memory'} += s/\bcuMemsetD32Async\b/hipMemsetD32Async/g; $ft{'memory'} += s/\bcuMemsetD32_v2\b/hipMemsetD32/g; $ft{'memory'} += s/\bcuMemsetD8\b/hipMemsetD8/g; - $ft{'memory'} += s/\bcuMemsetD8_v2\b/hipMemsetD8/g; $ft{'memory'} += s/\bcuMemsetD8Async\b/hipMemsetD8Async/g; - $ft{'memory'} += s/\bcuMemsetD16\b/hipMemsetD16/g; - $ft{'memory'} += s/\bcuMemsetD16_v2\b/hipMemsetD16/g; - $ft{'memory'} += s/\bcuMemsetD16Async\b/hipMemsetD16Async/g; + $ft{'memory'} += s/\bcuMemsetD8_v2\b/hipMemsetD8/g; $ft{'memory'} += s/\bcudaFree\b/hipFree/g; $ft{'memory'} += s/\bcudaFreeArray\b/hipFreeArray/g; $ft{'memory'} += s/\bcudaFreeHost\b/hipHostFree/g; @@ -625,6 +625,7 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcurandSetQuasiRandomGeneratorDimensions\b/hiprandSetQuasiRandomGeneratorDimensions/g; $ft{'library'} += s/\bcurandSetStream\b/hiprandSetStream/g; $ft{'library'} += s/\bcusparseCreate\b/hipsparseCreate/g; + $ft{'library'} += s/\bcusparseCreateCsrgemm2Info\b/hipsparseCreateCsrgemm2Info/g; $ft{'library'} += s/\bcusparseCreateCsrilu02Info\b/hipsparseCreateCsrilu02Info/g; $ft{'library'} += s/\bcusparseCreateCsrsv2Info\b/hipsparseCreateCsrsv2Info/g; $ft{'library'} += s/\bcusparseCreateHybMat\b/hipsparseCreateHybMat/g; @@ -633,6 +634,9 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcusparseDaxpyi\b/hipsparseDaxpyi/g; $ft{'library'} += s/\bcusparseDcsr2csc\b/hipsparseDcsr2csc/g; $ft{'library'} += s/\bcusparseDcsr2hyb\b/hipsparseDcsr2hyb/g; + $ft{'library'} += s/\bcusparseDcsrgemm\b/hipsparseDcsrgemm/g; + $ft{'library'} += s/\bcusparseDcsrgemm2\b/hipsparseDcsrgemm2/g; + $ft{'library'} += s/\bcusparseDcsrgemm2_bufferSizeExt\b/hipsparseDcsrgemm2_bufferSizeExt/g; $ft{'library'} += s/\bcusparseDcsrilu02\b/hipsparseDcsrilu02/g; $ft{'library'} += s/\bcusparseDcsrilu02_analysis\b/hipsparseDcsrilu02_analysis/g; $ft{'library'} += s/\bcusparseDcsrilu02_bufferSize\b/hipsparseDcsrilu02_bufferSize/g; @@ -646,6 +650,7 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcusparseDcsrsv2_solve\b/hipsparseDcsrsv2_solve/g; $ft{'library'} += s/\bcusparseDdoti\b/hipsparseDdoti/g; $ft{'library'} += s/\bcusparseDestroy\b/hipsparseDestroy/g; + $ft{'library'} += s/\bcusparseDestroyCsrgemm2Info\b/hipsparseDestroyCsrgemm2Info/g; $ft{'library'} += s/\bcusparseDestroyCsrilu02Info\b/hipsparseDestroyCsrilu02Info/g; $ft{'library'} += s/\bcusparseDestroyCsrsv2Info\b/hipsparseDestroyCsrsv2Info/g; $ft{'library'} += s/\bcusparseDestroyHybMat\b/hipsparseDestroyHybMat/g; @@ -665,6 +670,9 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcusparseSaxpyi\b/hipsparseSaxpyi/g; $ft{'library'} += s/\bcusparseScsr2csc\b/hipsparseScsr2csc/g; $ft{'library'} += s/\bcusparseScsr2hyb\b/hipsparseScsr2hyb/g; + $ft{'library'} += s/\bcusparseScsrgemm\b/hipsparseScsrgemm/g; + $ft{'library'} += s/\bcusparseScsrgemm2\b/hipsparseScsrgemm2/g; + $ft{'library'} += s/\bcusparseScsrgemm2_bufferSizeExt\b/hipsparseScsrgemm2_bufferSizeExt/g; $ft{'library'} += s/\bcusparseScsrilu02\b/hipsparseScsrilu02/g; $ft{'library'} += s/\bcusparseScsrilu02_analysis\b/hipsparseScsrilu02_analysis/g; $ft{'library'} += s/\bcusparseScsrilu02_bufferSize\b/hipsparseScsrilu02_bufferSize/g; @@ -693,7 +701,11 @@ sub simpleSubstitutions { $ft{'library'} += s/\bcusparseXcoosortByColumn\b/hipsparseXcoosortByColumn/g; $ft{'library'} += s/\bcusparseXcoosortByRow\b/hipsparseXcoosortByRow/g; $ft{'library'} += s/\bcusparseXcoosort_bufferSizeExt\b/hipsparseXcoosort_bufferSizeExt/g; + $ft{'library'} += s/\bcusparseXcscsort\b/hipsparseXcscsort/g; + $ft{'library'} += s/\bcusparseXcscsort_bufferSizeExt\b/hipsparseXcscsort_bufferSizeExt/g; $ft{'library'} += s/\bcusparseXcsr2coo\b/hipsparseXcsr2coo/g; + $ft{'library'} += s/\bcusparseXcsrgemm2Nnz\b/hipsparseXcsrgemm2Nnz/g; + $ft{'library'} += s/\bcusparseXcsrgemmNnz\b/hipsparseXcsrgemmNnz/g; $ft{'library'} += s/\bcusparseXcsrilu02_zeroPivot\b/hipsparseXcsrilu02_zeroPivot/g; $ft{'library'} += s/\bcusparseXcsrsort\b/hipsparseXcsrsort/g; $ft{'library'} += s/\bcusparseXcsrsort_bufferSizeExt\b/hipsparseXcsrsort_bufferSizeExt/g; @@ -816,6 +828,8 @@ sub simpleSubstitutions { $ft{'type'} += s/\bCUstream_st\b/ihipStream_t/g; $ft{'type'} += s/\bCUtexObject\b/hipTextureObject_t/g; $ft{'type'} += s/\bCUtexref_st\b/textureReference/g; + $ft{'type'} += s/\bcsrgemm2Info\b/csrgemm2Info/g; + $ft{'type'} += s/\bcsrgemm2Info_t\b/csrgemm2Info_t/g; $ft{'type'} += s/\bcsrilu02Info_t\b/csrilu02Info_t/g; $ft{'type'} += s/\bcsrsv2Info_t\b/csrsv2Info_t/g; $ft{'type'} += s/\bcuComplex\b/hipComplex/g; diff --git a/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md b/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md index 3a25acc20c..03f7cc2902 100644 --- a/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUSPARSE_API_supported_by_HIP.md @@ -83,6 +83,8 @@ | typedef |`bsrilu02Info_t` | | | | struct |`csru2csrInfo` | | | | typedef |`csru2csrInfo_t` | | | +| struct |`csrgemm2Info` | |`csrgemm2Info` | +| typedef |`csrgemm2Info_t` | |`csrgemm2Info_t` | | struct |`cusparseColorInfo` | | | | typedef |`cusparseColorInfo_t` | | | | struct |`pruneInfo` | 9.0 | | @@ -163,8 +165,8 @@ |`cusparseDestroyBsric02Info` | | |`cusparseCreateBsrilu02Info` | | |`cusparseDestroyBsrilu02Info` | | -|`cusparseCreateCsrgemm2Info` | | -|`cusparseDestroyCsrgemm2Info` | | +|`cusparseCreateCsrgemm2Info` |`hipsparseCreateCsrgemm2Info` | +|`cusparseDestroyCsrgemm2Info` |`hipsparseDestroyCsrgemm2Info` | |`cusparseCreatePruneInfo` | | 9.0 | |`cusparseDestroyPruneInfo` | | 9.0 | @@ -361,18 +363,18 @@ |`cusparseDcsrgeam2_bufferSizeExt` | | 9.2 | |`cusparseCcsrgeam2_bufferSizeExt` | | 9.2 | |`cusparseZcsrgeam2_bufferSizeExt` | | 9.2 | -|`cusparseXcsrgemmNnz` | | -|`cusparseScsrgemm` | | -|`cusparseDcsrgemm` | | +|`cusparseXcsrgemmNnz` |`hipsparseXcsrgemmNnz` | +|`cusparseScsrgemm` |`hipsparseScsrgemm` | +|`cusparseDcsrgemm` |`hipsparseDcsrgemm` | |`cusparseCcsrgemm` | | |`cusparseZcsrgemm` | | -|`cusparseXcsrgemm2Nnz` | | -|`cusparseScsrgemm2` | | -|`cusparseDcsrgemm2` | | +|`cusparseXcsrgemm2Nnz` |`hipsparseXcsrgemm2Nnz` | +|`cusparseScsrgemm2` |`hipsparseScsrgemm2` | +|`cusparseDcsrgemm2` |`hipsparseDcsrgemm2` | |`cusparseCcsrgemm2` | | |`cusparseZcsrgemm2` | | -|`cusparseScsrgemm2_bufferSizeExt` | | -|`cusparseDcsrgemm2_bufferSizeExt` | | +|`cusparseScsrgemm2_bufferSizeExt` |`hipsparseScsrgemm2_bufferSizeExt` | +|`cusparseDcsrgemm2_bufferSizeExt` |`hipsparseDcsrgemm2_bufferSizeExt` | |`cusparseCcsrgemm2_bufferSizeExt` | | |`cusparseZcsrgemm2_bufferSizeExt` | | @@ -670,8 +672,8 @@ |`cusparseXcoosortByColumn` |`hipsparseXcoosortByColumn` | |`cusparseXcsrsort_bufferSizeExt` |`hipsparseXcsrsort_bufferSizeExt` | |`cusparseXcsrsort` |`hipsparseXcsrsort` | -|`cusparseXcscsort_bufferSizeExt` | | -|`cusparseXcscsort` | | +|`cusparseXcscsort_bufferSizeExt` |`hipsparseXcscsort_bufferSizeExt` | +|`cusparseXcscsort` |`hipsparseXcscsort` | |`cusparseCreateCsru2csrInfo` | | |`cusparseDestroyCsru2csrInfo` | | |`cusparseScsru2csr_bufferSizeExt` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index 1383a6f9b4..0a3e8fbd02 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -63,8 +63,8 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseDestroyBsric02Info", {"hipsparseDestroyBsric02Info", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseCreateBsrilu02Info", {"hipsparseCreateBsrilu02Info", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDestroyBsrilu02Info", {"hipsparseDestroyBsrilu02Info", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseCreateCsrgemm2Info", {"hipsparseCreateCsrgemm2Info", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseDestroyCsrgemm2Info", {"hipsparseDestroyCsrgemm2Info", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCreateCsrgemm2Info", {"hipsparseCreateCsrgemm2Info", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDestroyCsrgemm2Info", {"hipsparseDestroyCsrgemm2Info", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseCreatePruneInfo", {"hipsparseCreatePruneInfo", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDestroyPruneInfo", {"hipsparseDestroyPruneInfo", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -289,20 +289,20 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseCcsrgeam2_bufferSizeExt", {"hipsparseCcsrgeam2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrgeam2_bufferSizeExt", {"hipsparseZcsrgeam2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseXcsrgemmNnz", {"hipsparseXcsrgemmNnz", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseScsrgemm", {"hipsparseScsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseDcsrgemm", {"hipsparseDcsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseXcsrgemmNnz", {"hipsparseXcsrgemmNnz", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseScsrgemm", {"hipsparseScsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDcsrgemm", {"hipsparseDcsrgemm", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseCcsrgemm", {"hipsparseCcsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrgemm", {"hipsparseZcsrgemm", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseXcsrgemm2Nnz", {"hipsparseXcsrgemm2Nnz", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseScsrgemm2", {"hipsparseScsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseDcsrgemm2", {"hipsparseDcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseXcsrgemm2Nnz", {"hipsparseXcsrgemm2Nnz", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseScsrgemm2", {"hipsparseScsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDcsrgemm2", {"hipsparseDcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseCcsrgemm2", {"hipsparseCcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrgemm2", {"hipsparseZcsrgemm2", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseScsrgemm2_bufferSizeExt", {"hipsparseScsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseDcsrgemm2_bufferSizeExt", {"hipsparseDcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseScsrgemm2_bufferSizeExt", {"hipsparseScsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseDcsrgemm2_bufferSizeExt", {"hipsparseDcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseCcsrgemm2_bufferSizeExt", {"hipsparseCcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsrgemm2_bufferSizeExt", {"hipsparseZcsrgemm2_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, @@ -633,8 +633,8 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseXcsrsort_bufferSizeExt", {"hipsparseXcsrsort_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseXcsrsort", {"hipsparseXcsrsort", "", CONV_LIB_FUNC, API_SPARSE}}, - {"cusparseXcscsort_bufferSizeExt", {"hipsparseXcscsort_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseXcscsort", {"hipsparseXcscsort", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseXcscsort_bufferSizeExt", {"hipsparseXcscsort_bufferSizeExt", "", CONV_LIB_FUNC, API_SPARSE}}, + {"cusparseXcscsort", {"hipsparseXcscsort", "", CONV_LIB_FUNC, API_SPARSE}}, {"cusparseCreateCsru2csrInfo", {"hipsparseCreateCsru2csrInfo", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDestroyCsru2csrInfo", {"hipsparseDestroyCsru2csrInfo", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp index f9e3920af4..1d3fe28c62 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_types.cpp @@ -62,6 +62,9 @@ const std::map CUDA_SPARSE_TYPE_NAME_MAP{ {"csru2csrInfo", {"csru2csrInfo", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, {"csru2csrInfo_t", {"csru2csrInfo_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, + {"csrgemm2Info", {"csrgemm2Info", "", CONV_TYPE, API_SPARSE}}, + {"csrgemm2Info_t", {"csrgemm2Info_t", "", CONV_TYPE, API_SPARSE}}, + {"cusparseColorInfo", {"hipsparseColorInfo", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseColorInfo_t", {"hipsparseColorInfo_t", "", CONV_TYPE, API_SPARSE, HIP_UNSUPPORTED}}, From 4f899d487c3042de028785c1ebeb90787eeb9623 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 7 Oct 2019 01:41:52 -0700 Subject: [PATCH 12/17] Fix PCI Domain ID query (#1424) * Fix PCI Domain ID query * Update BDF comment [ROCm/hip commit: 96530cba3b9ad80464c454d5ddf3d742614e816d] --- projects/hip/src/hip_hcc.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index f509409334..7150140e12 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -787,11 +787,13 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) { err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID, &bdf_id); DeviceErrorCheck(err); - // BDFID is 16bit uint: [8bit - BusID | 5bit - Device ID | 3bit - Function/DomainID] - prop->pciDomainID = bdf_id & 0x7; + // BDFID is 16bit uint: [8bit - BusID | 5bit - Device ID | 3bit - FunctionID] prop->pciDeviceID = (bdf_id >> 3) & 0x1F; prop->pciBusID = (bdf_id >> 8) & 0xFF; + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_DOMAIN, &prop->pciDomainID); + DeviceErrorCheck(err); + // Masquerade as a 3.0-level device. This will change as more HW functions are properly // supported. Application code should use the arch.has* to do detailed feature detection. prop->major = 3; From 5aedf2d61bd3e7e10022a969ce103cbb69d963ad Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 7 Oct 2019 14:12:08 +0530 Subject: [PATCH 13/17] [ci] Enable tests on ROCm 2.9 (#1503) [ROCm/hip commit: 87651cffc2d70457aea8ac022c27efd94cfdbe03] --- projects/hip/Jenkinsfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index 1873ff3a2b..41b260c6b1 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -295,13 +295,13 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) -// The following launches 3 builds in parallel: rocm-head, rocm-2.8.x and cuda-10.x -parallel rocm_2_8: +// The following launches 3 builds in parallel: rocm-head, rocm-2.9.x and cuda-10.x +parallel rocm_2_9: { node('hip-rocm') { - String hcc_ver = 'rocm-2.8.x' - String from_image = 'ci_test_nodes/rocm-2.8.x/ubuntu-16.04:latest' + String hcc_ver = 'rocm-2.9.x' + String from_image = 'ci_test_nodes/rocm-2.9.x/ubuntu-16.04:latest' String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' // Checkout source code, dependencies and version files From 512adfd735ef63e6893f4d69e81e674bf4c983d3 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 7 Oct 2019 14:33:04 +0300 Subject: [PATCH 14/17] [HIPIFY] Remove duplicates + C++ maps and docs affected, hipify-perl left unchanged [ROCm/hip commit: b744ec45485f2bdc407b58c0ebcdbfcf383a8ae0] --- .../CUDA_Runtime_API_functions_supported_by_HIP.md | 4 ---- .../docs/markdown/CURAND_API_supported_by_HIP.md | 14 -------------- .../src/CUDA2HIP_BLAS_API_functions.cpp | 6 ------ .../hipify-clang/src/CUDA2HIP_Device_functions.cpp | 5 ----- .../src/CUDA2HIP_SPARSE_API_functions.cpp | 2 +- 5 files changed, 1 insertion(+), 30 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index b4a399e191..2df5829853 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -440,8 +440,6 @@ | `cudaMemcpyFromSymbolAsync` | | | `cudaMemcpyToSymbol` | | | `cudaMemcpyToSymbolAsync` | | -| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` |`hipOccupancyMaxActiveBlocksPerMultiprocessor` | -| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` |`hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags`| | `cudaOccupancyMaxPotentialBlockSize` |`hipOccupancyMaxPotentialBlockSize` | | `cudaOccupancyMaxPotentialBlockSizeWithFlags` | | | `cudaOccupancyMaxPotentialBlockSizeVariableSMem` | | @@ -484,7 +482,6 @@ | struct |`cudaTextureDesc` | |`hipTextureDesc` | | struct |`textureReference` | |`textureReference` | | struct |`surfaceReference` | | | -| struct |`CUuuid_st` | | | | enum |***`cudaCGScope`*** | 9.0 | | | 0 |*`cudaCGScopeInvalid`* | 9.0 | | | 1 |*`cudaCGScopeGrid`* | 9.0 | | @@ -974,7 +971,6 @@ | typedef |`cudaArray_const_t` | |`hipArray_const_t` | | typedef |`cudaEvent_t` | |`hipEvent_t` | | struct |`CUevent_st` | |`ihipEvent_t` | -| typedef |`cudaGraphicsResource_t` | | | | struct |`cudaMipmappedArray` | |`hipMipmappedArray` | | typedef |`cudaMipmappedArray_t` | |`hipMipmappedArray_t` | | typedef |`cudaMipmappedArray_const_t` | |`hipMipmappedArray_const_t` | diff --git a/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md b/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md index a46a2c1afd..d4a3111c9c 100644 --- a/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md @@ -54,20 +54,6 @@ | uint | `curandDirectionVectors64_t` | | | | struct | `curandGenerator_st` | | `hiprandGenerator_st` | | struct* | `curandGenerator_t` | | `hiprandGenerator_t` | -| enum |***`curandOrdering`*** | | | -| enum |***`curandOrdering_t`*** | | | -| 100 |*`CURAND_ORDERING_PSEUDO_BEST`* | | | -| 101 |*`CURAND_ORDERING_PSEUDO_DEFAULT`* | | | -| 102 |*`CURAND_ORDERING_PSEUDO_SEEDED`* | | | -| 201 |*`CURAND_ORDERING_QUASI_DEFAULT`* | | | -| enum |***`curandDirectionVectorSet`*** | | | -| enum |***`curandDirectionVectorSet_t`*** | | | -| 101 |*`CURAND_DIRECTION_VECTORS_32_JOEKUO6`* | | | -| 102 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`* | | | -| 103 |*`CURAND_DIRECTION_VECTORS_64_JOEKUO6`* | | | -| 104 |*`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`* | | | -| uint | `curandDirectionVectors32_t` | | `hiprandDirectionVectors32_t` | -| uint | `curandDirectionVectors64_t` | | | | double | `curandDistribution_st` | | | | double | `curandHistogramM2V_st` | | | | double* | `curandDistribution_t` | | | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp index 9d7f559e1a..4efb0bcdbb 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp @@ -333,12 +333,6 @@ const std::map CUDA_BLAS_FUNCTION_MAP{ {"cublasCtrsm", {"hipblasCtrsm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasZtrsm", {"hipblasZtrsm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, - // TRSM - Batched Triangular Solver - {"cublasStrsmBatched", {"hipblasStrsmBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, - {"cublasDtrsmBatched", {"hipblasDtrsmBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, - {"cublasCtrsmBatched", {"hipblasCtrsmBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, - {"cublasZtrsmBatched", {"hipblasZtrsmBatched", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, - // TRMM {"cublasStrmm", {"hipblasStrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, {"cublasDtrmm", {"hipblasDtrmm", "", CONV_LIB_FUNC, API_BLAS, UNSUPPORTED}}, diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp index bd555156c3..5b4f560e15 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp @@ -354,7 +354,6 @@ const std::map CUDA_DEVICE_FUNC_MAP{ {"__vabsss2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__vadd2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__vaddss2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__vaddss2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__vaddus2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__vavgs2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__vavgu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, @@ -514,7 +513,6 @@ const std::map CUDA_DEVICE_FUNC_MAP{ {"__short_as_half", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__ushort_as_half", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__ldg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, - {"__ldg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__ldcg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__ldca", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__ldcs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, @@ -522,13 +520,11 @@ const std::map CUDA_DEVICE_FUNC_MAP{ {"__hne2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hle2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hge2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, - {"__hge2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hlt2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hgt2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hequ2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hneu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hleu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, - {"__hleu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hgeu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hltu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hgtu2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, @@ -611,7 +607,6 @@ const std::map CUDA_DEVICE_FUNC_MAP{ {"atomicMax", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"atomicInc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"atomicDec", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, - {"atomicDec", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"atomicAnd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"atomicOr", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"atomicXor", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp index 0a3e8fbd02..bc8cfe46a5 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_SPARSE_API_functions.cpp @@ -581,7 +581,7 @@ const std::map CUDA_SPARSE_FUNCTION_MAP{ {"cusparseScsr2csr_compress", {"hipsparseScsr2csr_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseDcsr2csr_compress", {"hipsparseDcsr2csr_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, - {"cusparseDcsr2csr_compress", {"hipsparseDcsr2csr_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, + {"cusparseCcsr2csr_compress", {"hipsparseCcsr2csr_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseZcsr2csr_compress", {"hipsparseZcsr2csr_compress", "", CONV_LIB_FUNC, API_SPARSE, HIP_UNSUPPORTED}}, {"cusparseScsr2hyb", {"hipsparseScsr2hyb", "", CONV_LIB_FUNC, API_SPARSE}}, From a05383fc4c727ca036c0a44520fcfd17b77586e6 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 7 Oct 2019 14:44:05 +0300 Subject: [PATCH 15/17] [HIPIFY] Sync with HIP + Update hipify-perl accordingly [ROCm/hip commit: 6aa7540be863f95f7ecce3654b59cf1952468658] --- projects/hip/bin/hipify-perl | 5 +++++ .../CUDA_Driver_API_functions_supported_by_HIP.md | 6 +++--- .../hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp | 8 ++++---- .../hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp | 2 +- 4 files changed, 13 insertions(+), 8 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index f13b865be4..1cfc82eae4 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -269,6 +269,7 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcudaMemcpy2DAsync\b/hipMemcpy2DAsync/g; $ft{'memory'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g; $ft{'memory'} += s/\bcudaMemcpy3D\b/hipMemcpy3D/g; + $ft{'memory'} += s/\bcudaMemcpy3DAsync\b/hipMemcpy3DAsync/g; $ft{'memory'} += s/\bcudaMemcpyAsync\b/hipMemcpyAsync/g; $ft{'memory'} += s/\bcudaMemcpyFromArray\b/hipMemcpyFromArray/g; $ft{'memory'} += s/\bcudaMemcpyFromSymbol\b/hipMemcpyFromSymbol/g; @@ -336,6 +337,10 @@ sub simpleSubstitutions { $ft{'occupancy'} += s/\bcudaOccupancyMaxActiveBlocksPerMultiprocessor\b/hipOccupancyMaxActiveBlocksPerMultiprocessor/g; $ft{'occupancy'} += s/\bcudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags\b/hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags/g; $ft{'occupancy'} += s/\bcudaOccupancyMaxPotentialBlockSize\b/hipOccupancyMaxPotentialBlockSize/g; + $ft{'texture'} += s/\bcuTexRefGetAddress\b/hipTexRefGetAddress/g; + $ft{'texture'} += s/\bcuTexRefGetAddressMode\b/hipTexRefGetAddressMode/g; + $ft{'texture'} += s/\bcuTexRefGetAddress_v2\b/hipTexRefGetAddress/g; + $ft{'texture'} += s/\bcuTexRefGetArray\b/hipTexRefGetArray/g; $ft{'texture'} += s/\bcuTexRefSetAddress\b/hipTexRefSetAddress/g; $ft{'texture'} += s/\bcuTexRefSetAddress2D\b/hipTexRefSetAddress2D/g; $ft{'texture'} += s/\bcuTexRefSetAddress2D_v2\b/hipTexRefSetAddress2D/g; diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 5a6ca9fc1e..21401b0626 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -1053,9 +1053,9 @@ | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| -| `cuTexRefGetAddress` | | -| `cuTexRefGetAddressMode` | | -| `cuTexRefGetArray` | | +| `cuTexRefGetAddress` |`hipTexRefGetAddress` | +| `cuTexRefGetAddressMode` |`hipTexRefGetAddressMode` | +| `cuTexRefGetArray` |`hipTexRefGetArray` | | `cuTexRefGetBorderColor` | | 8.0 | | `cuTexRefGetFilterMode` | | | `cuTexRefGetFlags` | | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 792d34b3ce..3c00660cae 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -530,10 +530,10 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // 5.21. Texture Reference Management // no analogues - {"cuTexRefGetAddress", {"hipTexRefGetAddress", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuTexRefGetAddress_v2", {"hipTexRefGetAddress", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuTexRefGetAddressMode", {"hipTexRefGetAddressMode", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, - {"cuTexRefGetArray", {"hipTexRefGetArray", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuTexRefGetAddress", {"hipTexRefGetAddress", "", CONV_TEXTURE, API_DRIVER}}, + {"cuTexRefGetAddress_v2", {"hipTexRefGetAddress", "", CONV_TEXTURE, API_DRIVER}}, + {"cuTexRefGetAddressMode", {"hipTexRefGetAddressMode", "", CONV_TEXTURE, API_DRIVER}}, + {"cuTexRefGetArray", {"hipTexRefGetArray", "", CONV_TEXTURE, API_DRIVER}}, {"cuTexRefGetBorderColor", {"hipTexRefGetBorderColor", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefGetFilterMode", {"hipTexRefGetFilterMode", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefGetFlags", {"hipTexRefGetFlags", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp index e3c8212aab..56f2c5a806 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp @@ -305,7 +305,7 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ {"cudaMemcpy3D", {"hipMemcpy3D", "", CONV_MEMORY, API_RUNTIME}}, // no analogue // NOTE: Not equal to cuMemcpy3DAsync due to different signatures - {"cudaMemcpy3DAsync", {"hipMemcpy3DAsync", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaMemcpy3DAsync", {"hipMemcpy3DAsync", "", CONV_MEMORY, API_RUNTIME}}, // no analogue // NOTE: Not equal to cuMemcpy3DPeer due to different signatures {"cudaMemcpy3DPeer", {"hipMemcpy3DPeer", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, From 69ea727e92375adc0c8ebff9f1857c357028fd30 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 8 Oct 2019 15:58:48 +0300 Subject: [PATCH 16/17] [HIPIFY][#1487][fix] Translate correctly kernel names prefixed with namespace + Modify CUDA2HIP_perl for the fix + Add ns_kernel_launch.cu test + Update hipify-perl by hipify-clang -perl [ROCm/hip commit: 3286ffdfc04dcf0bbbab00ee74a66530a0a74023] --- projects/hip/bin/hipify-perl | 24 ++++++++-------- .../hip/hipify-clang/src/CUDA2HIP_Perl.cpp | 25 +++++++++-------- .../unit_tests/namespace/ns_kernel_launch.cu | 28 +++++++++++++++++++ 3 files changed, 53 insertions(+), 24 deletions(-) create mode 100644 projects/hip/tests/hipify-clang/unit_tests/namespace/ns_kernel_launch.cu diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 1cfc82eae4..8116cfbe7c 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1638,34 +1638,34 @@ sub transformKernelLaunch { my $k = 0; # Handle the kern<...><<>>() syntax with empty args: - $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; + $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; # Handle the kern<<>>() syntax with empty args: - $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5)/g; + $k += s/([:|\w]+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5)/g; # Handle the kern<...><<>>(...) syntax with non-empty args: - $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; + $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; # Handle the kern<<>>(...) syntax with non-empty args: - $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5, /g; + $k += s/([:|\w]+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5, /g; # Handle the kern<...><<>>() syntax with empty args: - $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; + $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; # Handle the kern<<>>() syntax with empty args: - $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0)/g; + $k += s/([:|\w]+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0)/g; # Handle the kern<...><>>(...) syntax with non-empty args: - $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; + $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; # Handle the kern<<>>(...) syntax with non-empty args: - $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g; + $k += s/([:|\w]+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g; # Handle the kern<...><<>>() syntax with empty args: - $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0)/g; + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0)/g; # Handle the kern<<>>() syntax with empty args: - $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g; + $k += s/([:|\w]+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g; # Handle the kern<...><<>>(...) syntax with non-empty args: - $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0, /g; + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0, /g; # Handle the kern<<>>(...) syntax with non-empty args: - $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g; + $k += s/([:|\w]+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g; if ($k) { $ft{'kernel_launch'} += $k; diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index f0a7c093f8..9bbb52a2fa 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -213,35 +213,36 @@ namespace perl { *streamPtr.get() << tab << no_warns << endl; *streamPtr.get() << tab << my_k << endl_2; + string s_k = "$k += s/([:|\\w]+)\\s*"; *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 << s_k << "<(.+)>\\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 << s_k << "<<<\\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:" << 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 << s_k << "<(.+)>\\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 << s_k << "<<<\\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:" << 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 << s_k << "<(.+)>\\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 << s_k << "<<<\\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:" << 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 << s_k << "<(.+)>\\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 << s_k << "<<<\\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:" << 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 << s_k << "<(.+)>\\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 << s_k << "<<<\\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:" << 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 << s_k << "<(.+)>\\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 << s_k << "<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g;" << endl_2; *streamPtr.get() << tab << "if ($k) {" << endl; *streamPtr.get() << tab_2 << "$ft{'kernel_launch'} += $k;" << endl; diff --git a/projects/hip/tests/hipify-clang/unit_tests/namespace/ns_kernel_launch.cu b/projects/hip/tests/hipify-clang/unit_tests/namespace/ns_kernel_launch.cu new file mode 100644 index 0000000000..0d9fa1cf3f --- /dev/null +++ b/projects/hip/tests/hipify-clang/unit_tests/namespace/ns_kernel_launch.cu @@ -0,0 +1,28 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// CHECK: #include +#include + +__global__ void test_0() { + int a = 10; +} + +namespace first { + __global__ void test_1() { + int b = 20; + } + namespace second { + __global__ void test_2() { + int c = 30; + } + } +} + +int main() { + // CHECK: hipLaunchKernelGGL(::test_0, dim3(1), dim3(1), 0, 0); + ::test_0<<<1, 1>>>(); + // CHECK: hipLaunchKernelGGL(first::test_1, dim3(1), dim3(1), 0, 0); + first::test_1<<<1, 1>>>(); + // CHECK: hipLaunchKernelGGL(first::second::test_2, dim3(1), dim3(1), 0, 0); + first::second::test_2<<<1, 1>>>(); + return 0; +} From 7aa4b752591b20938b199f6e19857489ea0b0043 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 8 Oct 2019 22:27:48 +0300 Subject: [PATCH 17/17] [HIPIFY][doc] Add a few words about hipify-perl + Update testing and versions [ROCm/hip commit: de1b81b78a8e12a145f95fb079ef30ce909dbecd] --- projects/hip/hipify-clang/README.md | 159 ++++++++++++++++------------ 1 file changed, 93 insertions(+), 66 deletions(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index c7524ad0f1..ab52c288c8 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -14,6 +14,7 @@ * [Linux](#linux) * [Windows](#windows) - [Running and using hipify-clang](#running-and-using-hipify-clang) + * [hipify-perl](#perl) - [Disclaimer](#disclaimer) @@ -35,7 +36,7 @@ 1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**9.0.0**](http://releases.llvm.org/download.html#9.0.0). -2. [**CUDA**](https://developer.nvidia.com/cuda-downloads) at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**10.1 Update 2**](https://developer.nvidia.com/cuda-downloads). +2. [**CUDA**](https://developer.nvidia.com/cuda-downloads) of at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**10.1 Update 2**](https://developer.nvidia.com/cuda-downloads). | **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** | |:----------------------------------------------------------:|:-------------------------------------------------------------------:|:------------:|:---------:| @@ -253,70 +254,72 @@ Linux 5.2.0 - Platform OS 64 - hipify-clang binary bitness 64 - python 2.7.12 binary bitness ======================================== --- Testing: 61 tests, 12 threads -- -PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 61) -PASS: hipify :: unit_tests/headers/headers_test_03.cu (2 of 61) -PASS: hipify :: unit_tests/headers/headers_test_05.cu (3 of 61) -PASS: hipify :: unit_tests/headers/headers_test_02.cu (4 of 61) -PASS: hipify :: unit_tests/headers/headers_test_01.cu (5 of 61) -PASS: hipify :: unit_tests/device/math_functions.cu (6 of 61) -PASS: hipify :: unit_tests/headers/headers_test_06.cu (7 of 61) -PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 61) -PASS: hipify :: unit_tests/device/device_symbols.cu (9 of 61) -PASS: hipify :: unit_tests/headers/headers_test_04.cu (10 of 61) -PASS: hipify :: unit_tests/headers/headers_test_10.cu (11 of 61) -PASS: hipify :: unit_tests/headers/headers_test_11.cu (12 of 61) -PASS: hipify :: unit_tests/headers/headers_test_08.cu (13 of 61) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (14 of 61) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 61) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (16 of 61) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (17 of 61) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (18 of 61) -PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (19 of 61) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (20 of 61) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (21 of 61) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (22 of 61) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (23 of 61) -PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (24 of 61) -PASS: hipify :: unit_tests/headers/headers_test_09.cu (25 of 61) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (27 of 61) -PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (28 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (29 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (30 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (31 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (32 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (33 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (34 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (35 of 61) -PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (36 of 61) -PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (37 of 61) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (38 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (39 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (40 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (41 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (42 of 61) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (43 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (44 of 61) -PASS: hipify :: unit_tests/samples/MallocManaged.cpp (45 of 61) -PASS: hipify :: unit_tests/samples/allocators.cu (46 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (47 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (48 of 61) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (49 of 61) -PASS: hipify :: unit_tests/samples/coalescing.cu (50 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (51 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (52 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (53 of 61) -PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (54 of 61) -PASS: hipify :: unit_tests/samples/intro.cu (55 of 61) -PASS: hipify :: unit_tests/samples/axpy.cu (56 of 61) -PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (57 of 61) -PASS: hipify :: unit_tests/samples/square.cu (58 of 61) -PASS: hipify :: unit_tests/samples/static_shared_memory.cu (59 of 61) -PASS: hipify :: unit_tests/samples/vec_add.cu (60 of 61) -PASS: hipify :: unit_tests/samples/cudaRegister.cu (61 of 61) -Testing Time: 2.81s - Expected Passes : 61 +-- Testing: 63 tests, 12 threads -- +PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 63) +PASS: hipify :: unit_tests/headers/headers_test_01.cu (2 of 63) +PASS: hipify :: unit_tests/headers/headers_test_03.cu (3 of 63) +PASS: hipify :: unit_tests/headers/headers_test_02.cu (4 of 63) +PASS: hipify :: unit_tests/headers/headers_test_05.cu (5 of 63) +PASS: hipify :: unit_tests/device/math_functions.cu (6 of 63) +PASS: hipify :: unit_tests/device/atomics.cu (7 of 63) +PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 63) +PASS: hipify :: unit_tests/headers/headers_test_06.cu (9 of 63) +PASS: hipify :: unit_tests/headers/headers_test_04.cu (10 of 63) +PASS: hipify :: unit_tests/device/device_symbols.cu (11 of 63) +PASS: hipify :: unit_tests/headers/headers_test_10.cu (12 of 63) +PASS: hipify :: unit_tests/headers/headers_test_11.cu (13 of 63) +PASS: hipify :: unit_tests/headers/headers_test_08.cu (14 of 63) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 63) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (16 of 63) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (17 of 63) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (18 of 63) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (19 of 63) +PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (20 of 63) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (21 of 63) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (22 of 63) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (23 of 63) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (24 of 63) +PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (25 of 63) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (27 of 63) +PASS: hipify :: unit_tests/headers/headers_test_09.cu (28 of 63) +PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (29 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (30 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (31 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (32 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (33 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (34 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (35 of 63) +PASS: hipify :: unit_tests/namespace/ns_kernel_launch.cu (36 of 63) +PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (37 of 63) +PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (38 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (39 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (40 of 63) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (41 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (42 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (43 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (44 of 63) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (45 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (46 of 63) +PASS: hipify :: unit_tests/samples/allocators.cu (47 of 63) +PASS: hipify :: unit_tests/samples/MallocManaged.cpp (48 of 63) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (49 of 63) +PASS: hipify :: unit_tests/samples/coalescing.cu (50 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (51 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (52 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (53 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (54 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (55 of 63) +PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (56 of 63) +PASS: hipify :: unit_tests/samples/intro.cu (57 of 63) +PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (58 of 63) +PASS: hipify :: unit_tests/samples/axpy.cu (59 of 63) +PASS: hipify :: unit_tests/samples/square.cu (60 of 63) +PASS: hipify :: unit_tests/samples/vec_add.cu (61 of 63) +PASS: hipify :: unit_tests/samples/static_shared_memory.cu (62 of 63) +PASS: hipify :: unit_tests/samples/cudaRegister.cu (63 of 63) +Testing Time: 2.91s + Expected Passes : 63 [100%] Built target test-hipify ``` ### Windows @@ -331,7 +334,7 @@ LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn-7.0.5.15 - cudnn-7.6.4.38 Build system for the above configurations: -Python 3.6 - 3.7.4, cmake 3.12.3 - 3.15.3, Visual Studio 2017 (15.5.2) - 2019 (16.3.1). +Python 3.6 - 3.7.4, cmake 3.12.3 - 3.15.4, Visual Studio 2017 (15.5.2) - 2019 (16.3.2). Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`: @@ -380,6 +383,30 @@ may be useful. For a list of `hipify-clang` options, run `hipify-clang --help`. +### hipify-perl + +To produce a Perl-based script `hipify-perl`, run `hipify-clang --perl`. + +The `hipify-perl` script, unlike the `hipify-clang`, being based on regular expressions, and not on an abstract syntax tree, has several gaps: + +1. macros expansion; + +2. namespaces: + + - redefines of CUDA entities in user namespaces; + + - using directive; + +3. templates (some cases); + +4. device/host function calls distinguishing; + +5. header files correct injection; + +6. complicated argument lists parsing. + +Nonetheless, `hipify-perl` is easy in use and doesn't check the input source CUDA code for correctness. + ## Disclaimer The information contained herein is for informational purposes only, and is subject to change without notice. While every precaution has been taken in the preparation of this document, it may contain technical inaccuracies, omissions and typographical errors, and AMD is under no obligation to update or otherwise correct this information. Advanced Micro Devices, Inc. makes no representations or warranties with respect to the accuracy or completeness of the contents of this document, and assumes no liability of any kind, including the implied warranties of noninfringement, merchantability or fitness for particular purposes, with respect to the operation or use of AMD hardware, software or other products described herein. No license, including implied or arising by estoppel, to any intellectual property rights is granted by this document. Terms and limitations applicable to the purchase or use of AMD's products are as set forth in a signed agreement between the parties or in AMD's Standard Terms and Conditions of Sale.