From db3ebbcbaed0d9fff4f11fa485018298fe93cb09 Mon Sep 17 00:00:00 2001
From: wsttiger <wsttiger@gmail.com>
Date: Tue, 15 Aug 2017 13:22:20 -0500
Subject: [PATCH] Add missing CUDA tags

[ROCm/hip commit: 512214e930f59c5907984a9054ce89634959827a]
---
 projects/hip/bin/hipify-perl               | 9 ++++++---
 projects/hip/include/hip/hip_runtime_api.h | 1 +
 2 files changed, 7 insertions(+), 3 deletions(-)

diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl
index 4b6e6e89b9..a37f728bd1 100755
--- a/projects/hip/bin/hipify-perl
+++ b/projects/hip/bin/hipify-perl
@@ -247,6 +247,7 @@ while (@ARGV) {
         $ft{'err'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g;
         $ft{'err'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g;
         $ft{'err'} += s/\bcudaErrorInvalidDevice\b/hipErrorInvalidDevice/g;
+        $ft{'err'} += s/\bcudaErrorInvalidDevicePointer\b/hipErrorInvalidDevicePointer/g;
         $ft{'err'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g;
         $ft{'err'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g;
         $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g;
@@ -280,6 +281,7 @@ while (@ARGV) {
         $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g;
 
         $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g;
+        $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttribute_t/g;
 
         $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g;
         $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g;
@@ -339,6 +341,7 @@ while (@ARGV) {
         $ft{'event'} += s/\bcudaEventElapsedTime\b/hipEventElapsedTime/g;
         $ft{'event'} += s/\bcudaEventSynchronize\b/hipEventSynchronize/g;
         $ft{'event'} += s/\bcudaEventDisableTiming\b/hipEventDisableTiming/g;
+        $ft{'event'} += s/\bcudaEventQuery\b/hipEventQuery/g;
 
         #--------
         # Streams
@@ -489,15 +492,15 @@ while (@ARGV) {
             my $kernelName;
 
             # Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax:
-            $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, $6, /g;
+            $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*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, 0, /g;
+            $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\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), 0, 0, /g;
+            $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;
diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h
index 0cdace0e99..97733964bf 100644
--- a/projects/hip/include/hip/hip_runtime_api.h
+++ b/projects/hip/include/hip/hip_runtime_api.h
@@ -161,6 +161,7 @@ typedef enum hipError_t {
     hipErrorProfilerNotInitialized  = 6,
     hipErrorProfilerAlreadyStarted  = 7,
     hipErrorProfilerAlreadyStopped  = 8,
+    hipErrorInsufficientDriver      = 35,
     hipErrorInvalidImage            = 200,
     hipErrorInvalidContext          = 201,  ///< Produced when input context is invalid.
     hipErrorContextAlreadyCurrent   = 202,