From 2cd2afa84b1a36778a6f72726f12e8b31338def0 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 16 Sep 2019 17:36:55 +0300 Subject: [PATCH] [HIPIFY][perl][fix] Treat ::device_function as a device function + Do not treat somenamespace::device_function_name as a device function + Fix generation of warnUnsupportedDeviceFunctions function in hipify-clang + Update hipify-perl based on hipify-clang -perl generation + Update device test math_functions.cu for hipify-perl [Restrictions] - hipify-perl is yet unable to handle function declarations in user namespaces - hipify-perl is yet unable to handle using directive [ROCm/hip commit: 4f59ec25fef86e44adcf47ad4b1e25b97e21d1f5] --- projects/hip/bin/hipify-perl | 10 ++++++---- .../hip/hipify-clang/src/CUDA2HIP_Perl.cpp | 10 ++++++---- .../unit_tests/device/math_functions.cu | 18 +++++++++++++++--- 3 files changed, 27 insertions(+), 11 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 40583c41c3..b5fc365955 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1956,11 +1956,13 @@ sub warnUnsupportedDeviceFunctions "umul24" ) { - # match math at the beginning of a word, but not if it already has a namespace qualifier ('::') : - my $mt = m/[:]?[:]?\b($func)\b(\w*\()/g; - if ($mt) { + # match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...)); + # function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier); + my $mt_namespace = m/(\w+)::($func)\s*\(\s*.*\s*\)/g; + my $mt = m/($func)\s*\(\s*.*\s*\)/g; + if ($mt && !$mt_namespace) { $m += $mt; - print STDERR " warning: $fileName:#$line_num : unsupported device function : $_\n"; + print STDERR " warning: $fileName:$line_num: unsupported device function \"$func\": $_\n"; } } return $m; diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index e40b0df379..75afc0a534 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -53,11 +53,13 @@ namespace perl { *perlStreamPtr.get() << "\nsub warnUnsupportedDeviceFunctions\n" << "{\n" << space << "my $line_num = shift;\n" << space << "my $m = 0;\n" << space << "foreach $func (\n"; *perlStreamPtr.get() << sUnsupported.str() << "\n" << space << ")\n"; *perlStreamPtr.get() << space << "{\n"; - *perlStreamPtr.get() << double_space << "# match math at the beginning of a word, but not if it already has a namespace qualifier ('::') :\n"; - *perlStreamPtr.get() << double_space << "my $mt = m/[:]?[:]?\\b($func)\\b(\\w*\\()/g;\n"; - *perlStreamPtr.get() << double_space << "if ($mt) {\n"; + *perlStreamPtr.get() << double_space << "# match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));\n"; + *perlStreamPtr.get() << double_space << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);\n"; + *perlStreamPtr.get() << double_space << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n"; + *perlStreamPtr.get() << double_space << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n"; + *perlStreamPtr.get() << double_space << "if ($mt && !$mt_namespace) {\n"; *perlStreamPtr.get() << triple_space << "$m += $mt;\n"; - *perlStreamPtr.get() << triple_space << "print STDERR \" warning: $fileName:#$line_num : unsupported device function : $_\\n\";\n"; + *perlStreamPtr.get() << triple_space << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n"; *perlStreamPtr.get() << double_space << "}\n"; *perlStreamPtr.get() << space << "}\n"; *perlStreamPtr.get() << space << "return $m;\n"; diff --git a/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu b/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu index 3bc1c1e51d..c833ca0182 100644 --- a/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu +++ b/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu @@ -1,5 +1,5 @@ // RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args -// Test to warn only on device functions umin and umax as unsupported, but not on user defined ones. +// Synthetic test to warn only on device functions umin and umax as unsupported, but not on user defined ones. // ToDo: change lit testing in order to parse the output. #define LEN 1024 @@ -8,9 +8,11 @@ #include namespace my { + // user defined function unsigned int umin(unsigned int arg1, unsigned int arg2) { return (arg1 < arg2) ? arg1 : arg2; } + // user defined function unsigned int umax(unsigned int arg1, unsigned int arg2) { return (arg1 > arg2) ? arg1 : arg2; } @@ -18,8 +20,16 @@ namespace my { __global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsigned int u2) { - unsigned int _umin = umin(u1, u2); - unsigned int _umax = umax(u1, u2); + // device function call (warn if unsupported) + unsigned int _umin = umin ( u1, u2 ); + // device function call (warn if unsupported) + unsigned int _umax = umax ( u1, u2 ); + // device function call (warn if unsupported) + unsigned int _umin_global = ::umin ( u1, u2 ); + // device function call (warn if unsupported) + unsigned int _umax_global = ::umax(u1, u2); + if (_umin != _umin_global) return; + if (_umax != _umax_global) return; int i = threadIdx.x; A[i] = i + _umin; B[i] = i + _umax; @@ -29,7 +39,9 @@ __global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsig int main() { unsigned int u1 = 33; unsigned int u2 = 34; + // user defined function call unsigned int _min = my::umin(u1, u2); + // user defined function call unsigned int _max = my::umax(u1, u2); float *A, *B, *C; // CHECK: hipMalloc((void**)&A, SIZE);