Merge pull request #1481 from emankov/hipify-clang

[HIPIFY] Reconcile transformation (statistics) types between hipify-clang and hipify-perl

hipamd/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)) {

hipamd/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;
   }

hipamd/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] = {

hipamd/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;