From 790ab87461c449d2c66ffa0703fe5730f7f9b29b Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Tue, 3 Sep 2019 16:44:20 +0300
Subject: [PATCH] [HIPIFY][perl][#259] Fix

empty<<<1, 2>>> ( );     >>  hipLaunchKernelGGL(empty, dim3(1), dim3(2), 0, 0);
empty<<<1, 2, 0>>>();    >>
empty<<<1, 2, 0, 0>>>(); >>

instead of erroneous:    >> hipLaunchKernelGGL((empty), dim3(1), dim3(2), 0, 0, );


[ROCm/clr commit: e26ec02a99635a6e0529ca76b9f490d9e6f9c6f9]
---
 projects/clr/hipamd/bin/hipify-perl           | 27 ++++++++++---------
 .../hipify-clang/unit_tests/samples/axpy.cu   | 11 ++++++++
 2 files changed, 25 insertions(+), 13 deletions(-)

diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl
index 625d2b6fd2..f9186bfda2 100755
--- a/projects/clr/hipamd/bin/hipify-perl
+++ b/projects/clr/hipamd/bin/hipify-perl
@@ -403,24 +403,25 @@ while (@ARGV) {
         #--------
         # CUDA Launch Syntax
         # Note these only work if launch is on a single line.
-        # Handle the <<numBlocks, blockDim>>> syntax:
         {
             # match uses ? for <.*> which will be unitialized if this is not present in launch syntax.
             no warnings qw/uninitialized/;
             my $k = 0;
-            my $kernelName;
-            # Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax:
-            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, $6, /g;
-            $kernelName = $1 if $k;
-            # Handle the <<numBlocks, blockDim, sharedSize>>> syntax:
-            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, 0, /g;
-            $kernelName = $1 if $k;
-            # Handle the <<numBlocks, blockDim>>> syntax:
-            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), 0, 0, /g;
-            $kernelName = $1 if $k;
-            $ft{'kern'} += $k;
+            # Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax with empty args:
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, $6)/g;
+            # Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax with non-empty args:
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, $6, /g;
+            # Handle the <<numBlocks, blockDim, sharedSize>>> syntax with empty args:
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, 0)/g;
+            # Handle the <<numBlocks, blockDim, sharedSize>>> syntax with non-empty args:
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, 0, /g;
+            # Handle the <<numBlocks, blockDim>>> syntax with empty args:
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), 0, 0)/g;
+            # Handle the <<numBlocks, blockDim>>> syntax with non-empty args:
+            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), 0, 0, /g;
             if ($k) {
-                $Tkernels{$kernelName} ++;
+                $ft{'kern'} += $k;
+                $Tkernels{$1} ++;
             }
         }
         unless ($no_translate_textures) {
diff --git a/projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/axpy.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/axpy.cu
index 6bd1065c83..b315715ff9 100644
--- a/projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/axpy.cu
+++ b/projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/axpy.cu
@@ -19,6 +19,8 @@ __global__ void axpy(T a, T *x, T *y) {
   y[threadIdx.x] = a * x[threadIdx.x];
 }
 
+__global__ void empty() {
+}
 
 int main(int argc, char* argv[]) {
   const int kDataLen = 4;
@@ -64,6 +66,15 @@ int main(int argc, char* argv[]) {
   // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
   KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
 
+  // CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
+  empty<<<1, kDataLen>>> ( );
+
+  // CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
+  empty<<<1, kDataLen, 0>>>();
+
+  // CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
+  empty<<<1, kDataLen, 0, 0>>>();
+
   // CHECK: COMPLETE_LAUNCH;
   COMPLETE_LAUNCH;