[HIPIFY][perl][fix] Treat ::device_function as a device function

+ Do not treat somenamespace::device_function_name as a device function
+ Fix generation of warnUnsupportedDeviceFunctions function in hipify-clang
+ Update hipify-perl based on hipify-clang -perl generation
+ Update device test math_functions.cu for hipify-perl

[Restrictions]
- hipify-perl is yet unable to handle function declarations in user namespaces
- hipify-perl is yet unable to handle using directive


[ROCm/hip commit: 4f59ec25fe]

projects/hip/bin/hipify-perl

@@ -1956,11 +1956,13 @@ sub warnUnsupportedDeviceFunctions
         "umul24"
     )
     {
-        # match math at the beginning of a word, but not if it already has a namespace qualifier ('::') :
-        my $mt = m/[:]?[:]?\b($func)\b(\w*\()/g;
-        if ($mt) {
+        # match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));
+        # function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);
+        my $mt_namespace = m/(\w+)::($func)\s*\(\s*.*\s*\)/g;
+        my $mt = m/($func)\s*\(\s*.*\s*\)/g;
+        if ($mt && !$mt_namespace) {
             $m += $mt;
-            print STDERR "  warning: $fileName:#$line_num : unsupported device function : $_\n";
+            print STDERR "  warning: $fileName:$line_num: unsupported device function \"$func\": $_\n";
         }
     }
     return $m;

projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp

@@ -53,11 +53,13 @@ namespace perl {
       *perlStreamPtr.get() << "\nsub warnUnsupportedDeviceFunctions\n" << "{\n" << space << "my $line_num = shift;\n" << space << "my $m = 0;\n" << space << "foreach $func (\n";
       *perlStreamPtr.get() << sUnsupported.str() << "\n" << space << ")\n";
       *perlStreamPtr.get() << space << "{\n";
-      *perlStreamPtr.get() << double_space << "# match math at the beginning of a word, but not if it already has a namespace qualifier ('::') :\n";
-      *perlStreamPtr.get() << double_space << "my $mt = m/[:]?[:]?\\b($func)\\b(\\w*\\()/g;\n";
-      *perlStreamPtr.get() << double_space << "if ($mt) {\n";
+      *perlStreamPtr.get() << double_space << "# match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
+      *perlStreamPtr.get() << double_space << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);\n";
+      *perlStreamPtr.get() << double_space << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      *perlStreamPtr.get() << double_space << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      *perlStreamPtr.get() << double_space << "if ($mt && !$mt_namespace) {\n";
       *perlStreamPtr.get() << triple_space << "$m += $mt;\n";
-      *perlStreamPtr.get() << triple_space << "print STDERR \"  warning: $fileName:#$line_num : unsupported device function : $_\\n\";\n";
+      *perlStreamPtr.get() << triple_space << "print STDERR \"  warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
       *perlStreamPtr.get() << double_space << "}\n";
       *perlStreamPtr.get() << space << "}\n";
       *perlStreamPtr.get() << space << "return $m;\n";

projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu

@@ -1,5 +1,5 @@
 // RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
-// Test to warn only on device functions umin and umax as unsupported, but not on user defined ones.
+// Synthetic test to warn only on device functions umin and umax as unsupported, but not on user defined ones.
 // ToDo: change lit testing in order to parse the output.
 
 #define LEN 1024
@@ -8,9 +8,11 @@
 #include <algorithm>
 
 namespace my {
+  // user defined function
   unsigned int umin(unsigned int arg1, unsigned int arg2) {
     return (arg1 < arg2) ? arg1 : arg2;
   }
+  // user defined function
   unsigned int umax(unsigned int arg1, unsigned int arg2) {
     return (arg1 > arg2) ? arg1 : arg2;
   }
@@ -18,8 +20,16 @@ namespace my {
 
 __global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsigned int u2)
 {
-  unsigned int _umin = umin(u1, u2);
-  unsigned int _umax = umax(u1, u2);
+  // device function call (warn if unsupported)
+  unsigned int _umin = umin ( u1, u2 );
+  // device function call (warn if unsupported)
+  unsigned int _umax = umax ( u1, u2 );
+  // device function call (warn if unsupported)
+  unsigned int _umin_global = ::umin ( u1, u2 );
+  // device function call (warn if unsupported)
+  unsigned int _umax_global = ::umax(u1, u2);
+  if (_umin != _umin_global) return;
+  if (_umax != _umax_global) return;
   int i = threadIdx.x;
   A[i] = i + _umin;
   B[i] = i + _umax;
@@ -29,7 +39,9 @@ __global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsig
 int main() {
   unsigned int u1 = 33;
   unsigned int u2 = 34;
+  // user defined function call
   unsigned int _min = my::umin(u1, u2);
+  // user defined function call
   unsigned int _max = my::umax(u1, u2);
   float *A, *B, *C;
   // CHECK: hipMalloc((void**)&A, SIZE);