AI/rocm-systems
2
0
Tükrözés 0
Kód Hibajegyek Egyesítési kérések Actions Packages Projektek Kiadások Wiki Tevékenység

[HIPIFY][perl] Code clean-up before continuing generation

Forráskód böngészése 
[ROCm/clr commit: cdd1888293]
This commit is contained in:
Evgeny Mankov
2019-09-28 20:44:15 +03:00
szülő ad563cef42
commit f876daa083
1 fájl változott, egészen pontosan 21 új sor hozzáadva és 38 régi sor törölve
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/bin/hipify-perl
+21 -38
Fájl megtekintése
@@ -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 (<INFILE>)
{
$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<T>::type s[];'
#'extern __shared__ typename mapper<Float>::type s_data[];'
#'extern __attribute__((used)) __shared__ typename mapper<Float>::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);