From 042d68e5fbbe0ed9e7100c14a267a33e6f2e4b9d Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Sun, 29 Sep 2019 17:33:37 +0300 Subject: [PATCH] [HIPIFY] Reconcile transformation (statistics) types between hipify-clang and hipify-perl + Update hipify-perl accordingly [ROCm/hip commit: 76a439f4c04f942ff94024dc7afdc2800c942700] --- projects/hip/bin/hipify-perl | 26 +++++++++---------- .../hip/hipify-clang/src/HipifyAction.cpp | 4 +-- projects/hip/hipify-clang/src/Statistics.cpp | 8 ++++-- projects/hip/hipify-clang/src/Statistics.h | 2 ++ 4 files changed, 23 insertions(+), 17 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 2e8d8d836e..ef77fe0ea9 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -68,7 +68,7 @@ $no_output = 1 if $examine; push(@whitelist, split(',', $whitelist)); -@statNames = ("error", "init", "version", "device", "context", "module", "memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "kernel_func", "extern_shared", "kern_launch"); +@statNames = ("error", "init", "version", "device", "context", "module", "memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "device_function", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "extern_shared", "kernel_launch"); sub totalStats { my %count = %{ shift() }; @@ -91,23 +91,23 @@ sub printStats { printf STDERR "%s:%d ", $stat, $counts{$stat}; } printf STDERR ")\n warn:%d LOC:%d", $warnings, $loc; -} +}; sub addStats { - my $dest_ref = shift(); - my %adder = %{ shift() }; + my $dest_ref = shift(); + my %adder = %{ shift() }; foreach $key (keys %adder) { $dest_ref->{$key} += $adder{$key}; } -} +}; sub clearStats { - my $dest_ref = shift() ; + my $dest_ref = shift(); my @statNames = @{ shift() }; - foreach $stat (@statNames) { - $dest_ref->{$stat} = 0; + foreach $stat(@statNames) { + $dest_ref->{$stat} = 0; } -} +}; # Count of transforms in all files my %tt; @@ -1697,7 +1697,7 @@ while (@ARGV) { $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g; if ($k) { - $ft{'kern_launch'} += $k; + $ft{'kernel_launch'} += $k; $Tkernels{$1} ++; } } @@ -1706,7 +1706,7 @@ while (@ARGV) { $convertedTags{$1}++; } } - my $hasDeviceCode = $countKeywords + $ft{'kernel_func'}; + 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 (/<<<.*>>>/)) { @@ -1738,13 +1738,13 @@ while (@ARGV) { } } if ($hasDeviceCode > 0) { - $ft{'kernel_func'} += countSupportedDeviceFunctions(); + $ft{'device_function'} += countSupportedDeviceFunctions(); } transformHostFunctions(); # 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'}; + my $kernStuff = $hasDeviceCode + $ft{'kernel_launch'} + $ft{'device_function'}; my $totalCalls = $apiCalls + $kernStuff; $is_dos = m/\r\n$/; if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) { diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index e331cfbec5..cb0a5eedc9 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -380,7 +380,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc ct::Replacement Rep(*SM, launchStart, length, OS.str()); clang::FullSourceLoc fullSL(launchStart, *SM); insertReplacement(Rep, fullSL); - hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_EXECUTION, ApiTypes::API_RUNTIME}; + hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME}; Statistics::current().incrementCounter(counter, refName.str()); return true; } @@ -423,7 +423,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match ct::Replacement Rep(*SM, slStart, repLength, repName); clang::FullSourceLoc fullSL(slStart, *SM); insertReplacement(Rep, fullSL); - hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME}; + hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_EXTERN_SHARED, ApiTypes::API_RUNTIME}; Statistics::current().incrementCounter(counter, refName.str()); return true; } diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp index d3efd4a5d5..9751763be3 100644 --- a/projects/hip/hipify-clang/src/Statistics.cpp +++ b/projects/hip/hipify-clang/src/Statistics.cpp @@ -63,7 +63,9 @@ const char *counterNames[NUM_CONV_TYPES] = { "type", // CONV_TYPE "literal", // CONV_LITERAL "numeric_literal", // CONV_NUMERIC_LITERAL - "define" // CONV_DEFINE + "define", // CONV_DEFINE + "extern_shared", // CONV_EXTERN_SHARED + "kernel_launch" // CONV_KERNEL_LAUNCH }; const char *counterTypes[NUM_CONV_TYPES] = { @@ -102,7 +104,9 @@ const char *counterTypes[NUM_CONV_TYPES] = { "CONV_TYPE", "CONV_LITERAL", "CONV_NUMERIC_LITERAL", - "CONV_DEFINE" + "CONV_DEFINE", + "CONV_EXTERN_SHARED", + "CONV_KERNEL_LAUNCH" }; const char *apiNames[NUM_API_TYPES] = { diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h index 91f493f4a5..051f680fb1 100644 --- a/projects/hip/hipify-clang/src/Statistics.h +++ b/projects/hip/hipify-clang/src/Statistics.h @@ -119,6 +119,8 @@ enum ConvTypes { CONV_LITERAL, CONV_NUMERIC_LITERAL, CONV_DEFINE, + CONV_EXTERN_SHARED, + CONV_KERNEL_LAUNCH, CONV_LAST }; constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;