From f876daa083feffe1639d04f37cd1825af11acec1 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Sat, 28 Sep 2019 20:44:15 +0300 Subject: [PATCH] [HIPIFY][perl] Code clean-up before continuing generation [ROCm/clr commit: cdd1888293041f10164d11a81163a2c05a778924] --- projects/clr/hipamd/bin/hipify-perl | 59 ++++++++++------------------- 1 file changed, 21 insertions(+), 38 deletions(-) diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl index c3dead1506..2e8d8d836e 100755 --- a/projects/clr/hipamd/bin/hipify-perl +++ b/projects/clr/hipamd/bin/hipify-perl @@ -93,13 +93,11 @@ sub printStats { printf STDERR ")\n warn:%d LOC:%d", $warnings, $loc; } -# Add adder stats to dest. Used to add stats for current file to a running total for all files: sub addStats { my $dest_ref = shift(); my %adder = %{ shift() }; foreach $key (keys %adder) { $dest_ref->{$key} += $adder{$key}; - #printf ("D{$key} += %d => %d\n", $adder{$key}, $dest{$key}); } } @@ -111,7 +109,7 @@ sub clearStats { } } -# count of transforms in all files: +# Count of transforms in all files my %tt; clearStats(\%tt, \@statNames); $Twarnings = 0; @@ -142,18 +140,17 @@ 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 - typically that is what we want in this case: - # count of transforms in this file, init to 0 here: + # 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; # keywords like __global__, __shared__ - not converted by hipify-perl, but counted here. - my $warnings = 0; - my $warningsCublas = 0; - my $warningsCurand = 0; - my %warningTags; # hash with counts of particular unknown keywords. + my $countKeywords = 0; + my $warnings = 0; + my %warningTags; my $lineCount = 0; - undef $/; # Read whole file at once, so we can match newlines. + undef $/; + # Read whole file at once, so we can match newlines while () { $ft{'error'} += s/\bcudaGetErrorName\b/hipGetErrorName/g; @@ -1645,13 +1642,11 @@ while (@ARGV) { # CUDA extern __shared__ syntax # Note these only work if declaration is on a single line. { - # match uses ? for <.*> which will be unitialized if this is not present in launch syntax. + # Match uses ? for <.*> which will be unitialized if this is not present in launch syntax no warnings qw/uninitialized/; my $k = 0; - # Match extern __shared__ type foo[]; syntax # Replace as HIP_DYNAMIC_SHARED() macro - $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g; - # test patterns for the regular expression above: + # Match patterns for the below regular expression: #'extern __shared__ double foo[];' #'extern __shared__ unsigned int foo[];' #'extern volatile __shared__ double foo[];' @@ -1662,12 +1657,12 @@ while (@ARGV) { #'extern __shared__ blah::type s[];' #'extern __shared__ typename mapper::type s_data[];' #'extern __attribute__((used)) __shared__ typename mapper::type s_data[];' + $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; } - # CUDA Launch Syntax. Note these only work if launch is on a single line. { - # match uses ? for <.*> which will be unitialized if this is not present in launch syntax. + # Match uses ? for <.*> which will be unitialized if this is not present in launch syntax no warnings qw/uninitialized/; my $k = 0; @@ -1709,28 +1704,25 @@ while (@ARGV) { if ($print_stats) { while (/(\bhip[A-Z]\w+\b)/g) { $convertedTags{$1}++; - #print STDERR "HIP: $1 : ", $translateTags{$1}, "\n"; } } - # guess that we are in device code , or at least in a file that calls device code. - # will almost certainly call one of the coordiante functions - could be fooled by clever macros but usually works: my $hasDeviceCode = $countKeywords + $ft{'kernel_func'}; unless ($quiet_warnings) { - #print STDERR "Check WARNINGs\n"; - # copy into array of lines, process line-by-line to show warnings: - if ($hasDeviceCode or (/\bcuda/) or (/<<<.*>>>/)) { + # Copy into array of lines, process line-by-line to show warnings + if ($hasDeviceCode or (/\bcu/) or (/\bCU_/) or (/\bCUDA_/) or (/<<<.*>>>/)) { my @lines = split /\n/, $_; - my $tmp = $_; # copies the whole file, could be a little smarter here... + # Copy the whole file + my $tmp = $_; my $line_num = 0; foreach (@lines) { $line_num ++; - # remove any whitelisted words: + # Remove any whitelisted words foreach $w (@whitelist) { s/\b$w\b/ZAP/ } my $tag; if ((/(\bcuda[A-Z]\w+)/) or (/<<<.*>>>/)) { - # flag any remaining code that look like cuda API calls, may want to add these to hipify + # Flag any remaining code that look like cuda API calls: may want to add these to hipify $tag = (defined $1) ? $1 : "Launch"; } if (defined $tag) { @@ -1745,16 +1737,11 @@ while (@ARGV) { $_ = $tmp; } } - - # To limit bogus translations, try to make sure we are in a kernel: if ($hasDeviceCode > 0) { $ft{'kernel_func'} += countSupportedDeviceFunctions(); } - transformHostFunctions(); - - # Print it! - # TODO - would like to move this code outside loop but it uses $_ which contains the whole file. + # TODO: would like to move this code outside loop but it uses $_ which contains the whole file unless ($no_output) { my $apiCalls = $ft{'error'} + $ft{'init'} + $ft{'version'} + $ft{'device'} + $ft{'context'} + $ft{'module'} + $ft{'memory'} + $ft{'addressing'} + $ft{'stream'} + $ft{'event'} + $ft{'external_resource_interop'} + $ft{'stream_memory'} + $ft{'execution'} + $ft{'graph'} + $ft{'occupancy'} + $ft{'texture'} + $ft{'surface'} + $ft{'peer'} + $ft{'graphics'} + $ft{'profiler'} + $ft{'openGL'} + $ft{'D3D9'} + $ft{'D3D10'} + $ft{'D3D11'} + $ft{'VDPAU'} + $ft{'EGL'} + $ft{'thread'} + $ft{'complex'} + $ft{'library'} + $ft{'device_library'} + $ft{'include'} + $ft{'include_cuda_main_header'} + $ft{'type'} + $ft{'literal'} + $ft{'numeric_literal'} + $ft{'define'}; my $kernStuff = $hasDeviceCode + $ft{'kern_launch'} + $ft{'kernel_func'}; @@ -1762,9 +1749,6 @@ while (@ARGV) { $is_dos = m/\r\n$/; if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) { # TODO: implement hipify-clang's logic with header files AMAP - # If this file makes kernel builtin calls, and does not include the cuda_runtime.h, - # then add an #include to match "magic" includes provided by NVCC. - # This logic can miss cases where cuda_runtime.h is included by another include file. print $OUTFILE '#include "hip/hip_runtime.h"' . ($is_dos ? "\r\n" : "\n"); } print $OUTFILE "$_"; @@ -1772,12 +1756,11 @@ while (@ARGV) { $lineCount = $_ =~ tr/\n//; } my $totalConverted = totalStats(\%ft); - #printf "TOTAL-CONV=%d\n", $totalConverted; if (($totalConverted+$warnings) and $print_stats) { printStats(" info: converted", \@statNames, \%ft, $warnings, $lineCount); print STDERR " in '$fileName'\n"; } - # Update totals for all files: + # Update totals for all files addStats(\%tt, \%ft); $Twarnings += $warnings; $TlineCount += $lineCount; @@ -1785,7 +1768,7 @@ while (@ARGV) { $TwarningTags{$key} += $warningTags{$key}; } } -#-- Print total stats for all files processed: +# Print total stats for all files processed: if ($print_stats and ($fileCount > 1)) { print STDERR "\n"; printStats(" info: TOTAL-converted", \@statNames, \%tt, $Twarnings, $TlineCount);