From bf5546f35e6a18bc33d9a3d8a5586e2680234051 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Mon, 30 Sep 2019 17:47:46 +0300
Subject: [PATCH] [HIPIFY][perl] Generate transformExternShared function

+ Update hipify-perl accordingly
---
 hipamd/bin/hipify-perl                    | 37 ++++++++---------------
 hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp | 14 +++++++--
 2 files changed, 25 insertions(+), 26 deletions(-)

diff --git a/hipamd/bin/hipify-perl b/hipamd/bin/hipify-perl
index 080a368d07..877349d847 100755
--- a/hipamd/bin/hipify-perl
+++ b/hipamd/bin/hipify-perl
@@ -91,7 +91,7 @@ sub printStats {
         printf STDERR "%s:%d ", $stat, $counts{$stat};
     }
     printf STDERR ")\n  warn:%d LOC:%d", $warnings, $loc;
-};
+}
 
 sub addStats {
     my $dest_ref  = shift();
@@ -99,7 +99,7 @@ sub addStats {
     foreach $key (keys %adder) {
         $dest_ref->{$key} += $adder{$key};
     }
-};
+}
 
 sub clearStats {
     my $dest_ref  = shift();
@@ -107,7 +107,7 @@ sub clearStats {
     foreach $stat(@statNames) {
         $dest_ref->{$stat} = 0;
     }
-};
+}
 
 sub simpleSubstitutions {
     $ft{'error'} += s/\bcudaGetErrorName\b/hipGetErrorName/g;
@@ -1594,6 +1594,14 @@ sub simpleSubstitutions {
     $ft{'define'} += s/\bcudaTextureTypeCubemapLayered\b/hipTextureTypeCubemapLayered/g;
 }
 
+# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro
+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;
+}
+
 # Count of transforms in all files
 my %tt;
 clearStats(\%tt, \@statNames);
@@ -1643,27 +1651,8 @@ while (@ARGV) {
         $countKeywords += m/__global__/;
         $countKeywords += m/__shared__/;
 
-        # 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
-            no warnings qw/uninitialized/;
-            my $k = 0;
-            # Replace as HIP_DYNAMIC_SHARED() macro
-            # Match patterns for the below regular expression:
-            #'extern __shared__ double foo[];'
-            #'extern __shared__ unsigned int foo[];'
-            #'extern volatile __shared__ double foo[];'
-            #'extern volatile __shared__ unsigned int sdata[];'
-            #'extern __shared__ volatile unsigned int sdata[];'
-            #'extern __shared__ T s[];'
-            #'extern __shared__ T::type s[];'
-            #'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;
-        }
+        transformExternShared();
+
         # 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
diff --git a/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp b/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
index e38a702f01..0ccd69da55 100644
--- a/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
+++ b/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
@@ -143,13 +143,13 @@ namespace perl {
     *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;
+    *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 << "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() << double_tab << "$dest_ref->" << (i ? "{$stat} = 0;" : "{$key} += $adder{$key};") << std::endl << tab << "}" << std::endl << "}" << std::endl;
     }
   }
 
@@ -179,6 +179,15 @@ namespace perl {
     *streamPtr.get() << "}" << std::endl;
   }
 
+  void generateExternShared(std::unique_ptr<std::ostream>& 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 << "no warnings qw/uninitialized/;" << std::endl;
+    *streamPtr.get() << tab << "my $k = 0;" << 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 generateHostFunctions(std::unique_ptr<std::ostream>& streamPtr) {
     *streamPtr.get() << std::endl << sSub << " transformHostFunctions" << " {" << std::endl << tab << sMy;
     std::set<std::string> &funcSet = DeviceSymbolFunctions0;
@@ -295,6 +304,7 @@ namespace perl {
     *streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << std::endl;
     generateStatFunctions(streamPtr);
     generateSimpleSubstitutions(streamPtr);
+    generateExternShared(streamPtr);
     generateHostFunctions(streamPtr);
     generateDeviceFunctions(streamPtr);
     streamPtr.get()->flush();