From 3d744756a1dc3b936656f79d44c3a5c241e3a6bd Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Tue, 17 Sep 2019 16:17:28 +0300
Subject: [PATCH 1/9] [HIPIFY] Add supported device functions (from
 math_functions.h)

+ [perl] Add generation of sub countSupportedDeviceFunctions
+ [FIX] Do not perform any replacement for supported device functions
+ [perl] Sync hipify-perl accordingly


[ROCm/clr commit: 33cd4dffd519ad4491b8f592dea9ebf30aaf3524]
---
 projects/clr/hipamd/bin/hipify-perl           | 190 +++++++++++++++++-
 .../src/CUDA2HIP_Device_functions.cpp         | 175 ++++++++++++++++
 .../hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp |  74 +++++--
 .../hipamd/hipify-clang/src/HipifyAction.cpp  |   7 +-
 .../hipamd/hipify-clang/src/HipifyAction.h    |   2 +-
 5 files changed, 417 insertions(+), 31 deletions(-)

diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl
index b5fc365955..b75b68d3b4 100755
--- a/projects/clr/hipamd/bin/hipify-perl
+++ b/projects/clr/hipamd/bin/hipify-perl
@@ -1812,15 +1812,191 @@ if ($count_conversions) {
 sub countSupportedDeviceFunctions
 {
     my $m = 0;
-    # TODO: list all of the supported functions
-    # TODO: split the list on math, device, and maybe fp16
     foreach $func (
-        # Synchronization:
-        "__syncthreads",
+        "abs",
+        "acos",
+        "acosf",
+        "acosh",
+        "acoshf",
+        "asin",
+        "asinf",
+        "asinh",
+        "asinhf",
+        "atan",
+        "atan2",
+        "atan2f",
+        "atanf",
+        "atanh",
+        "atanhf",
+        "cbrt",
+        "cbrtf",
+        "ceil",
+        "ceilf",
+        "copysign",
+        "copysignf",
+        "cos",
+        "cosf",
+        "cosh",
+        "coshf",
+        "cospi",
+        "cospif",
+        "cyl_bessel_i0",
+        "cyl_bessel_i0f",
+        "cyl_bessel_i1",
+        "cyl_bessel_i1f",
+        "erf",
+        "erfc",
+        "erfcf",
+        "erfcinv",
+        "erfcinvf",
+        "erfcx",
+        "erfcxf",
+        "erff",
+        "erfinv",
+        "erfinvf",
+        "exp",
+        "exp10",
+        "exp10f",
+        "exp2",
+        "exp2f",
+        "expf",
+        "expm1",
+        "expm1f",
+        "fabs",
+        "fabsf",
+        "fdim",
+        "fdimf",
+        "floor",
+        "floorf",
+        "fma",
+        "fmaf",
+        "fmax",
+        "fmaxf",
+        "fmin",
+        "fminf",
+        "fmod",
+        "fmodf",
+        "frexp",
+        "frexpf",
+        "hypot",
+        "hypotf",
+        "ilogb",
+        "ilogbf",
+        "isfinite",
+        "isinf",
+        "isnan",
+        "j0",
+        "j0f",
+        "j1",
+        "j1f",
+        "jn",
+        "jnf",
+        "labs",
+        "ldexp",
+        "ldexpf",
+        "lgamma",
+        "lgammaf",
+        "llabs",
+        "llrint",
+        "llrintf",
+        "llround",
+        "llroundf",
+        "log",
+        "log10",
+        "log10f",
+        "log1p",
+        "log1pf",
+        "log2",
+        "log2f",
+        "logb",
+        "logbf",
+        "logf",
+        "lrint",
+        "lrintf",
+        "lround",
+        "lroundf",
+        "max",
+        "min",
+        "modf",
+        "modff",
+        "nan",
+        "nanf",
+        "nearbyint",
+        "nearbyintf",
+        "nextafter",
+        "nextafterf",
+        "norm",
+        "norm3d",
+        "norm3df",
+        "norm4d",
+        "norm4df",
+        "normcdf",
+        "normcdff",
+        "normcdfinv",
+        "normcdfinvf",
+        "normf",
+        "pow",
+        "powf",
+        "rcbrt",
+        "rcbrtf",
+        "remainder",
+        "remainderf",
+        "remquo",
+        "remquof",
+        "rhypot",
+        "rhypotf",
+        "rint",
+        "rintf",
+        "rnorm",
+        "rnorm3d",
+        "rnorm3df",
+        "rnorm4d",
+        "rnorm4df",
+        "rnormf",
+        "round",
+        "roundf",
+        "rsqrt",
+        "rsqrtf",
+        "scalbln",
+        "scalblnf",
+        "scalbn",
+        "scalbnf",
+        "signbit",
+        "sin",
+        "sincos",
+        "sincosf",
+        "sincospi",
+        "sincospif",
+        "sinf",
+        "sinh",
+        "sinhf",
+        "sinpi",
+        "sinpif",
+        "sqrt",
+        "sqrtf",
+        "tan",
+        "tanf",
+        "tanh",
+        "tanhf",
+        "tgamma",
+        "tgammaf",
+        "trunc",
+        "truncf",
+        "y0",
+        "y0f",
+        "y1",
+        "y1f",
+        "yn",
+        "ynf"
     )
     {
-        # match device func at the beginning of a word, but not if it already has a namespace qualifier ('::') :
-        $m += m/[:]?[:]?\b($func)\b(\w*\()/g;
+        # match device function from the list, 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;
+        }
     }
     return $m;
 }
@@ -1956,7 +2132,7 @@ sub warnUnsupportedDeviceFunctions
         "umul24"
     )
     {
-        # match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));
+        # match device function from the list, 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;
diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp
index af48279a95..b2271366b1 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp
@@ -25,6 +25,181 @@ THE SOFTWARE.
 // Maps CUDA header names to HIP header names
 const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP{
   // math functions
+  {"abs",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"labs",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"llabs",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fabs",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fabsf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"min",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fminf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fmin",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"max",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fmaxf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fmax",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sin",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cos",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sincos",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sincosf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"tan",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sqrt",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rsqrt",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rsqrtf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"log2",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"exp2",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"exp2f",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"exp10",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"exp10f",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"expm1",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"expm1f",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"log2f",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"log10",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"log",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"log1p",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"log1pf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"floor",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"exp",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cosh",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sinh",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"tanh",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"acosh",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"acoshf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"asinh",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"asinhf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"atanh",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"atanhf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"ldexp",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"ldexpf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"logb",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"logbf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"ilogb",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"ilogbf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"scalbn",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"scalbnf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"scalbln",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"scalblnf",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"frexp",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"frexpf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"round",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"roundf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"lround",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"lroundf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"llround",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"llroundf",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rint",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rintf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"lrint",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"lrintf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"llrint",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"llrintf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"nearbyint",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"nearbyintf",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"ceil",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"trunc",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"truncf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fdim",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fdimf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"atan2",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"atan",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"acos",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"asin",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"hypot",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rhypot",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"hypotf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rhypotf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"norm3d",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rnorm3d",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"norm4d",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rnorm4d",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"norm",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rnorm",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rnormf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"normf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"norm3df",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rnorm3df",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"norm4df",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rnorm4df",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cbrt",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cbrtf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rcbrt",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"rcbrtf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sinpi",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sinpif",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cospi",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cospif",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sincospi",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sincospif",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"pow",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"modf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fmod",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"remainder",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"remainderf",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"remquo",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"remquof",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"j0",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"j0f",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"j1",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"j1f",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"jn",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"jnf",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"y0",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"y0f",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"y1",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"y1f",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"yn",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"ynf",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cyl_bessel_i0",   {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cyl_bessel_i0f",  {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cyl_bessel_i1",   {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cyl_bessel_i1f",  {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erf",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erff",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfinv",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfinvf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfc",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfcf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"lgamma",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfcinv",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfcinvf",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"normcdfinv",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"normcdfinvf",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"normcdf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"normcdff",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfcx",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"erfcxf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"lgammaf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"tgamma",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"tgammaf",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"copysign",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"copysignf",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"nextafter",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"nextafterf",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"nan",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"nanf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fma",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fmaf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"acosf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"asinf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"atanf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"atan2f",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"cosf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sinf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"tanf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"coshf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sinhf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"tanhf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"expf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"logf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"log10f",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"modff",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"powf",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"sqrtf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"ceilf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"floorf",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fmodf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"signbit",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"isfinite",        {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"isnan",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"isinf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
   {"umin",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
   {"llmin",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
   {"ullmin",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
index 75afc0a534..25705e5bd6 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
@@ -39,32 +39,64 @@ namespace perl {
   const std::string space = "    ";
   const std::string double_space = space + space;
   const std::string triple_space = double_space + space;
+  const std::string sSub = "sub";
+  const std::string sReturn_0 = "return 0;\n";
 
-  void generateUnsupportedDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
-    unsigned int num = 0;
+  void generateDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
+    unsigned int countUnsupported = 0;
+    unsigned int countSupported = 0;
+    std::stringstream sSupported;
     std::stringstream sUnsupported;
     for (auto& ma : CUDA_DEVICE_FUNC_MAP) {
-      if (Statistics::isUnsupported(ma.second)) {
-        sUnsupported << (num ? ",\n" : "") << double_space << "\"" << ma.first.str() << "\"";
-        num++;
+      bool isUnsupported = Statistics::isUnsupported(ma.second);
+      (isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << double_space << "\"" << ma.first.str() << "\"";
+      if (isUnsupported) {
+        countUnsupported++;
+      } else {
+        countSupported++;
       }
     }
-    if (num) {
-      *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 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 \\\"$func\\\": $_\\n\";\n";
-      *perlStreamPtr.get() << double_space << "}\n";
-      *perlStreamPtr.get() << space << "}\n";
-      *perlStreamPtr.get() << space << "return $m;\n";
-      *perlStreamPtr.get() << "}\n";
+    std::stringstream subCountSupported;
+    std::stringstream subWarnUnsupported;
+    std::stringstream subCommon;
+    std::string sCommon = space + "my $m = 0;\n" + space + "foreach $func (\n";
+    subCountSupported << "\n" << sSub << " countSupportedDeviceFunctions\n" << "{\n" << (countSupported ? sCommon : space + sReturn_0);
+    subWarnUnsupported << "\n" << sSub << " warnUnsupportedDeviceFunctions\n" << "{\n" << (countUnsupported ? space + "my $line_num = shift;\n" + sCommon : space + sReturn_0);
+    if (countSupported) {
+      subCountSupported << sSupported.str() << "\n" << space << ")\n";
     }
+    if (countUnsupported) {
+      subWarnUnsupported << sUnsupported.str() << "\n" << space << ")\n";
+    }
+    if (countSupported || countUnsupported) {
+      subCommon << space << "{\n";
+      subCommon << double_space << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
+      subCommon << 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";
+      subCommon << double_space << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      subCommon << double_space << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      subCommon << double_space << "if ($mt && !$mt_namespace) {\n";
+      subCommon << triple_space << "$m += $mt;\n";
+    }
+    if (countSupported) {
+      subCountSupported << subCommon.str();
+    }
+    if (countUnsupported) {
+      subWarnUnsupported << subCommon.str();
+      subWarnUnsupported << triple_space << "print STDERR \"  warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
+    }
+    if (countSupported || countUnsupported) {
+      sCommon = double_space + "}\n" + space + "}\n" + space + "return $m;\n";
+    }
+    if (countSupported) {
+      subCountSupported << sCommon;
+    }
+    if (countUnsupported) {
+      subWarnUnsupported << sCommon;
+    }
+    subCountSupported << "}\n";
+    subWarnUnsupported << "}\n";
+    *perlStreamPtr.get() << subCountSupported.str();
+    *perlStreamPtr.get() << subWarnUnsupported.str();
   }
 
   bool generate(bool Generate) {
@@ -125,7 +157,7 @@ namespace perl {
         }
       }
     }
-    generateUnsupportedDeviceFunctions(perlStreamPtr);
+    generateDeviceFunctions(perlStreamPtr);
     perlStreamPtr.get()->flush();
     bool ret = true;
     EC = sys::fs::copy_file(tmpFile, dstPerlMap);
diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
index 80d80b5f2c..e52c4cd2e9 100644
--- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
@@ -83,7 +83,7 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
 
 void HipifyAction::FindAndReplace(llvm::StringRef name,
                                   clang::SourceLocation sl,
-                                  const std::map<llvm::StringRef, hipCounter>& repMap) {
+                                  const std::map<llvm::StringRef, hipCounter>& repMap, bool bReplace) {
   const auto found = repMap.find(name);
   if (found == repMap.end()) {
     // So it's an identifier, but not CUDA? Boring.
@@ -100,6 +100,9 @@ void HipifyAction::FindAndReplace(llvm::StringRef name,
     DE.Report(sl, ID) << sWarn;
     return;
   }
+  if (!bReplace) {
+    return;
+  }
   StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
   clang::SourceManager& SM = getCompilerInstance().getSourceManager();
   ct::Replacement Rep(SM, sl, name.size(), repName.str());
@@ -392,7 +395,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
 bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
   if (const clang::CallExpr *call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaDeviceFuncCall")) {
     const clang::FunctionDecl *funcDcl = call->getDirectCallee();
-    FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP);
+    FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false);
   }
   return true;
 }
diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h
index d38eddca0a..208e6fb0b2 100644
--- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h
+++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h
@@ -100,5 +100,5 @@ protected:
   void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override;
   std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
   bool Exclude(const hipCounter & hipToken);
-  void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map<llvm::StringRef, hipCounter>& repMap);
+  void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map<llvm::StringRef, hipCounter>& repMap, bool bReplace = true);
 };

From 0c6cf5338e1826cb998a5a916519775e458555f1 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Wed, 18 Sep 2019 11:53:23 +0300
Subject: [PATCH 2/9] [HIPIFY][#1430] Build fix for LLVM < 8.0.0

[ROCm/clr commit: 774a6c56671c98f53c0c0875d56c6b16e6d7af2e]
---
 projects/clr/hipamd/hipify-clang/src/main.cpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/projects/clr/hipamd/hipify-clang/src/main.cpp b/projects/clr/hipamd/hipify-clang/src/main.cpp
index 4b426b41f7..2214567df3 100644
--- a/projects/clr/hipamd/hipify-clang/src/main.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/main.cpp
@@ -34,6 +34,9 @@ THE SOFTWARE.
 #include "ArgParse.h"
 #include "StringUtils.h"
 #include "llvm/Support/Debug.h"
+#if LLVM_VERSION_MAJOR < 8
+#include "llvm/Support/Path.h"
+#endif
 
 constexpr auto DEBUG_TYPE = "cuda2hip";
 

From c2f3ca9225a75eb109c134706ffb83dab060d930 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Wed, 18 Sep 2019 17:07:50 +0300
Subject: [PATCH 3/9] [HIPIFY] Add supported device functions (from
 device_functions.h)

+ [perl] Sync hipify-perl accordingly


[ROCm/clr commit: 05d71d4b8b8d3257bbef1e28e06ac4d8f54e8fe2]
---
 projects/clr/hipamd/bin/hipify-perl           | 112 ++++++++++++++++
 .../src/CUDA2HIP_Device_functions.cpp         | 121 +++++++++++++++++-
 2 files changed, 226 insertions(+), 7 deletions(-)

diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl
index b75b68d3b4..650b13c429 100755
--- a/projects/clr/hipamd/bin/hipify-perl
+++ b/projects/clr/hipamd/bin/hipify-perl
@@ -1294,6 +1294,7 @@ while (@ARGV) {
         $ft{'numeric_literal'} += s/\bCU_CTX_SCHED_MASK\b/hipDeviceScheduleMask/g;
         $ft{'numeric_literal'} += s/\bCU_CTX_SCHED_SPIN\b/hipDeviceScheduleSpin/g;
         $ft{'numeric_literal'} += s/\bCU_CTX_SCHED_YIELD\b/hipDeviceScheduleYield/g;
+        $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY\b/hipDeviceAttributeCanMapHostMemory/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_CLOCK_RATE\b/hipDeviceAttributeClockRate/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR\b/hipDeviceAttributeComputeCapabilityMajor/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR\b/hipDeviceAttributeComputeCapabilityMinor/g;
@@ -1301,8 +1302,10 @@ while (@ARGV) {
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS\b/hipDeviceAttributeConcurrentKernels/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH\b/hipDeviceAttributeCooperativeLaunch/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH\b/hipDeviceAttributeCooperativeMultiDeviceLaunch/g;
+        $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_ECC_ENABLED\b/hipDeviceAttributeEccEnabled/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH\b/hipDeviceAttributeMemoryBusWidth/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_INTEGRATED\b/hipDeviceAttributeIntegrated/g;
+        $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT\b/hipDeviceAttributeKernelExecTimeout/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE\b/hipDeviceAttributeL2CacheSize/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH\b/hipDeviceAttributeMaxTexture1DWidth/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT\b/hipDeviceAttributeMaxTexture2DHeight/g;
@@ -1316,6 +1319,7 @@ while (@ARGV) {
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X\b/hipDeviceAttributeMaxGridDimX/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y\b/hipDeviceAttributeMaxGridDimY/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z\b/hipDeviceAttributeMaxGridDimZ/g;
+        $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_PITCH\b/hipDeviceAttributeMaxPitch/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK\b/hipDeviceAttributeMaxRegistersPerBlock/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK\b/hipDeviceAttributeMaxSharedMemoryPerBlock/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g;
@@ -1328,6 +1332,7 @@ while (@ARGV) {
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID\b/hipDeviceAttributePciDeviceId/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_REGISTERS_PER_BLOCK\b/hipDeviceAttributeMaxRegistersPerBlock/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_SHARED_MEMORY_PER_BLOCK\b/hipDeviceAttributeMaxSharedMemoryPerBlock/g;
+        $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT\b/hipDeviceAttributeTextureAlignment/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY\b/hipDeviceAttributeTotalConstantMemory/g;
         $ft{'numeric_literal'} += s/\bCU_DEVICE_ATTRIBUTE_WARP_SIZE\b/hipDeviceAttributeWarpSize/g;
         $ft{'numeric_literal'} += s/\bCU_EVENT_BLOCKING_SYNC\b/hipEventBlockingSync/g;
@@ -1813,6 +1818,111 @@ sub countSupportedDeviceFunctions
 {
     my $m = 0;
     foreach $func (
+        "__brev",
+        "__brevll",
+        "__byte_perm",
+        "__clz",
+        "__clzll",
+        "__cosf",
+        "__double2int_rz",
+        "__double2ll_rz",
+        "__double2uint_rz",
+        "__double2ull_rz",
+        "__exp10f",
+        "__expf",
+        "__fadd_rd",
+        "__fadd_rn",
+        "__fadd_ru",
+        "__fadd_rz",
+        "__fdiv_rd",
+        "__fdiv_rn",
+        "__fdiv_ru",
+        "__fdiv_rz",
+        "__fdividef",
+        "__ffs",
+        "__ffsll",
+        "__float2int_rd",
+        "__float2int_rn",
+        "__float2int_ru",
+        "__float2int_rz",
+        "__float2ll_rd",
+        "__float2ll_rn",
+        "__float2ll_ru",
+        "__float2ll_rz",
+        "__float2uint_rd",
+        "__float2uint_rn",
+        "__float2uint_ru",
+        "__float2uint_rz",
+        "__float2ull_rd",
+        "__float2ull_rn",
+        "__float2ull_ru",
+        "__float2ull_rz",
+        "__float_as_int",
+        "__float_as_uint",
+        "__fmaf_rd",
+        "__fmaf_rn",
+        "__fmaf_ru",
+        "__fmaf_rz",
+        "__fmul_rd",
+        "__fmul_rn",
+        "__fmul_ru",
+        "__fmul_rz",
+        "__frcp_rd",
+        "__frcp_rn",
+        "__frcp_ru",
+        "__frcp_rz",
+        "__frsqrt_rn",
+        "__fsqrt_rd",
+        "__fsqrt_rn",
+        "__fsqrt_ru",
+        "__fsqrt_rz",
+        "__fsub_rd",
+        "__fsub_rn",
+        "__fsub_ru",
+        "__fsub_rz",
+        "__hadd",
+        "__int2float_rd",
+        "__int2float_rn",
+        "__int2float_ru",
+        "__int2float_rz",
+        "__int_as_float",
+        "__ll2float_rd",
+        "__ll2float_rn",
+        "__ll2float_ru",
+        "__ll2float_rz",
+        "__log10f",
+        "__log2f",
+        "__logf",
+        "__mul24",
+        "__mul64hi",
+        "__mulhi",
+        "__popc",
+        "__popcll",
+        "__powf",
+        "__rhadd",
+        "__sad",
+        "__saturatef",
+        "__sincosf",
+        "__sinf",
+        "__syncthreads",
+        "__tanf",
+        "__threadfence",
+        "__threadfence_block",
+        "__uhadd",
+        "__uint2float_rd",
+        "__uint2float_rn",
+        "__uint2float_ru",
+        "__uint2float_rz",
+        "__uint_as_float",
+        "__ull2float_rd",
+        "__ull2float_rn",
+        "__ull2float_ru",
+        "__ull2float_rz",
+        "__umul24",
+        "__umul64hi",
+        "__umulhi",
+        "__urhadd",
+        "__usad",
         "abs",
         "acos",
         "acosf",
@@ -1866,6 +1976,8 @@ sub countSupportedDeviceFunctions
         "fabsf",
         "fdim",
         "fdimf",
+        "fdivide",
+        "fdividef",
         "floor",
         "floorf",
         "fma",
diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp
index b2271366b1..791e477ecb 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Device_functions.cpp
@@ -235,13 +235,120 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP{
   {"int2float",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
   {"uint2float",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
   // device functions
-  {"__prof_trigger",  {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
-  {"__trap",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
-  {"__brkpt",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
-  {"__pm0",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
-  {"__pm1",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
-  {"__pm2",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
-  {"__pm3",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
+  {"__mulhi",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__umulhi",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__mul64hi",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__umul64hi",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__int_as_float",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float_as_int",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__uint_as_float",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float_as_uint",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__syncthreads",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__threadfence",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__threadfence_block", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__saturatef",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__sad",               {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__usad",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__mul24",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__umul24",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fdividef",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fdividef",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"fdivide",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__sinf",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__cosf",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__tanf",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__sincosf",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__expf",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__exp10f",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__log2f",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__log10f",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__logf",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__powf",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2int_rn",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2int_rz",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2int_ru",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2int_rd",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2uint_rn",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2uint_rz",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2uint_ru",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2uint_rd",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__int2float_rn",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__int2float_rz",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__int2float_ru",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__int2float_rd",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__uint2float_rn",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__uint2float_rz",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__uint2float_ru",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__uint2float_rd",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ll_rn",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ll_rz",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ll_ru",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ll_rd",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ull_rn",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ull_rz",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ull_ru",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__float2ull_rd",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ll2float_rn",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ll2float_rz",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ll2float_ru",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ll2float_rd",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ull2float_rn",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ull2float_rz",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ull2float_ru",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ull2float_rd",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fadd_rn",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fadd_rz",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fadd_ru",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fadd_rd",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsub_rn",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsub_rz",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsub_ru",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsub_rd",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmul_rn",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmul_rz",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmul_ru",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmul_rd",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmaf_rn",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmaf_rz",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmaf_ru",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fmaf_rd",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__frcp_rn",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__frcp_rz",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__frcp_ru",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__frcp_rd",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsqrt_rn",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsqrt_rz",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsqrt_ru",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fsqrt_rd",          {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__frsqrt_rn",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fdiv_rn",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fdiv_rz",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fdiv_ru",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__fdiv_rd",           {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__clz",               {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ffs",               {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__popc",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__brev",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__clzll",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__ffsll",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__popcll",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__brevll",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__byte_perm",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__hadd",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__rhadd",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__uhadd",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__urhadd",            {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__double2int_rz",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__double2uint_rz",    {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__double2ll_rz",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__double2ull_rz",     {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
+  {"__prof_trigger",      {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
+  {"__trap",              {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
+  {"__brkpt",             {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
+  {"__pm0",               {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
+  {"__pm1",               {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
+  {"__pm2",               {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
+  {"__pm3",               {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
   // SIMD functions
   {"__vabs2",         {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
   {"__vabsss2",       {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},

From b6aa917b43b03ad294686ce22dc51e986adbce6e Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Wed, 18 Sep 2019 17:57:31 +0300
Subject: [PATCH 4/9] [HIPIFY][#1437] Fix: cudaFuncGetAttributes to
 hipFuncGetAttributes is supported

+ Update hipify-perl and CUDA_Runtime_API_functions_supported_by_HIP.md accordingly


[ROCm/clr commit: a49f9924a7a4d7cab24c9d9c59f2f04283aea91c]
---
 projects/clr/hipamd/bin/hipify-perl                             | 1 +
 .../markdown/CUDA_Runtime_API_functions_supported_by_HIP.md     | 2 +-
 .../hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp  | 2 +-
 3 files changed, 3 insertions(+), 2 deletions(-)

diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl
index 650b13c429..8094b96add 100755
--- a/projects/clr/hipamd/bin/hipify-perl
+++ b/projects/clr/hipamd/bin/hipify-perl
@@ -363,6 +363,7 @@ while (@ARGV) {
         $ft{'execution'} += s/\bcuFuncGetAttribute\b/hipFuncGetAttribute/g;
         $ft{'execution'} += s/\bcuLaunchKernel\b/hipModuleLaunchKernel/g;
         $ft{'execution'} += s/\bcudaConfigureCall\b/hipConfigureCall/g;
+        $ft{'execution'} += s/\bcudaFuncGetAttributes\b/hipFuncGetAttributes/g;
         $ft{'execution'} += s/\bcudaLaunch\b/hipLaunchByPtr/g;
         $ft{'execution'} += s/\bcudaLaunchCooperativeKernel\b/hipLaunchCooperativeKernel/g;
         $ft{'execution'} += s/\bcudaLaunchCooperativeKernelMultiDevice\b/hipLaunchCooperativeKernelMultiDevice/g;
diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
index dfb80a3537..b4a399e191 100644
--- a/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+++ b/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
@@ -101,7 +101,7 @@
 
 |   **CUDA**                                                |   **HIP**                             |**CUDA version\***|
 |-----------------------------------------------------------|---------------------------------------|:----------------:|
-| `cudaFuncGetAttributes`                                   |                                       |
+| `cudaFuncGetAttributes`                                   |`hipFuncGetAttributes`                 |
 | `cudaFuncSetAttribute`                                    |                                       | 9.0              |
 | `cudaFuncSetCacheConfig`                                  |`hipFuncSetCacheConfig`                |
 | `cudaFuncSetSharedMemConfig`                              |                                       |
diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
index f15220d0d1..e3c8212aab 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
@@ -179,7 +179,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
 
   // 5.7. Execution Control
   // no analogue
-  {"cudaFuncGetAttributes",                                   {"hipFuncGetAttributes",                                   "", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
+  {"cudaFuncGetAttributes",                                   {"hipFuncGetAttributes",                                   "", CONV_EXECUTION, API_RUNTIME}},
   // no analogue
   {"cudaFuncSetAttribute",                                    {"hipFuncSetAttribute",                                    "", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
   // no analogue

From 2ee59279d6cda1b4baa9a429721332808e46b291 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Thu, 19 Sep 2019 19:33:42 +0300
Subject: [PATCH 5/9] [HIPIFY][#1435] Add HIP_SYMBOL wrapper to the templated
 Device Symbol argument of the following functions:

cudaMemcpyToSymbol, cudaMemcpyToSymbolAsync, cudaGetSymbolSize, cudaGetSymbolAddress, cudaMemcpyFromSymbol, cudaMemcpyFromSymbolAsync

+ Add a corresponding cudaSymbolFuncCall matcher.
+ Add device_symbols.cu test for the above 6 functions, update existed.
+ Fix dim3() type cast issue, update affected tests.

TODO: Do the same in hipify-perl


[ROCm/clr commit: d4f8c6bc4bbb0f1aa49350443583734562a2ebfc]
---
 .../hipamd/hipify-clang/src/HipifyAction.cpp  |  87 +++++++++-
 .../hipamd/hipify-clang/src/HipifyAction.h    |   1 +
 .../unit_tests/device/device_symbols.cu       | 152 ++++++++++++++++++
 .../cuRAND/benchmark_curand_kernel.cpp        |   4 +-
 .../libraries/cuRAND/poisson_api_example.cu   |   2 +-
 5 files changed, 241 insertions(+), 5 deletions(-)
 create mode 100644 projects/clr/hipamd/tests/hipify-clang/unit_tests/device/device_symbols.cu

diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
index e52c4cd2e9..6cb75f8911 100644
--- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
@@ -20,6 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+#include <algorithm>
+#include <set>
 #include "HipifyAction.h"
 #include "clang/Basic/SourceLocation.h"
 #include "clang/Frontend/CompilerInstance.h"
@@ -34,6 +36,25 @@ THE SOFTWARE.
 namespace ct = clang::tooling;
 namespace mat = clang::ast_matchers;
 
+const std::string sCudaMemcpyToSymbol = "cudaMemcpyToSymbol";
+const std::string sCudaMemcpyToSymbolAsync = "cudaMemcpyToSymbolAsync";
+const std::string sCudaGetSymbolSize = "cudaGetSymbolSize";
+const std::string sCudaGetSymbolAddress = "cudaGetSymbolAddress";
+const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol";
+const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync";
+
+const std::set<std::string> DeviceSymbolFunctions0 {
+  {sCudaMemcpyToSymbol},
+  {sCudaMemcpyToSymbolAsync}
+};
+
+const std::set<std::string> DeviceSymbolFunctions1 {
+  {sCudaGetSymbolSize},
+  {sCudaGetSymbolAddress},
+  {sCudaMemcpyFromSymbol},
+  {sCudaMemcpyFromSymbolAsync}
+};
+
 void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
   clang::SourceManager& SM = getCompilerInstance().getSourceManager();
   size_t begin = 0;
@@ -316,8 +337,12 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
 
   // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero.
   // Copy the two dimensional arguments verbatim.
-  OS << "dim3(" << readSourceText(*SM, config->getArg(0)->getSourceRange()) << "), ";
-  OS << "dim3(" << readSourceText(*SM, config->getArg(1)->getSourceRange()) << "), ";
+  std::string sDim3 = "dim3(";
+  for (unsigned int i = 0; i < 2; ++i) {
+    const std::string sArg = readSourceText(*SM, config->getArg(i)->getSourceRange()).str();
+    bool bDim3 = std::equal(sDim3.begin(), sDim3.end(), sArg.c_str());
+    OS << (bDim3 ? "" : sDim3) << sArg << (bDim3 ? "" : ")") << ", ";
+  }
   // The stream/memory arguments default to zero if omitted.
   OS << stringifyZeroDefaultedArg(*SM, config->getArg(2)) << ", ";
   OS << stringifyZeroDefaultedArg(*SM, config->getArg(3));
@@ -395,11 +420,50 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
 bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
   if (const clang::CallExpr *call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaDeviceFuncCall")) {
     const clang::FunctionDecl *funcDcl = call->getDirectCallee();
+    if (!funcDcl) {
+      return true;
+    }
     FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false);
   }
   return true;
 }
 
+bool HipifyAction::cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
+  if (const clang::CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaSymbolFuncCall")) {
+    if (!call->getNumArgs()) {
+      return true;
+    }
+    const clang::FunctionDecl* funcDcl = call->getDirectCallee();
+    if (!funcDcl) {
+      return true;
+    }
+    std::string sName = funcDcl->getDeclName().getAsString();
+    unsigned int argNum = 0;
+    if (DeviceSymbolFunctions0.find(sName) != DeviceSymbolFunctions0.end()) {
+      argNum = 0;
+    } else if (call->getNumArgs() > 1 && DeviceSymbolFunctions1.find(sName) != DeviceSymbolFunctions1.end()) {
+      argNum = 1;
+    } else {
+      return true;
+    }
+    clang::SmallString<40> XStr;
+    llvm::raw_svector_ostream OS(XStr);
+    clang::SourceRange sr = call->getArg(argNum)->getSourceRange();
+    clang::SourceManager* SM = Result.SourceManager;
+    const std::string sSymbol = "HIP_SYMBOL";
+    OS << sSymbol << "(" << readSourceText(*SM, sr) << ")";
+    clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() });
+    clang::SourceLocation s = replacementRange.getBegin();
+    clang::SourceLocation e = replacementRange.getEnd();
+    clang::LangOptions DefaultLangOptions;
+    size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(e, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(s);
+    ct::Replacement Rep(*SM, s, length, OS.str());
+    clang::FullSourceLoc fullSL(s, *SM);
+    insertReplacement(Rep, fullSL);
+  }
+  return true;
+}
+
 void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) {
   llcompat::insertReplacement(*replacements, rep);
   if (PrintStats) {
@@ -423,6 +487,24 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
     ).bind("cudaSharedIncompleteArrayVar"),
     this
   );
+  Finder->addMatcher(
+    mat::callExpr(
+      mat::isExpansionInMainFile(),
+      mat::callee(
+        mat::functionDecl(
+          mat::hasAnyName(
+            sCudaGetSymbolAddress,
+            sCudaGetSymbolSize,
+            sCudaMemcpyFromSymbol,
+            sCudaMemcpyFromSymbolAsync,
+            sCudaMemcpyToSymbol,
+            sCudaMemcpyToSymbolAsync
+          )
+        )
+      )
+    ).bind("cudaSymbolFuncCall"),
+    this
+  );
   Finder->addMatcher(
     mat::callExpr(
       mat::isExpansionInMainFile(),
@@ -560,5 +642,6 @@ void HipifyAction::ExecuteAction() {
 void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
   if (cudaLaunchKernel(Result)) return;
   if (cudaSharedIncompleteArrayVar(Result)) return;
+  if (cudaSymbolFuncCall(Result)) return;
   if (cudaDeviceFuncCall(Result)) return;
 }
diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h
index 208e6fb0b2..6c04e2c0cc 100644
--- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h
+++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h
@@ -71,6 +71,7 @@ public:
   bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result);
   bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result);
   bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
+  bool cudaSymbolFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
   // Called by the preprocessor for each include directive during the non-raw lexing pass.
   void InclusionDirective(clang::SourceLocation hash_loc,
                           const clang::Token &include_token,
diff --git a/projects/clr/hipamd/tests/hipify-clang/unit_tests/device/device_symbols.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/device/device_symbols.cu
new file mode 100644
index 0000000000..b58abeda46
--- /dev/null
+++ b/projects/clr/hipamd/tests/hipify-clang/unit_tests/device/device_symbols.cu
@@ -0,0 +1,152 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+
+/*
+Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda_runtime.h>
+#include <iostream>
+#include <assert.h>
+
+#define NUM 1024
+#define SIZE 1024 * 4
+
+__device__ int globalIn[NUM];
+__device__ int globalOut[NUM];
+
+__global__ void Assign(int* Out) {
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    Out[tid] = globalIn[tid];
+    globalOut[tid] = globalIn[tid];
+}
+
+__device__ __constant__ int globalConst[NUM];
+
+__global__ void checkAddress(int* addr, bool* out) {
+    *out = (globalConst == addr);
+}
+
+int main() {
+    int *A, *Am, *B, *Ad, *C, *Cm;
+    A = new int[NUM];
+    B = new int[NUM];
+    C = new int[NUM];
+    for (int i = 0; i < NUM; ++i) {
+        A[i] = -1 * i;
+        B[i] = 0;
+        C[i] = 0;
+    }
+    // CHECK: hipMalloc((void**)&Ad, SIZE);
+    cudaMalloc((void**)&Ad, SIZE);
+    // CHECK: hipHostMalloc((void**)&Am, SIZE);
+    cudaMallocHost((void**)&Am, SIZE);
+    // CHECK: hipHostMalloc((void**)&Cm, SIZE);
+    cudaMallocHost((void**)&Cm, SIZE);
+    for (int i = 0; i < NUM; ++i) {
+        Am[i] = -1 * i;
+        Cm[i] = 0;
+    }
+    // CHECK: hipStream_t stream = NULL;
+    cudaStream_t stream = NULL;
+    // CHECK: hipStreamCreate(&stream);
+    cudaStreamCreate(&stream);
+    // CHECK: hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream);
+    cudaMemcpyToSymbolAsync(globalIn, Am, SIZE, 0, cudaMemcpyHostToDevice, stream);
+    // CHECK: hipStreamSynchronize(stream);
+    cudaStreamSynchronize(stream);
+    // CHECK: hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    Assign<<<dim3(1, 1, 1), dim3(NUM, 1, 1)>>>(Ad);
+    // CHECK: hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
+    cudaMemcpy(B, Ad, SIZE, cudaMemcpyDeviceToHost);
+    // CHECK: hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
+    cudaMemcpyFromSymbolAsync(Cm, globalOut, SIZE, 0, cudaMemcpyDeviceToHost, stream);
+    // CHECK: hipStreamSynchronize(stream);
+    cudaStreamSynchronize(stream);
+    for (int i = 0; i < NUM; ++i) {
+        assert(Am[i] == B[i]);
+        assert(Am[i] == Cm[i]);
+    }
+    for (int i = 0; i < NUM; ++i) {
+        A[i] = -2 * i;
+        B[i] = 0;
+    }
+    // CHECK: hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice);
+    cudaMemcpyToSymbol(globalIn, A, SIZE, 0, cudaMemcpyHostToDevice);
+    // CHECK: hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    Assign<<<dim3(1, 1, 1), dim3(NUM, 1, 1)>>>(Ad);
+    // CHECK: hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
+    cudaMemcpy(B, Ad, SIZE, cudaMemcpyDeviceToHost);
+    // CHECK: hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost);
+    cudaMemcpyFromSymbol(C, globalOut, SIZE, 0, cudaMemcpyDeviceToHost);
+    for (int i = 0; i < NUM; ++i) {
+        assert(A[i] == B[i]);
+        assert(A[i] == C[i]);
+    }
+    for (int i = 0; i < NUM; ++i) {
+        A[i] = -3 * i;
+        B[i] = 0;
+    }
+    // CHECK: hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream);
+    cudaMemcpyToSymbolAsync(globalIn, A, SIZE, 0, cudaMemcpyHostToDevice, stream);
+    // CHECK: hipStreamSynchronize(stream);
+    cudaStreamSynchronize(stream);
+    // CHECK: hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
+    Assign<<<dim3(1, 1, 1), dim3(NUM, 1, 1)>>>(Ad);
+    // CHECK: hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
+    cudaMemcpy(B, Ad, SIZE, cudaMemcpyDeviceToHost);
+    // CHECK: hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
+    cudaMemcpyFromSymbolAsync(C, globalOut, SIZE, 0, cudaMemcpyDeviceToHost, stream);
+    // CHECK: hipStreamSynchronize(stream);
+    cudaStreamSynchronize(stream);
+    for (int i = 0; i < NUM; ++i) {
+        assert(A[i] == B[i]);
+        assert(A[i] == C[i]);
+    }
+    bool *checkOkD;
+    bool checkOk = false;
+    size_t symbolSize = 0;
+    int *symbolAddress;
+    // CHECK: hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst));
+    cudaGetSymbolSize(&symbolSize, globalConst);
+    // CHECK: hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst));
+    cudaGetSymbolAddress((void**) &symbolAddress, globalConst);
+    // CHECK: hipMalloc((void**)&checkOkD, sizeof(bool));
+    cudaMalloc((void**)&checkOkD, sizeof(bool));
+    // CHECK: hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
+    checkAddress<<<dim3(1, 1, 1), dim3(1, 1, 1)>>>(symbolAddress, checkOkD);
+    // CHECK: hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost);
+    cudaMemcpy(&checkOk, checkOkD, sizeof(bool), cudaMemcpyDeviceToHost);
+    // CHECK: hipFree(checkOkD);
+    cudaFree(checkOkD);
+    assert(checkOk);
+    assert(symbolSize == SIZE);
+    // CHECK: hipHostFree(Am);
+    cudaFreeHost(Am);
+    // CHECK: hipHostFree(Cm);
+    cudaFreeHost(Cm);
+    // CHECK: hipFree(Ad);
+    cudaFree(Ad);
+    delete[] A;
+    delete[] B;
+    delete[] C;
+    return 0;
+}
diff --git a/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
index 74bf4b5a0d..ece384f04b 100644
--- a/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
+++ b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
@@ -304,7 +304,7 @@ struct runner<curandStateSobol32_t>
         CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice));
 
         const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
-        // CHECK: hipLaunchKernelGGL((init_kernel), dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, directions, offset);
+        // CHECK: hipLaunchKernelGGL((init_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, directions, offset);
         init_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, directions, offset);
         // CHECK: CUDA_CALL(hipPeekAtLastError());
         // CHECK: CUDA_CALL(hipDeviceSynchronize());
@@ -329,7 +329,7 @@ struct runner<curandStateSobol32_t>
                   const Extra extra)
     {
         const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
-        // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
+        // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
         generate_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions, generate_func, extra);
     }
 };
diff --git a/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu
index f4fd05ba48..567de05e6e 100644
--- a/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu
+++ b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/poisson_api_example.cu
@@ -247,7 +247,7 @@ API_TYPE set_API_type()
 void settings()
 {
     add_cachiers(cashiers_load);
-    // CHECK: hipMemcpyToSymbol("cashiers_load", cashiers_load_h,
+    // CHECK: hipMemcpyToSymbol(HIP_SYMBOL("cashiers_load"), cashiers_load_h,
     // CHECK: HOURS * sizeof(int), 0, hipMemcpyHostToDevice);
     cudaMemcpyToSymbol("cashiers_load", cashiers_load_h,
             HOURS * sizeof(int), 0, cudaMemcpyHostToDevice);

From 4868880d033c9993b851a862440f5fdb846ff6ca Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Fri, 20 Sep 2019 09:16:10 +0300
Subject: [PATCH 6/9] [HIPIFY][doc] Update README.md due to LLVM 9.0.0 release

+ LLVM 9.0.0 is latest stable release, no patches are needed
+ The latest CUDA 10.1 Update 2 is supported
+ Tested on Windows and Linux


[ROCm/clr commit: 2a27572463aa1a7e2ad768e1d73cfc4b0316dac5]
---
 projects/clr/hipamd/hipify-clang/README.md | 107 ++++++++++-----------
 1 file changed, 53 insertions(+), 54 deletions(-)

diff --git a/projects/clr/hipamd/hipify-clang/README.md b/projects/clr/hipamd/hipify-clang/README.md
index fdb9e080f7..eb64d94367 100644
--- a/projects/clr/hipamd/hipify-clang/README.md
+++ b/projects/clr/hipamd/hipify-clang/README.md
@@ -32,11 +32,10 @@
 ## <a name="dependencies"></a> Dependencies
 
 `hipify-clang` requires:
-1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**6.0.1**](http://releases.llvm.org/download.html#6.0.1) on **Windows**, and [**8.0.1**](http://releases.llvm.org/download.html#8.0.1) on **Linux**.
 
-2. **CUDA** at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**9.0**](https://developer.nvidia.com/cuda-90-download-archive) on **Windows**, and [**10.0**](https://developer.nvidia.com/cuda-10.0-download-archive) on **Linux**.
+1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**9.0.0**](http://releases.llvm.org/download.html#9.0.0).
 
-If the target CUDA is [9.1](https://developer.nvidia.com/cuda-91-download-archive), [9.2](https://developer.nvidia.com/cuda-92-download-archive) or [10.0](https://developer.nvidia.com/cuda-10.0-download-archive), to work on Windows you may apply patches* for LLVM: [7.0.0](patches/patch_for_clang_7.0.0_bug_38811.zip), [7.0.1](patches/patch_for_clang_7.0.1_bug_38811.zip), [7.1.0](patches/patch_for_clang_7.1.0_bug_38811.zip), [8.0.0](patches/patch_for_clang_8.0.0_bug_38811.zip), [8.0.1](patches/patch_for_clang_8.0.1_bug_38811.zip).
+2. [**CUDA**](https://developer.nvidia.com/cuda-downloads) at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**10.1 Update 2**](https://developer.nvidia.com/cuda-downloads).
 
 | **LLVM release version**                                   | **CUDA latest supported version**                                   | **Windows**  | **Linux** |
 |:----------------------------------------------------------:|:-------------------------------------------------------------------:|:------------:|:---------:|
@@ -50,20 +49,20 @@ If the target CUDA is [9.1](https://developer.nvidia.com/cuda-91-download-archiv
 | [5.0.1](http://releases.llvm.org/download.html#5.0.1)      | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive)    | +            | +         |
 | [5.0.2](http://releases.llvm.org/download.html#5.0.2)      | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive)    | +            | +         |
 | [6.0.0](http://releases.llvm.org/download.html#6.0.0)      | [9.0](https://developer.nvidia.com/cuda-90-download-archive)        | +            | +         |
-| [**6.0.1**](http://releases.llvm.org/download.html#6.0.1)  | [**9.0**](https://developer.nvidia.com/cuda-90-download-archive)    | + <br/> **LATEST STABLE RELEASE** | + |
+| [6.0.1](http://releases.llvm.org/download.html#6.0.1)      | [9.0](https://developer.nvidia.com/cuda-90-download-archive)        | +            | +         |
 | [7.0.0](http://releases.llvm.org/download.html#7.0.0)      | [9.2](https://developer.nvidia.com/cuda-92-download-archive)        | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
 | [7.0.1](http://releases.llvm.org/download.html#7.0.1)      | [9.2](https://developer.nvidia.com/cuda-92-download-archive)        | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
 | [7.1.0](http://releases.llvm.org/download.html#7.1.0)      | [9.2](https://developer.nvidia.com/cuda-92-download-archive)        | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.1.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
-| [8.0.0](http://releases.llvm.org/download.html#8.0.0)      | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + |
-| [**8.0.1**](http://releases.llvm.org/download.html#8.0.1)  | [**10.0**](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.1_bug_38811.zip)*</br> | + <br/> **LATEST STABLE RELEASE** |
-| 9.0.0                                                      | 10.1 | <br/> LLVM 9.0.0 <br/> is not yet released | <br/> LLVM 9.0.0 <br/> is not yet released |
+| [8.0.0](http://releases.llvm.org/download.html#8.0.0)      | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive)     | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + |
+| [8.0.1](http://releases.llvm.org/download.html#8.0.1)      | [10.0](https://developer.nvidia.com/cuda-10.0-download-archive)     | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.1_bug_38811.zip)*</br> | + |
+| [**9.0.0**](http://releases.llvm.org/download.html#9.0.0)  | [**10.1**](https://developer.nvidia.com/cuda-downloads)             | + <br/> **LATEST STABLE RELEASE** | + <br/> **LATEST STABLE RELEASE** |
 
 `*` Download the patch and unpack it into your LLVM distributive directory; a few header files will be overwritten; rebuilding of LLVM is not needed.
 
 In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
 
 Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set
-[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.12/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.1\dist`
+[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.12/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\9.0.0\dist`
 
 ## <a name="build-and-install"></a> Build and install
 
@@ -83,10 +82,10 @@ cmake \
 
 make -j install
 ```
-On Windows, the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make.`
+On Windows, the following option should be specified for `cmake` at first place: `-G "Visual Studio 16 2019 Win64"`; the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make.`
 
 Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested; `LLVM+CLANG` should be built in `Debug` mode as well.
-64-bit build mode `-Thost=x64` is supported as well; `LLVM+CLANG` should be built in 64-bit mode as well.
+64-bit build mode (`-Thost=x64` on Windows) is also supported; `LLVM+CLANG` should be built in 64-bit mode as well.
 
 The binary can then be found at `./dist/bin/hipify-clang`.
 
@@ -97,8 +96,8 @@ The binary can then be found at `./dist/bin/hipify-clang`.
 **LLVM+CLANG should be built from sources, pre-built binaries are not exhaustive for testing.**
 
 To run it:
-1. Download [`LLVM`](http://releases.llvm.org/6.0.1/llvm-6.0.1.src.tar.xz)+[`CLANG`](http://releases.llvm.org/6.0.1/cfe-6.0.1.src.tar.xz) sources.
-2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html):
+1. download [`LLVM`](http://releases.llvm.org/9.0.0/llvm-9.0.0.src.tar.xz)+[`CLANG`](http://releases.llvm.org/9.0.0/cfe-9.0.0.src.tar.xz) sources; 
+2. build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html):
    ```shell
    cd llvm
    mkdir build dist
@@ -119,7 +118,7 @@ To run it:
      - **Windows**:
    ```shell
         cmake \
-         -G "Visual Studio 15 2017 Win64" \
+         -G "Visual Studio 16 2019 Win64" \
          -DCMAKE_INSTALL_PREFIX=../dist \
          -DLLVM_SOURCE_DIR=../llvm \
          -DLLVM_TARGETS_TO_BUILD="NVPTX" \
@@ -128,26 +127,26 @@ To run it:
          ../llvm
    ```
 
-&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp; Run `Visual Studio 15 2017`, open the generated `LLVM.sln`, build all, build project `INSTALL`.
+&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp; Run `Visual Studio 16 2019`, open the generated `LLVM.sln`, build all, build project `INSTALL`.
 
 
 3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.0 is installed.
 
     * Having multiple CUDA installations to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
 
-        - Linux: `-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.0`
+        - Linux: `-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1`
 
-        - Windows: `-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`
+        - Windows: `-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1"`
 
-          `-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"`
+          `-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1"`
 
 4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of the version corresponding to CUDA's version is installed.
 
     * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option:
 
-        - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.3.30`
+        - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.3.30`
 
-        - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.3.30`
+        - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.3.30`
 
 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
 
@@ -155,15 +154,15 @@ To run it:
 
     * Install `lit` into `python`:
 
-        - Linux: `python /srv/git/LLVM/8.0.1/llvm/utils/lit/setup.py install`
+        - Linux: `python /srv/git/LLVM/9.0.0/llvm/utils/lit/setup.py install`
 
-        - Windows: `python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`
+        - Windows: `python f:/LLVM/9.0.0/llvm/utils/lit/setup.py install`
 
     * Starting with LLVM 6.0.1 path to `llvm-lit` python script should be specified by the `LLVM_EXTERNAL_LIT` option:
 
-        - Linux: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/8.0.1/build/bin/llvm-lit`
+        - Linux: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit`
 
-        - Windows: `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`
+        - Windows: `-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py`
 
 7. Set `HIPIFY_CLANG_TESTS` option turned on: `-DHIPIFY_CLANG_TESTS=1`.
 
@@ -175,7 +174,7 @@ To run it:
 
      - Linux: `make test-hipify`.
 
-     - Windows: run `Visual Studio 15 2017`, open the generated `hipify-clang.sln`, build project `test-hipify`.
+     - Windows: run `Visual Studio 16 2019`, open the generated `hipify-clang.sln`, build project `test-hipify`.
 
 ### <a name="linux"></a >Linux
 
@@ -183,7 +182,7 @@ On Linux the following configurations are tested:
 
 Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.3.30
 
-Ubuntu 16-18: LLVM 8.0.0 - 8.0.1, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.6.3.30
+Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.3.30
 
 Build system for the above configurations:
 
@@ -196,16 +195,16 @@ cmake
  -DHIPIFY_CLANG_TESTS=1 \
  -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_INSTALL_PREFIX=../dist \
- -DCMAKE_PREFIX_PATH=/srv/git/LLVM/8.0.1/dist \
- -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.0 \
- -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.3.30 \
- -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/8.0.1/build/bin/llvm-lit \
+ -DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.0/dist \
+ -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1 \
+ -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.3.30 \
+ -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \
  ..
 ```
 *A corresponding successful output:*
 ```shell
--- The C compiler identification is GNU 5.4.0
--- The CXX compiler identification is GNU 5.4.0
+-- The C compiler identification is GNU 7.4.0
+-- The CXX compiler identification is GNU 7.4.0
 -- Check for working C compiler: /usr/bin/cc
 -- Check for working C compiler: /usr/bin/cc -- works
 -- Detecting C compiler ABI info
@@ -218,14 +217,14 @@ cmake
 -- Detecting CXX compiler ABI info - done
 -- Detecting CXX compile features
 -- Detecting CXX compile features - done
--- Found LLVM 8.0.1:
---    - CMake module path: /srv/git/LLVM/8.0.1/dist/lib/cmake/llvm
---    - Include path     : /srv/git/LLVM/8.0.1/dist/include
---    - Binary path      : /srv/git/LLVM/8.0.1/dist/bin
+-- Found LLVM 9.0.0:
+--    - CMake module path: /srv/git/LLVM/9.0.0/dist/lib/cmake/llvm
+--    - Include path     : /srv/git/LLVM/9.0.0/dist/include
+--    - Binary path      : /srv/git/LLVM/9.0.0/dist/bin
 -- Linker detection: GNU ld
 -- Found PythonInterp: /usr/bin/python2.7 (found suitable version "2.7.12", minimum required is "2.7")
 -- Found lit: /usr/local/bin/lit
--- Found FileCheck: /srv/git/LLVM/8.0.1/dist/bin/FileCheck
+-- Found FileCheck: /srv/git/LLVM/9.0.0/dist/bin/FileCheck
 -- Looking for pthread.h
 -- Looking for pthread.h - found
 -- Looking for pthread_create
@@ -235,7 +234,7 @@ cmake
 -- Looking for pthread_create in pthread
 -- Looking for pthread_create in pthread - found
 -- Found Threads: TRUE
--- Found CUDA: /usr/local/cuda-10.0 (found version "10.0")
+-- Found CUDA: /usr/local/cuda-10.1 (found version "10.1")
 -- Configuring done
 -- Generating done
 -- Build files have been written to: /srv/git/HIP/hipify-clang/build
@@ -247,8 +246,8 @@ make test-hipify
 ```shell
 Running HIPify regression tests
 ========================================
-CUDA 10.0 - will be used for testing
-LLVM 8.0.1 - will be used for testing
+CUDA 10.1 - will be used for testing
+LLVM 9.0.0 - will be used for testing
 x86_64 - Platform architecture
 Linux 5.2.0 - Platform OS
 64 - hipify-clang binary bitness
@@ -326,38 +325,38 @@ LLVM 5.0.0 - 5.0.2, CUDA 8.0,  cudnn-5.1.10 - cudnn-7.1.4.18
 
 LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.3.30
 
-LLVM 7.0.0 - 8.0.1 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.6.3.30
+LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn-7.0.5.15 - cudnn-7.6.3.30
 
 Build system for the above configurations:
 
 Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.5).
 
-Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`:
+Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`:
 
 ```shell
 cmake
- -G "Visual Studio 15 2017 Win64" \
+ -G "Visual Studio 16 2019 Win64" \
  -DHIPIFY_CLANG_TESTS=1 \
  -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_INSTALL_PREFIX=../dist \
- -DCMAKE_PREFIX_PATH=f:/LLVM/6.0.1/dist \
- -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \
- -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \
- -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.3.30 \
- -DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py \
+ -DCMAKE_PREFIX_PATH=f:/LLVM/9.0.0/dist \
+ -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1" \
+ -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1" \
+ -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.3.30 \
+ -DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \
  -Thost=x64
  ..
 ```
 *A corresponding successful output:*
 ```shell
--- Found LLVM 6.0.1:
---    - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm
---    - Include path     : F:/LLVM/6.0.1/dist/include
---    - Binary path      : F:/LLVM/6.0.1/dist/bin
+-- Found LLVM 9.0.0:
+--    - CMake module path: F:/LLVM/9.0.0/dist/lib/cmake/llvm
+--    - Include path     : F:/LLVM/9.0.0/dist/include
+--    - Binary path      : F:/LLVM/9.0.0/dist/bin
 -- Found PythonInterp: C:/Program Files/Python37/python.exe (found suitable version "3.7.4", minimum required is "3.6")
 -- Found lit: C:/Program Files/Python36/Scripts/lit.exe
--- Found FileCheck: F:/LLVM/6.0.1/dist/bin/FileCheck.exe
--- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0")
+-- Found FileCheck: F:/LLVM/9.0.0/dist/bin/FileCheck.exe
+-- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1 (found version "10.1")
 -- Configuring done
 -- Generating done
 -- Build files have been written to: f:/HIP/hipify-clang/build
@@ -370,7 +369,7 @@ To process a file, `hipify-clang` needs access to the same headers that would be
 For example:
 
 ```shell
-./hipify-clang square.cu --cuda-path=/usr/local/cuda-10.0 -I /usr/local/cuda-10.0/samples/common/inc
+./hipify-clang square.cu --cuda-path=/usr/local/cuda-10.1 -I /usr/local/cuda-10.1/samples/common/inc
 ```
 
 `hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you

From ef23cdda4198f650b3c407126155d97114a9d5a3 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Sat, 21 Sep 2019 07:33:17 +0300
Subject: [PATCH 7/9] [HIPIFY][#1435][perl] Add HIP_SYMBOL wrapper to the
 templated Device Symbol argument of the following functions:

cudaMemcpyToSymbol, cudaMemcpyToSymbolAsync, cudaGetSymbolSize, cudaGetSymbolAddress, cudaMemcpyFromSymbol, cudaMemcpyFromSymbolAsync

+ Perl part of [#1441]
+ Implement function generateSymbolFunctions() in hipify-clang for that purposes
+ Update hipify-perl

TODO: Eliminate dim3() issue in hipify-perl as well


[ROCm/clr commit: 72a658295f143e292064d3f40b84436e127a9cf6]
---
 projects/clr/hipamd/bin/hipify-perl           | 26 ++++++++++++
 .../hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp | 41 ++++++++++++++++++-
 .../hipify-clang/src/CUDA2HIP_Scripting.h     |  3 ++
 .../hipamd/hipify-clang/src/HipifyAction.cpp  |  4 +-
 4 files changed, 70 insertions(+), 4 deletions(-)

diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl
index 8094b96add..02a540b133 100755
--- a/projects/clr/hipamd/bin/hipify-perl
+++ b/projects/clr/hipamd/bin/hipify-perl
@@ -160,6 +160,8 @@ while (@ARGV) {
     undef $/; # Read whole file at once, so we can match newlines.
     while (<INFILE>)
     {
+        # chomp;
+        # next if /^(\s*(#.*)?)?$/;
         $ft{'error'} += s/\bcudaGetErrorName\b/hipGetErrorName/g;
         $ft{'error'} += s/\bcudaGetErrorString\b/hipGetErrorString/g;
         $ft{'error'} += s/\bcudaGetLastError\b/hipGetLastError/g;
@@ -1761,6 +1763,8 @@ while (@ARGV) {
             $ft{'kernel_func'} += countSupportedDeviceFunctions();
         }
 
+        $ft{'memory'} += transformSymbolFunctions();
+
         # Print it!
         # TODO - would like to move this code outside loop but it uses $_ which contains the whole file.
         unless ($no_output) {
@@ -1815,6 +1819,28 @@ if ($count_conversions) {
     }
 }
 
+sub transformSymbolFunctions
+{
+    my $m = 0;
+    foreach $func (
+        "hipMemcpyToSymbol",
+        "hipMemcpyToSymbolAsync"
+    )
+    {
+        $m += s/(?<!\/\/ CHECK: )($func)\s*\(\s*([^,]+)\s*,/$func\(HIP_SYMBOL\($2\),/g
+    }
+    foreach $func (
+        "hipGetSymbolAddress",
+        "hipGetSymbolSize",
+        "hipMemcpyFromSymbol",
+        "hipMemcpyFromSymbolAsync"
+    )
+    {
+        $m += s/(?<!\/\/ CHECK: )($func)\s*\(\s*([^,]+)\s*,\s*([^,\)]+)\s*(,\s*|\))\s*/$func\($2, HIP_SYMBOL\($3\)$4/g;
+    }
+    return $m;
+}
+
 sub countSupportedDeviceFunctions
 {
     my $m = 0;
diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
index 25705e5bd6..cc1a88aa9d 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
@@ -41,6 +41,42 @@ namespace perl {
   const std::string triple_space = double_space + space;
   const std::string sSub = "sub";
   const std::string sReturn_0 = "return 0;\n";
+  const std::string sReturn_m = "return $m;\n";
+  const std::string sForeach = "foreach $func (\n";
+  const std::string sMy = "my $m = 0;\n";
+
+  void generateSymbolFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
+    *perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << space << sMy;
+    std::string sCommon = space + sForeach;
+    *perlStreamPtr.get() << sCommon;
+    unsigned int count = 0;
+    for (auto& dsf : DeviceSymbolFunctions0) {
+      const auto found = CUDA_RENAMES_MAP().find(dsf);
+      if (found != CUDA_RENAMES_MAP().end()) {
+        *perlStreamPtr.get() << (count ? ",\n" : "") << double_space << "\"" << found->second.hipName.str() << "\"";
+        count++;
+      }
+    }
+    *perlStreamPtr.get() << "\n" << space << ")\n";
+    *perlStreamPtr.get() << space << "{\n";
+    *perlStreamPtr.get() << double_space << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(HIP_SYMBOL\\($2\\),/g\n";
+    *perlStreamPtr.get() << space << "}\n";
+    *perlStreamPtr.get() << sCommon;
+    count = 0;
+    for (auto& dsf : DeviceSymbolFunctions1) {
+      const auto found = CUDA_RENAMES_MAP().find(dsf);
+      if (found != CUDA_RENAMES_MAP().end()) {
+        *perlStreamPtr.get() << (count ? ",\n" : "") << double_space << "\"" << found->second.hipName.str() << "\"";
+        count++;
+      }
+    }
+    *perlStreamPtr.get() << "\n" << space << ")\n";
+    *perlStreamPtr.get() << space << "{\n";
+    *perlStreamPtr.get() << double_space << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, HIP_SYMBOL\\($3\\)$4/g;\n";
+    *perlStreamPtr.get() << space << "}\n";
+    *perlStreamPtr.get() << space << sReturn_m;
+    *perlStreamPtr.get() << "}\n";
+  }
 
   void generateDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
     unsigned int countUnsupported = 0;
@@ -59,7 +95,7 @@ namespace perl {
     std::stringstream subCountSupported;
     std::stringstream subWarnUnsupported;
     std::stringstream subCommon;
-    std::string sCommon = space + "my $m = 0;\n" + space + "foreach $func (\n";
+    std::string sCommon = space + sMy + space + sForeach;
     subCountSupported << "\n" << sSub << " countSupportedDeviceFunctions\n" << "{\n" << (countSupported ? sCommon : space + sReturn_0);
     subWarnUnsupported << "\n" << sSub << " warnUnsupportedDeviceFunctions\n" << "{\n" << (countUnsupported ? space + "my $line_num = shift;\n" + sCommon : space + sReturn_0);
     if (countSupported) {
@@ -85,7 +121,7 @@ namespace perl {
       subWarnUnsupported << triple_space << "print STDERR \"  warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
     }
     if (countSupported || countUnsupported) {
-      sCommon = double_space + "}\n" + space + "}\n" + space + "return $m;\n";
+      sCommon = double_space + "}\n" + space + "}\n" + space + sReturn_m;
     }
     if (countSupported) {
       subCountSupported << sCommon;
@@ -157,6 +193,7 @@ namespace perl {
         }
       }
     }
+    generateSymbolFunctions(perlStreamPtr);
     generateDeviceFunctions(perlStreamPtr);
     perlStreamPtr.get()->flush();
     bool ret = true;
diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Scripting.h b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Scripting.h
index c585a47d47..9469f25845 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Scripting.h
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Scripting.h
@@ -22,6 +22,9 @@ THE SOFTWARE.
 
 #pragma once
 
+extern std::set<std::string> DeviceSymbolFunctions0;
+extern std::set<std::string> DeviceSymbolFunctions1;
+
 namespace perl {
 
   bool generate(bool Generate = true);
diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
index 6cb75f8911..4acbb7e78e 100644
--- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp
@@ -43,12 +43,12 @@ const std::string sCudaGetSymbolAddress = "cudaGetSymbolAddress";
 const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol";
 const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync";
 
-const std::set<std::string> DeviceSymbolFunctions0 {
+std::set<std::string> DeviceSymbolFunctions0 {
   {sCudaMemcpyToSymbol},
   {sCudaMemcpyToSymbolAsync}
 };
 
-const std::set<std::string> DeviceSymbolFunctions1 {
+std::set<std::string> DeviceSymbolFunctions1 {
   {sCudaGetSymbolSize},
   {sCudaGetSymbolAddress},
   {sCudaMemcpyFromSymbol},

From fa6a69e37560a9232ec23e38d325f246d62f7f11 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Sun, 22 Sep 2019 22:34:07 +0300
Subject: [PATCH 8/9] [HIPIFY][perl] generateSymbolFunctions() small
 refactoring

[ROCm/clr commit: d6b0ab0fd24d0857f158b116f8726a20b3092be0]
---
 .../hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp | 40 ++++++++-----------
 1 file changed, 16 insertions(+), 24 deletions(-)

diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
index cc1a88aa9d..8f1d8b83e9 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
@@ -48,32 +48,24 @@ namespace perl {
   void generateSymbolFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
     *perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << space << sMy;
     std::string sCommon = space + sForeach;
-    *perlStreamPtr.get() << sCommon;
-    unsigned int count = 0;
-    for (auto& dsf : DeviceSymbolFunctions0) {
-      const auto found = CUDA_RENAMES_MAP().find(dsf);
-      if (found != CUDA_RENAMES_MAP().end()) {
-        *perlStreamPtr.get() << (count ? ",\n" : "") << double_space << "\"" << found->second.hipName.str() << "\"";
-        count++;
+    std::set<std::string> &funcSet = DeviceSymbolFunctions0;
+    for (int i = 0; i < 2; ++i) {
+      *perlStreamPtr.get() << sCommon;
+      if (i == 1) funcSet = DeviceSymbolFunctions1;
+      unsigned int count = 0;
+      for (auto& f : funcSet) {
+        const auto found = CUDA_RUNTIME_FUNCTION_MAP.find(f);
+        if (found != CUDA_RUNTIME_FUNCTION_MAP.end()) {
+          *perlStreamPtr.get() << (count ? ",\n" : "") << double_space << "\"" << found->second.hipName.str() << "\"";
+          count++;
+        }
       }
+      *perlStreamPtr.get() << "\n" << space << ")\n";
+      *perlStreamPtr.get() << space << "{\n" << double_space;
+      if (i ==0) *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(HIP_SYMBOL\\($2\\),/g\n";
+      else *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, HIP_SYMBOL\\($3\\)$4/g;\n";
+      *perlStreamPtr.get() << space << "}\n";
     }
-    *perlStreamPtr.get() << "\n" << space << ")\n";
-    *perlStreamPtr.get() << space << "{\n";
-    *perlStreamPtr.get() << double_space << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(HIP_SYMBOL\\($2\\),/g\n";
-    *perlStreamPtr.get() << space << "}\n";
-    *perlStreamPtr.get() << sCommon;
-    count = 0;
-    for (auto& dsf : DeviceSymbolFunctions1) {
-      const auto found = CUDA_RENAMES_MAP().find(dsf);
-      if (found != CUDA_RENAMES_MAP().end()) {
-        *perlStreamPtr.get() << (count ? ",\n" : "") << double_space << "\"" << found->second.hipName.str() << "\"";
-        count++;
-      }
-    }
-    *perlStreamPtr.get() << "\n" << space << ")\n";
-    *perlStreamPtr.get() << space << "{\n";
-    *perlStreamPtr.get() << double_space << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, HIP_SYMBOL\\($3\\)$4/g;\n";
-    *perlStreamPtr.get() << space << "}\n";
     *perlStreamPtr.get() << space << sReturn_m;
     *perlStreamPtr.get() << "}\n";
   }

From 6dfc9a99d65a6c37272566be068d90ddecb65739 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Sun, 22 Sep 2019 23:43:07 +0300
Subject: [PATCH 9/9] [HIPIFY][perl] CUDA2HIP_Perl code cleanup

[ROCm/clr commit: a13540b97a8dbc670836779ba57658f8fd43ab48]
---
 .../hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp | 97 +++++++------------
 1 file changed, 37 insertions(+), 60 deletions(-)

diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
index 8f1d8b83e9..26a72886c7 100644
--- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
@@ -36,9 +36,9 @@ using namespace llvm;
 
 namespace perl {
 
-  const std::string space = "    ";
-  const std::string double_space = space + space;
-  const std::string triple_space = double_space + space;
+  const std::string tab = "    ";
+  const std::string double_tab = tab + tab;
+  const std::string triple_tab = double_tab + tab;
   const std::string sSub = "sub";
   const std::string sReturn_0 = "return 0;\n";
   const std::string sReturn_m = "return $m;\n";
@@ -46,28 +46,25 @@ namespace perl {
   const std::string sMy = "my $m = 0;\n";
 
   void generateSymbolFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
-    *perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << space << sMy;
-    std::string sCommon = space + sForeach;
+    *perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << tab << sMy;
     std::set<std::string> &funcSet = DeviceSymbolFunctions0;
     for (int i = 0; i < 2; ++i) {
-      *perlStreamPtr.get() << sCommon;
+      *perlStreamPtr.get() << tab + sForeach;
       if (i == 1) funcSet = DeviceSymbolFunctions1;
       unsigned int count = 0;
       for (auto& f : funcSet) {
         const auto found = CUDA_RUNTIME_FUNCTION_MAP.find(f);
         if (found != CUDA_RUNTIME_FUNCTION_MAP.end()) {
-          *perlStreamPtr.get() << (count ? ",\n" : "") << double_space << "\"" << found->second.hipName.str() << "\"";
+          *perlStreamPtr.get() << (count ? ",\n" : "") << double_tab << "\"" << found->second.hipName.str() << "\"";
           count++;
         }
       }
-      *perlStreamPtr.get() << "\n" << space << ")\n";
-      *perlStreamPtr.get() << space << "{\n" << double_space;
+      *perlStreamPtr.get() << "\n" << tab << ")\n" << tab << "{\n" << double_tab;
       if (i ==0) *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(HIP_SYMBOL\\($2\\),/g\n";
       else *perlStreamPtr.get() << "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, HIP_SYMBOL\\($3\\)$4/g;\n";
-      *perlStreamPtr.get() << space << "}\n";
+      *perlStreamPtr.get() << tab << "}\n";
     }
-    *perlStreamPtr.get() << space << sReturn_m;
-    *perlStreamPtr.get() << "}\n";
+    *perlStreamPtr.get() << tab << sReturn_m << "}\n";
   }
 
   void generateDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
@@ -77,50 +74,43 @@ namespace perl {
     std::stringstream sUnsupported;
     for (auto& ma : CUDA_DEVICE_FUNC_MAP) {
       bool isUnsupported = Statistics::isUnsupported(ma.second);
-      (isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << double_space << "\"" << ma.first.str() << "\"";
-      if (isUnsupported) {
-        countUnsupported++;
-      } else {
-        countSupported++;
-      }
+      (isUnsupported ? sUnsupported : sSupported) << ((isUnsupported && countUnsupported) || (!isUnsupported && countSupported) ? ",\n" : "") << double_tab << "\"" << ma.first.str() << "\"";
+      if (isUnsupported) countUnsupported++;
+      else countSupported++;
     }
     std::stringstream subCountSupported;
     std::stringstream subWarnUnsupported;
     std::stringstream subCommon;
-    std::string sCommon = space + sMy + space + sForeach;
-    subCountSupported << "\n" << sSub << " countSupportedDeviceFunctions\n" << "{\n" << (countSupported ? sCommon : space + sReturn_0);
-    subWarnUnsupported << "\n" << sSub << " warnUnsupportedDeviceFunctions\n" << "{\n" << (countUnsupported ? space + "my $line_num = shift;\n" + sCommon : space + sReturn_0);
+    std::string sCommon = tab + sMy + tab + sForeach;
+    subCountSupported << "\n" << sSub << " countSupportedDeviceFunctions\n" << "{\n" << (countSupported ? sCommon : tab + sReturn_0);
+    subWarnUnsupported << "\n" << sSub << " warnUnsupportedDeviceFunctions\n" << "{\n" << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0);
     if (countSupported) {
-      subCountSupported << sSupported.str() << "\n" << space << ")\n";
+      subCountSupported << sSupported.str() << "\n" << tab << ")\n";
     }
     if (countUnsupported) {
-      subWarnUnsupported << sUnsupported.str() << "\n" << space << ")\n";
+      subWarnUnsupported << sUnsupported.str() << "\n" << tab << ")\n";
     }
     if (countSupported || countUnsupported) {
-      subCommon << space << "{\n";
-      subCommon << double_space << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
-      subCommon << 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";
-      subCommon << double_space << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
-      subCommon << double_space << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
-      subCommon << double_space << "if ($mt && !$mt_namespace) {\n";
-      subCommon << triple_space << "$m += $mt;\n";
+      subCommon << tab << "{\n";
+      subCommon << double_tab << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
+      subCommon << double_tab << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);\n";
+      subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      subCommon << double_tab << "if ($mt && !$mt_namespace) {\n";
+      subCommon << triple_tab << "$m += $mt;\n";
     }
     if (countSupported) {
       subCountSupported << subCommon.str();
     }
     if (countUnsupported) {
       subWarnUnsupported << subCommon.str();
-      subWarnUnsupported << triple_space << "print STDERR \"  warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
+      subWarnUnsupported << triple_tab << "print STDERR \"  warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
     }
     if (countSupported || countUnsupported) {
-      sCommon = double_space + "}\n" + space + "}\n" + space + sReturn_m;
-    }
-    if (countSupported) {
-      subCountSupported << sCommon;
-    }
-    if (countUnsupported) {
-      subWarnUnsupported << sCommon;
+      sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_m;
     }
+    if (countSupported) subCountSupported << sCommon;
+    if (countUnsupported) subWarnUnsupported << sCommon;
     subCountSupported << "}\n";
     subWarnUnsupported << "}\n";
     *perlStreamPtr.get() << subCountSupported.str();
@@ -128,19 +118,13 @@ namespace perl {
   }
 
   bool generate(bool Generate) {
-    if (!Generate) {
-      return true;
-    }
+    if (!Generate) return true;
     std::string dstPerlMap = OutputPerlMapFilename, dstPerlMapDir = OutputPerlMapDir;
-    if (dstPerlMap.empty()) {
-      dstPerlMap = "hipify-perl-map";
-    }
+    if (dstPerlMap.empty()) dstPerlMap = "hipify-perl-map";
     std::error_code EC;
     if (!dstPerlMapDir.empty()) {
       std::string sOutputPerlMapDirAbsPath = getAbsoluteDirectoryPath(OutputPerlMapDir, EC, "output hipify-perl map");
-      if (EC) {
-        return false;
-      }
+      if (EC) return false;
       dstPerlMap = sOutputPerlMapDirAbsPath + "/" + dstPerlMap;
     }
     SmallString<128> tmpFile;
@@ -153,18 +137,16 @@ namespace perl {
     std::unique_ptr<std::ostream> perlStreamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
     std::string sConv = "my $conversions = ";
     *perlStreamPtr.get() << "@statNames = (";
-    for (int i = 0; i < NUM_CONV_TYPES - 1; i++) {
+    for (int i = 0; i < NUM_CONV_TYPES - 1; ++i) {
       *perlStreamPtr.get() << "\"" << counterNames[i] << "\", ";
       sConv += "$ft{'" + std::string(counterNames[i]) + "'} + ";
     }
     *perlStreamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");\n\n";
     *perlStreamPtr.get() << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};\n\n";
-    for (int i = 0; i < NUM_CONV_TYPES; i++) {
+    for (int i = 0; i < NUM_CONV_TYPES; ++i) {
       if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
         for (auto& ma : CUDA_INCLUDE_MAP) {
-          if (Statistics::isUnsupported(ma.second)) {
-            continue;
-          }
+          if (Statistics::isUnsupported(ma.second)) continue;
           if (i == ma.second.type) {
             std::string sCUDA = ma.first.str();
             std::string sHIP = ma.second.hipName.str();
@@ -173,12 +155,9 @@ namespace perl {
             *perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;\n";
           }
         }
-      }
-      else {
+      } else {
         for (auto& ma : CUDA_RENAMES_MAP()) {
-          if (Statistics::isUnsupported(ma.second)) {
-            continue;
-          }
+          if (Statistics::isUnsupported(ma.second)) continue;
           if (i == ma.second.type) {
             *perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;\n";
           }
@@ -194,9 +173,7 @@ namespace perl {
       llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstPerlMap << "\n";
       ret = false;
     }
-    if (!SaveTemps) {
-      sys::fs::remove(tmpFile);
-    }
+    if (!SaveTemps) sys::fs::remove(tmpFile);
     return ret;
   }
 }