Merge 'master' into 'amd-master'

Change-Id: I59f0b6a69cd389ed93b01bb902f97f7033a3499c

bin/hipify-perl

@@ -1,4 +1,5 @@
-#!/usr/bin/perl  -w
+#!/usr/bin/perl -w
+
 ##
 # Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
 #
@@ -20,60 +21,55 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 ##
-#usage hipify-perl [OPTIONS] INPUT_FILE
-use Getopt::Long;
 
-my $warn_whitelist ="";
+#usage hipify-perl [OPTIONS] INPUT_FILE
+
+use Getopt::Long;
+my $whitelist = "";
+
 GetOptions(
-      "print-stats" => \$print_stats                     # print the command-line, like a header.
+      "examine" => \$examine                  # Combines -no-output and -print-stats options.
-    , "count-conversions" => \$count_conversions         # count conversions.
+    , "inplace" => \$inplace                  # Modify input file inplace, replacing input with hipified output, save backup in .prehip file.
-    , "quiet-warnings" => \$quiet_warnings               # don't print warnings on unknown CUDA functions.
+    , "no-output" => \$no_output              # Don't write any translated output to stdout.
-    , "warn-whitelist=s"=> \$warn_whitelist
+    , "print-stats" => \$print_stats          # Print translation statistics.
-    , "no-output" => \$no_output                         # don't write any translated output to stdout.
+    , "quiet-warnings" => \$quiet_warnings    # Don't print warnings on unknown CUDA functions.
-    , "inplace" => \$inplace                             # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file.
+    , "whitelist=s" => \$whitelist            # TODO: test it beforehand
-                                                         # If .prehip file exists, use that as input to hip.
-    , "n" => \$n                                         # combination of print_stats + no-output.
 );
 
-$print_stats = 1 if $n;
+$print_stats = 1 if $examine;
-$no_output   = 1 if $n;
+$no_output = 1 if $examine;
 
-# These uses of cuda[A-Z] are commonly used in CUDA code but don't actually map to any CUDA API:
+# Whitelist of cuda[A-Z] identifiers, which are commonly used in CUDA sources but don't map to any CUDA API:
-# TODO - use a hash lookup for these.
+@whitelist = (
-@warn_whitelist = (
+      "cudaCloverField"
-         "cudaDevice"
+    , "cudaColorSpinorField"
-        ,"cudaDevice_t"
+    , "cudaDevice"
-        ,"cudaIDs"
+    , "cudaDeviceId"
-        ,"cudaGridDim"
+    , "cudaDevice_t"
-        ,"cudaDimGrid"
+    , "cudaDevices"
-        ,"cudaDimBlock"
+    , "cudaDimBlock"
-        ,"cudaDeviceId"
+    , "cudaDimGrid"
-        ,"cudaDevices",
+    , "cudaFatLink"
-        ,"cudaGradOutput",
+    , "cudaGauge"
-        ,"cudaInput",
+    , "cudaGaugeField"
-        ,"cudaOutput",
+    , "cudaGradInput"
-        ,"cudaGradInput",
+    , "cudaGradOutput"
-        ,"cudaIndices",
+    , "cudaGridDim"
-        ,"cudaGaugeField"
+    , "cudaIDs"
-        ,"cudaMom"
+    , "cudaInGauge"
-        ,"cudaGauge"
+    , "cudaIndices"
-        ,"cudaInGauge"
+    , "cudaInput"
-        ,"cudaColorSpinorField"
+    , "cudaMom"
-        ,"cudaSiteLink"
+    , "cudaOutput"
-        ,"cudaFatLink"
+    , "cudaParam"
-        ,"cudaStaple"
+    , "cudaSiteLink"
-        ,"cudaCloverField"
+    , "cudaStaple"
-        ,"cudaParam"
+);
-    );
-#print "WW=@warn_whitelist\n";
 
-# Allow users to add their own functions.
+push(@whitelist, split(',', $whitelist));
-push (@warn_whitelist, split(',',$warn_whitelist));
 
-#Stats tracking code:
+@statNames = ("error", "init", "version", "device", "context", "module", "memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "device_function", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "extern_shared", "kernel_launch");
-@statNames = ("error", "init", "version", "device", "context", "module", "memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "kernel_func", "extern_shared", "kern_launch");
 
-#Compute total of all individual counts:
 sub totalStats {
     my %count = %{ shift() };
     my $total = 0;
@@ -90,32 +86,30 @@ sub printStats {
     my $warnings  = shift();
     my $loc       = shift();
     my $total     = totalStats(\%counts);
-    printf STDERR "%s %d CUDA->HIP refs( ", $label, $total;
+    printf STDERR "%s %d CUDA->HIP refs ( ", $label, $total;
     foreach $stat (@statNames) {
         printf STDERR "%s:%d ", $stat, $counts{$stat};
     }
-    printf STDERR ") warn:%d LOC:%d", $warnings, $loc;
+    printf STDERR ")\n  warn:%d LOC:%d", $warnings, $loc;
-}
+};
 
-# Add adder stats to dest. Used to add stats for current file to a running total for all files:
 sub addStats {
-    my $dest_ref  =  shift();
+    my $dest_ref  = shift();
-    my %adder = %{ shift() };
+    my %adder     = %{ shift() };
     foreach $key (keys %adder) {
         $dest_ref->{$key} += $adder{$key};
-        #printf ("D{$key} += %d => %d\n", $adder{$key}, $dest{$key});
     }
-}
+};
 
 sub clearStats {
-    my $dest_ref  =  shift() ;
+    my $dest_ref  = shift();
     my @statNames = @{ shift() };
-    foreach $stat (@statNames) {
+    foreach $stat(@statNames) {
-        $dest_ref->{$stat}  = 0;
+        $dest_ref->{$stat} = 0;
     }
-}
+};
 
-# count of transforms in all files:
+# Count of transforms in all files
 my %tt;
 clearStats(\%tt, \@statNames);
 $Twarnings = 0;
@@ -146,22 +140,19 @@ while (@ARGV) {
         open(INFILE,"<", $fileName) or die "error: could not open $fileName";
         $OUTFILE = STDOUT;
     }
-    # Note : \b is used in perl to indicate the start of a word - typically that is what we want in this case:
+    # Note : \b is used in perl to indicate the start of a word
-    # count of transforms in this file, init to 0 here:
+    # Count of transforms in this file
     my %ft;
     clearStats(\%ft, \@statNames);
     my $countIncludes = 0;
-    my $countKeywords = 0; # keywords like __global__, __shared__ - not converted by hipify-perl, but counted here.
+    my $countKeywords = 0;
     my $warnings = 0; 
-    my $warningsCublas = 0;
+    my %warningTags;
-    my $warningsCurand = 0;
-    my %warningTags;  # hash with counts of particular unknown keywords.
     my $lineCount = 0;
-    undef $/; # Read whole file at once, so we can match newlines.
+    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;
@@ -1651,13 +1642,11 @@ while (@ARGV) {
         # CUDA extern __shared__ syntax
         # Note these only work if declaration is on a single line.
         {
-            # match uses ? for <.*> which will be unitialized if this is not present in launch syntax.
+            # Match uses ? for <.*> which will be unitialized if this is not present in launch syntax
             no warnings qw/uninitialized/;
             my $k = 0;
-            # Match extern __shared__ type foo[]; syntax
             # Replace as HIP_DYNAMIC_SHARED() macro
-            $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g;
+            # Match patterns for the below regular expression:
-            # test patterns for the regular expression above:
             #'extern __shared__ double foo[];'
             #'extern __shared__ unsigned int foo[];'
             #'extern volatile __shared__ double foo[];'
@@ -1668,12 +1657,12 @@ while (@ARGV) {
             #'extern __shared__ blah<T>::type s[];'
             #'extern __shared__ typename mapper<Float>::type s_data[];'
             #'extern __attribute__((used)) __shared__ typename mapper<Float>::type s_data[];'
+            $k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g;
             $ft{'extern_shared'} += $k;
         }
-
         # CUDA Launch Syntax. Note these only work if launch is on a single line.
         {
-            # match uses ? for <.*> which will be unitialized if this is not present in launch syntax.
+            # Match uses ? for <.*> which will be unitialized if this is not present in launch syntax
             no warnings qw/uninitialized/;
             my $k = 0;
 
@@ -1708,35 +1697,32 @@ while (@ARGV) {
             $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g;
 
             if ($k) {
-                $ft{'kern_launch'} += $k;
+                $ft{'kernel_launch'} += $k;
                 $Tkernels{$1} ++;
             }
         }
-        if ($count_conversions) {
+        if ($print_stats) {
             while (/(\bhip[A-Z]\w+\b)/g) {
                 $convertedTags{$1}++;
-                #print STDERR "HIP: $1 : ", $translateTags{$1}, "\n";
             }
         }
-        # guess that we are in device code , or at least in a file that calls device code.
+        my $hasDeviceCode = $countKeywords + $ft{'device_function'};
-        # will almost certainly call one of the coordiante functions - could be fooled by clever macros but usually works:
-        my $hasDeviceCode = $countKeywords + $ft{'kernel_func'};
         unless ($quiet_warnings) {
-            #print STDERR "Check WARNINGs\n";
+            # Copy into array of lines, process line-by-line to show warnings
-            # copy into array of lines, process line-by-line to show warnings:
+            if ($hasDeviceCode or (/\bcu/) or (/\bCU_/) or (/\bCUDA_/) or (/<<<.*>>>/)) {
-            if ($hasDeviceCode or (/\bcuda/) or (/<<<.*>>>/) or (/(\bcublas[A-Z]\w+)/) or (/(\bcurand[A-Z]\w+)/) ) {
                 my @lines = split /\n/, $_;
-                my $tmp = $_;  # copies the whole file, could be a little smarter here...
+                # Copy the whole file
+                my $tmp = $_;
                 my $line_num = 0;
                 foreach (@lines) {
                     $line_num ++;
-                    # remove any whitelisted words:
+                    # Remove any whitelisted words
-                    foreach $w (@warn_whitelist) {
+                    foreach $w (@whitelist) {
                         s/\b$w\b/ZAP/
                     }
                     my $tag;
                     if ((/(\bcuda[A-Z]\w+)/) or (/<<<.*>>>/)) {
-                        # flag any remaining code that look like cuda API calls, may want to add these to hipify
+                        # Flag any remaining code that look like cuda API calls: may want to add these to hipify
                         $tag = (defined $1) ? $1 : "Launch";
                     }
                     if (defined $tag) {
@@ -1751,26 +1737,18 @@ while (@ARGV) {
                 $_ = $tmp;
             }
         }
-
-        # To limit bogus translations, try to make sure we are in a kernel:
         if ($hasDeviceCode > 0) {
-            $ft{'kernel_func'} += countSupportedDeviceFunctions();
+            $ft{'device_function'} += countSupportedDeviceFunctions();
         }
-
         transformHostFunctions();
-
+        # TODO: would like to move this code outside loop but it uses $_ which contains the whole file
-        # Print it!
-        # TODO - would like to move this code outside loop but it uses $_ which contains the whole file.
         unless ($no_output) {
             my $apiCalls   = $ft{'error'} + $ft{'init'} + $ft{'version'} + $ft{'device'} + $ft{'context'} + $ft{'module'} + $ft{'memory'} + $ft{'addressing'} + $ft{'stream'} + $ft{'event'} + $ft{'external_resource_interop'} + $ft{'stream_memory'} + $ft{'execution'} + $ft{'graph'} + $ft{'occupancy'} + $ft{'texture'} + $ft{'surface'} + $ft{'peer'} + $ft{'graphics'} + $ft{'profiler'} + $ft{'openGL'} + $ft{'D3D9'} + $ft{'D3D10'} + $ft{'D3D11'} + $ft{'VDPAU'} + $ft{'EGL'} + $ft{'thread'} + $ft{'complex'} + $ft{'library'} + $ft{'device_library'} + $ft{'include'} + $ft{'include_cuda_main_header'} + $ft{'type'} + $ft{'literal'} + $ft{'numeric_literal'} + $ft{'define'};
-            my $kernStuff  = $hasDeviceCode + $ft{'kern_launch'} + $ft{'kernel_func'};
+            my $kernStuff  = $hasDeviceCode + $ft{'kernel_launch'} + $ft{'device_function'};
             my $totalCalls = $apiCalls + $kernStuff;
             $is_dos = m/\r\n$/;
             if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) {
                 # TODO: implement hipify-clang's logic with header files AMAP
-                # If this file makes kernel builtin calls, and does not include the cuda_runtime.h,
-                # then add an #include to match "magic" includes provided by NVCC.
-                # This logic can miss cases where cuda_runtime.h is included by another include file.
                 print $OUTFILE '#include "hip/hip_runtime.h"' . ($is_dos ? "\r\n" : "\n");
             }
             print $OUTFILE  "$_";
@@ -1778,12 +1756,11 @@ while (@ARGV) {
         $lineCount = $_ =~ tr/\n//;
     }
     my $totalConverted = totalStats(\%ft);
-    #printf "TOTAL-CONV=%d\n", $totalConverted;
     if (($totalConverted+$warnings) and $print_stats) {
-        printStats("info: converted", \@statNames, \%ft, $warnings, $lineCount);
+        printStats("  info: converted", \@statNames, \%ft, $warnings, $lineCount);
         print STDERR " in '$fileName'\n";
     }
-    # Update totals for all files:
+    # Update totals for all files
     addStats(\%tt, \%ft);
     $Twarnings += $warnings;
     $TlineCount += $lineCount;
@@ -1791,10 +1768,10 @@ while (@ARGV) {
         $TwarningTags{$key} += $warningTags{$key};
     }
 }
-#-- Print total stats for all files processed:
+# Print total stats for all files processed:
 if ($print_stats and ($fileCount > 1)) {
     print STDERR "\n";
-    printStats("info: TOTAL-converted", \@statNames, \%tt, $Twarnings, $TlineCount);
+    printStats("  info: TOTAL-converted", \@statNames, \%tt, $Twarnings, $TlineCount);
     print STDERR "\n";
     foreach my $key (sort { $TwarningTags{$b} <=> $TwarningTags{$a} } keys %TwarningTags) {
         printf STDERR "  warning: unconverted %s : %d\n", $key, $TwarningTags{$key};
@@ -1807,7 +1784,7 @@ if ($print_stats and ($fileCount > 1)) {
     print STDERR "\n";
     print STDERR "\n";
 }
-if ($count_conversions) {
+if ($print_stats) {
     foreach my $key (sort { $convertedTags{$b} <=> $convertedTags{$a} } keys %convertedTags) {
         printf STDERR "  %s %d\n", $key, $convertedTags{$key};
     }

docs/markdown/hip_porting_driver_api.md

@@ -105,11 +105,13 @@ hip-clang emits a global variable `__hip_gpubin_handle` of void** type with link
 #### Kernel Launching
 hip-clang supports kernel launching by CUDA `<<<>>>` syntax, hipLaunchKernel, and hipLaunchKernelGGL. The latter two are macros which expand to CUDA `<<<>>>` syntax.
 
-In host code, hip-clang emits a stub function with the same name and arguments as the kernel. In the body of this function, hipSetupArgument is called for each kernel argument, then hipLaunchByPtr is called with a function pointer to the stub function.
-
 When the executable or shared library is loaded by the dynamic linker, the initilization functions are called. In the initialization functions, when `__hipRegisterFatBinary` is called, the code objects containing all kernels are loaded; when `__hipRegisterFunction` is called, the stub functions are associated with the corresponding kernels in code objects.
 
-In the host code, for the `<<<>>>` statement, hip-clang first emits call of hipConfigureCall to set up the threads and grids, then emits call of the stub function with the given arguments. In the stub function, when the runtime host API function hipLaunchByPtr is called, the real kernel associated with the stub function is launched.
+hip-clang implements two sets of kernel launching APIs.
+
+By default, in the host code, for the `<<<>>>` statement, hip-clang first emits call of hipConfigureCall to set up the threads and grids, then emits call of the stub function with the given arguments. In the stub function, hipSetupArgument is called for each kernel argument, then hipLaunchByPtr is called with a function pointer to the stub function. In hipLaunchByPtr, the real kernel associated with the stub function is launched.
+
+If HIP program is compiled with -fhip-new-launch-api, in the host code, for the `<<<>>>` statement, hip-clang first emits call of `__hipPushCallConfiguration` to save the grid dimension, block dimension, shared memory usage and stream to a stack, then emits call of the stub function with the given arguments. In the stub function, `__hipPopCallConfiguration` is called to get the saved grid dimension, block dimension, shared memory usage and stream, then hipLaunchKernel is called with a function pointer to the stub function. In hipLaunchKernel, the real kernel associated with the stub function is launched.
 
 ### NVCC Implementation Notes
 

docs/markdown/hip_profiling.md

@@ -130,7 +130,7 @@ void FunctionFoo(...)
 
 The HIP marker API is only supported on ROCm platform.  The marker macros are defined on CUDA platforms and will compile, but are silently ignored at runtime.
 
-This [HIP sample](samples/2_Cookbook/2_Profiler/) shows the profiler marker API used in a small application.
+This [HIP sample](https://github.com/ROCm-Developer-Tools/HIP/tree/master/samples/2_Cookbook/2_Profiler) shows the profiler marker API used in a small application.
 
 More information on the marker API can be found in the profiler header file and PDF in a ROCm installation:
 - /opt/rocm/profiler/CXLActivityLogger/include/CXLActivityLogger.h
@@ -236,7 +236,7 @@ Here's an example for one API followed by a description for the sections of the
 - The second line shows the completion of the API, including the numeric return value (`ret= 0`) as well as an string representation for the error code (`hipSuccess`).  If the returned error code is non-zero, then the csecond line message is shown in red (unless HIP_TRACE_API_COLOR is "none" - see below).
 
 
-Heres a specific example showing the output of the [square](samples/0_Intro/square) program running on HIP:
+Heres a specific example showing the output of the [square](https://github.com/ROCm-Developer-Tools/HIP/tree/master/samples/0_Intro/square) program running on HIP:
 
 ```
 $ HIP_TRACE_API=1  ./square.hip.out 

hipify-clang/README.md

@@ -144,9 +144,9 @@ To run it:
 
     * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option:
 
-        - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.3.30`
+        - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38`
 
-        - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.3.30`
+        - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38`
 
 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
 
@@ -180,9 +180,9 @@ To run it:
 
 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 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.4.38
 
-Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, 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.4.38
 
 Build system for the above configurations:
 
@@ -197,7 +197,7 @@ cmake
  -DCMAKE_INSTALL_PREFIX=../dist \
  -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 \
+ -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38 \
  -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \
  ..
 ```
@@ -325,9 +325,9 @@ On Windows 10 the following configurations are tested:
 
 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 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.4.38
 
-LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, 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.4.38
 
 Build system for the above configurations:
 
@@ -344,7 +344,7 @@ cmake
  -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 \
+ -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38 \
  -DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \
  -Thost=x64
  ..

hipify-clang/src/CUDA2HIP_Perl.cpp

@@ -36,6 +36,29 @@ using namespace llvm;
 
 namespace perl {
 
+  const std::string sCopyright =
+    "##\n"
+    "# Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.\n"
+    "#\n"
+    "# Permission is hereby granted, free of charge, to any person obtaining a copy\n"
+    "# of this software and associated documentation files (the \"Software\"), to deal\n"
+    "# in the Software without restriction, including without limitation the rights\n"
+    "# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell\n"
+    "# copies of the Software, and to permit persons to whom the Software is\n"
+    "# furnished to do so, subject to the following conditions:\n"
+    "#\n"
+    "# The above copyright notice and this permission notice shall be included in\n"
+    "# all copies or substantial portions of the Software.\n"
+    "#\n"
+    "# THE SOFTWARE IS PROVIDED \"AS IS\", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR\n"
+    "# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,\n"
+    "# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE\n"
+    "# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER\n"
+    "# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,\n"
+    "# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN\n"
+    "# THE SOFTWARE.\n"
+    "##\n";
+
   const std::string tab = "    ";
   const std::string double_tab = tab + tab;
   const std::string triple_tab = double_tab + tab;
@@ -44,46 +67,130 @@ namespace perl {
   const std::string sReturn_m = "return $m;\n";
   const std::string sForeach = "foreach $func (\n";
   const std::string sMy = "my $m = 0;\n";
+  const std::string sCudaDevice = "cudaDevice";
+  const std::string sCudaDeviceId = "cudaDeviceId";
+  const std::string sCudaDevices = "cudaDevices";
+  const std::string sCudaDevice_t = "cudaDevice_t";
+  const std::string sCudaIDs = "cudaIDs";
+  const std::string sCudaGridDim = "cudaGridDim";
+  const std::string sCudaDimGrid = "cudaDimGrid";
+  const std::string sCudaDimBlock = "cudaDimBlock";
+  const std::string sCudaGradInput = "cudaGradInput";
+  const std::string sCudaGradOutput = "cudaGradOutput";
+  const std::string sCudaInput = "cudaInput";
+  const std::string sCudaOutput = "cudaOutput";
+  const std::string sCudaIndices = "cudaIndices";
+  const std::string sCudaGaugeField = "cudaGaugeField";
+  const std::string sCudaMom = "cudaMom";
+  const std::string sCudaGauge = "cudaGauge";
+  const std::string sCudaInGauge = "cudaInGauge";
+  const std::string sCudaColorSpinorField = "cudaColorSpinorField";
+  const std::string sCudaSiteLink = "cudaSiteLink";
+  const std::string sCudaFatLink = "cudaFatLink";
+  const std::string sCudaStaple = "cudaStaple";
+  const std::string sCudaCloverField = "cudaCloverField";
+  const std::string sCudaParam = "cudaParam";
 
-  void generateHostFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
+  const std::set<std::string> Whitelist{
-    *perlStreamPtr.get() << "\n" << sSub << " transformHostFunctions\n" << "{\n" << tab << sMy;
+    {sCudaDevice}, {sCudaDevice_t}, {sCudaIDs}, {sCudaGridDim}, {sCudaDimGrid}, {sCudaDimBlock}, {sCudaDeviceId}, {sCudaDevices},
+    {sCudaGradInput}, {sCudaGradOutput}, {sCudaInput}, {sCudaOutput}, {sCudaIndices}, {sCudaGaugeField}, {sCudaMom}, {sCudaGauge},
+    {sCudaInGauge}, {sCudaColorSpinorField}, {sCudaSiteLink}, {sCudaFatLink}, {sCudaStaple}, {sCudaCloverField}, {sCudaParam}
+  };
+
+  void generateHeader(std::unique_ptr<std::ostream>& streamPtr) {
+    *streamPtr.get() << "#!/usr/bin/perl -w" << std::endl << std::endl;
+    *streamPtr.get() << sCopyright << std::endl;
+    *streamPtr.get() << "#usage hipify-perl [OPTIONS] INPUT_FILE" << std::endl << std::endl;
+    *streamPtr.get() << "use Getopt::Long;" << std::endl;
+    *streamPtr.get() << "my $whitelist = \"\";" << std::endl << std::endl;
+    *streamPtr.get() << "GetOptions(" << std::endl;
+    *streamPtr.get() << tab << "  \"examine\" => \\$examine                  # Combines -no-output and -print-stats options." << std::endl;
+    *streamPtr.get() << tab << ", \"inplace\" => \\$inplace                  # Modify input file inplace, replacing input with hipified output, save backup in .prehip file." << std::endl;
+    *streamPtr.get() << tab << ", \"no-output\" => \\$no_output              # Don't write any translated output to stdout." << std::endl;
+    *streamPtr.get() << tab << ", \"print-stats\" => \\$print_stats          # Print translation statistics." << std::endl;
+    *streamPtr.get() << tab << ", \"quiet-warnings\" => \\$quiet_warnings    # Don't print warnings on unknown CUDA functions." << std::endl;
+    *streamPtr.get() << tab << ", \"whitelist=s\" => \\$whitelist            # TODO: test it beforehand" << std::endl;
+    *streamPtr.get() << ");" << std::endl << std::endl;
+    *streamPtr.get() << "$print_stats = 1 if $examine;" << std::endl;
+    *streamPtr.get() << "$no_output = 1 if $examine;" << std::endl << std::endl;
+    *streamPtr.get() << "# Whitelist of cuda[A-Z] identifiers, which are commonly used in CUDA sources but don't map to any CUDA API:" << std::endl;
+    *streamPtr.get() << "@whitelist = (";
+    unsigned int num = 0;
+    for (const std::string &m : Whitelist) {
+      *streamPtr.get() << std::endl << tab << (num ? ", " : "  ") << "\"" << m << "\"";
+      ++num;
+    }
+    *streamPtr.get() << std::endl << ");" << std::endl << std::endl;
+    *streamPtr.get() << "push(@whitelist, split(',', $whitelist));" << std::endl << std::endl;
+  }
+
+  void generateStatFunctions(std::unique_ptr<std::ostream>& streamPtr) {
+    *streamPtr.get() << std::endl << sSub << " totalStats" << " {" << std::endl;
+    *streamPtr.get() << tab << "my %count = %{ shift() };" << std::endl;
+    *streamPtr.get() << tab << "my $total = 0;" << std::endl;
+    *streamPtr.get() << tab << "foreach $key (keys %count) {" << std::endl;
+    *streamPtr.get() << double_tab << "$total += $count{$key};" << std::endl << tab << "}" << std::endl;
+    *streamPtr.get() << tab << "return $total;" << std::endl << "};" << std::endl;
+    *streamPtr.get() << std::endl << sSub << " printStats" << " {" << std::endl;
+    *streamPtr.get() << tab << "my $label     = shift();" << std::endl;
+    *streamPtr.get() << tab << "my @statNames = @{ shift() };" << std::endl;
+    *streamPtr.get() << tab << "my %counts    = %{ shift() };" << std::endl;
+    *streamPtr.get() << tab << "my $warnings  = shift();" << std::endl;
+    *streamPtr.get() << tab << "my $loc       = shift();" << std::endl;
+    *streamPtr.get() << tab << "my $total     = totalStats(\\%counts);" << std::endl;
+    *streamPtr.get() << tab << "printf STDERR \"%s %d CUDA->HIP refs ( \", $label, $total;" << std::endl;
+    *streamPtr.get() << tab << "foreach $stat (@statNames) {" << std::endl;
+    *streamPtr.get() << double_tab << "printf STDERR \"%s:%d \", $stat, $counts{$stat};" << std::endl;
+    *streamPtr.get() << tab << "}" << std::endl;
+    *streamPtr.get() << tab << "printf STDERR \")\\n  warn:%d LOC:%d\", $warnings, $loc;" << std::endl << "};" << std::endl;
+    for (int i = 0; i < 2; ++i) {
+      *streamPtr.get() << std::endl << sSub << " " << (i ? "clearStats" : "addStats") << " {" << std::endl;
+      *streamPtr.get() << tab << "my $dest_ref  = shift();" << std::endl;
+      *streamPtr.get() << tab << (i ? "my @statNames = @{ shift() };" : "my %adder     = %{ shift() };") << std::endl;
+      *streamPtr.get() << tab << "foreach " << (i ? "$stat(@statNames)" : "$key (keys %adder)") << " {" << std::endl;
+      *streamPtr.get() << double_tab << "$dest_ref->" << (i ? "{$stat} = 0;" : "{$key} += $adder{$key};") << std::endl << tab << "}" << std::endl << "};" << std::endl;
+    }
+  }
+
+  void generateHostFunctions(std::unique_ptr<std::ostream>& streamPtr) {
+    *streamPtr.get() << std::endl << sSub << " transformHostFunctions" << "{" << std::endl << tab << sMy;
     std::set<std::string> &funcSet = DeviceSymbolFunctions0;
     const std::string s0 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
     const std::string s1 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, ";
     for (int i = 0; i < 4; ++i) {
-      *perlStreamPtr.get() << tab + sForeach;
+      *streamPtr.get() << tab + sForeach;
       switch (i) {
-        case 1: funcSet = DeviceSymbolFunctions1; break;
+      case 1: funcSet = DeviceSymbolFunctions1; break;
-        case 2: funcSet = ReinterpretFunctions0; break;
+      case 2: funcSet = ReinterpretFunctions0; break;
-        case 3: funcSet = ReinterpretFunctions1; break;
+      case 3: funcSet = ReinterpretFunctions1; break;
-        default: funcSet = DeviceSymbolFunctions0;
+      default: funcSet = DeviceSymbolFunctions0;
       }
       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_tab << "\"" << found->second.hipName.str() << "\"";
+          *streamPtr.get() << (count ? ",\n" : "") << double_tab << "\"" << found->second.hipName.str() << "\"";
           count++;
         }
       }
-      *perlStreamPtr.get() << "\n" << tab << ")\n" << tab << "{\n" << double_tab;
+      *streamPtr.get() << std::endl << tab << ")" << std::endl << tab << "{" << std::endl << double_tab;
       switch (i) {
-        case 0:
+      case 0:
-        default:
+      default:
-          *perlStreamPtr.get() << s0 << sHIP_SYMBOL << "\\($2\\),/g\n"; break;
+        *streamPtr.get() << s0 << sHIP_SYMBOL << "\\($2\\),/g" << std::endl; break;
-        case 1:
+      case 1:
-          *perlStreamPtr.get() << s1 << sHIP_SYMBOL << "\\($3\\)$4/g;\n"; break;
+        *streamPtr.get() << s1 << sHIP_SYMBOL << "\\($3\\)$4/g;" << std::endl; break;
-        case 2:
+      case 2:
-          *perlStreamPtr.get() << s0 << s_reinterpret_cast << "\\($2\\),/g\n"; break;
+        *streamPtr.get() << s0 << s_reinterpret_cast << "\\($2\\),/g" << std::endl; break;
-        case 3:
+      case 3:
-          *perlStreamPtr.get() << s1 << s_reinterpret_cast << "\\($3\\)$4/g;\n"; break;
+        *streamPtr.get() << s1 << s_reinterpret_cast << "\\($3\\)$4/g;" << std::endl; break;
       }
-      *perlStreamPtr.get() << tab << "}\n";
+      *streamPtr.get() << tab << "}" << std::endl;
     }
-    *perlStreamPtr.get() << tab << sReturn_m << "}\n";
+    *streamPtr.get() << tab << sReturn_m << "}" << std::endl;
   }
 
-  void generateDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
+  void generateDeviceFunctions(std::unique_ptr<std::ostream>& streamPtr) {
     unsigned int countUnsupported = 0;
     unsigned int countSupported = 0;
     std::stringstream sSupported;
@@ -98,39 +205,39 @@ namespace perl {
     std::stringstream subWarnUnsupported;
     std::stringstream subCommon;
     std::string sCommon = tab + sMy + tab + sForeach;
-    subCountSupported << "\n" << sSub << " countSupportedDeviceFunctions\n" << "{\n" << (countSupported ? sCommon : tab + sReturn_0);
+    subCountSupported << std::endl << sSub << " countSupportedDeviceFunctions" << " {" << std::endl << (countSupported ? sCommon : tab + sReturn_0);
-    subWarnUnsupported << "\n" << sSub << " warnUnsupportedDeviceFunctions\n" << "{\n" << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0);
+    subWarnUnsupported << std::endl << sSub << " warnUnsupportedDeviceFunctions" << " {" << std::endl << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0);
     if (countSupported) {
-      subCountSupported << sSupported.str() << "\n" << tab << ")\n";
+      subCountSupported << sSupported.str() << std::endl << tab << ")" << std::endl;
     }
     if (countUnsupported) {
-      subWarnUnsupported << sUnsupported.str() << "\n" << tab << ")\n";
+      subWarnUnsupported << sUnsupported.str() << std::endl << tab << ")" << std::endl;
     }
     if (countSupported || countUnsupported) {
-      subCommon << tab << "{\n";
+      subCommon << tab << "{" << std::endl;
-      subCommon << double_tab << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
+      subCommon << double_tab << "# match device function from the list, except those, which have a namespace prefix (aka somenamespace::umin(...));" << std::endl;
-      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 << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);" << std::endl;
-      subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl;
-      subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
+      subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl;
-      subCommon << double_tab << "if ($mt && !$mt_namespace) {\n";
+      subCommon << double_tab << "if ($mt && !$mt_namespace) {" << std::endl;
-      subCommon << triple_tab << "$m += $mt;\n";
+      subCommon << triple_tab << "$m += $mt;" << std::endl;
     }
     if (countSupported) {
       subCountSupported << subCommon.str();
     }
     if (countUnsupported) {
       subWarnUnsupported << subCommon.str();
-      subWarnUnsupported << triple_tab << "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\";" << std::endl;
     }
     if (countSupported || countUnsupported) {
       sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_m;
     }
     if (countSupported) subCountSupported << sCommon;
     if (countUnsupported) subWarnUnsupported << sCommon;
-    subCountSupported << "}\n";
+    subCountSupported << "}" << std::endl;
-    subWarnUnsupported << "}\n";
+    subWarnUnsupported << "}" << std::endl;
-    *perlStreamPtr.get() << subCountSupported.str();
+    *streamPtr.get() << subCountSupported.str();
-    *perlStreamPtr.get() << subWarnUnsupported.str();
+    *streamPtr.get() << subWarnUnsupported.str();
   }
 
   bool generate(bool Generate) {
@@ -150,15 +257,17 @@ namespace perl {
       llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n";
       return false;
     }
-    std::unique_ptr<std::ostream> perlStreamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
+    std::unique_ptr<std::ostream> streamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
+    generateHeader(streamPtr);
     std::string sConv = "my $conversions = ";
-    *perlStreamPtr.get() << "@statNames = (";
+    *streamPtr.get() << "@statNames = (";
     for (int i = 0; i < NUM_CONV_TYPES - 1; ++i) {
-      *perlStreamPtr.get() << "\"" << counterNames[i] << "\", ";
+      *streamPtr.get() << "\"" << counterNames[i] << "\", ";
       sConv += "$ft{'" + std::string(counterNames[i]) + "'} + ";
     }
-    *perlStreamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");\n\n";
+    *streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << std::endl;
-    *perlStreamPtr.get() << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};\n\n";
+    generateStatFunctions(streamPtr);
+    *streamPtr.get() << std::endl << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};" << std::endl << std::endl;
     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) {
@@ -168,21 +277,22 @@ namespace perl {
             std::string sHIP = ma.second.hipName.str();
             sCUDA = std::regex_replace(sCUDA, std::regex("/"), "\\/");
             sHIP = std::regex_replace(sHIP, std::regex("/"), "\\/");
-            *perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;\n";
+            *streamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;" << std::endl;
           }
         }
-      } else {
+      }
+      else {
         for (auto& ma : CUDA_RENAMES_MAP()) {
           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";
+            *streamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << std::endl;
           }
         }
       }
     }
-    generateHostFunctions(perlStreamPtr);
+    generateHostFunctions(streamPtr);
-    generateDeviceFunctions(perlStreamPtr);
+    generateDeviceFunctions(streamPtr);
-    perlStreamPtr.get()->flush();
+    streamPtr.get()->flush();
     bool ret = true;
     EC = sys::fs::copy_file(tmpFile, dstPerlMap);
     if (EC) {

hipify-clang/src/HipifyAction.cpp

@@ -380,7 +380,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
   ct::Replacement Rep(*SM, launchStart, length, OS.str());
   clang::FullSourceLoc fullSL(launchStart, *SM);
   insertReplacement(Rep, fullSL);
-  hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_EXECUTION, ApiTypes::API_RUNTIME};
+  hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
   Statistics::current().incrementCounter(counter, refName.str());
   return true;
 }
@@ -423,7 +423,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
     ct::Replacement Rep(*SM, slStart, repLength, repName);
     clang::FullSourceLoc fullSL(slStart, *SM);
     insertReplacement(Rep, fullSL);
-    hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME};
+    hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_EXTERN_SHARED, ApiTypes::API_RUNTIME};
     Statistics::current().incrementCounter(counter, refName.str());
     return true;
   }

hipify-clang/src/Statistics.cpp

@@ -63,7 +63,9 @@ const char *counterNames[NUM_CONV_TYPES] = {
   "type", // CONV_TYPE
   "literal", // CONV_LITERAL
   "numeric_literal", // CONV_NUMERIC_LITERAL
-  "define" // CONV_DEFINE
+  "define", // CONV_DEFINE
+  "extern_shared", // CONV_EXTERN_SHARED
+  "kernel_launch" // CONV_KERNEL_LAUNCH
 };
 
 const char *counterTypes[NUM_CONV_TYPES] = {
@@ -102,7 +104,9 @@ const char *counterTypes[NUM_CONV_TYPES] = {
   "CONV_TYPE",
   "CONV_LITERAL",
   "CONV_NUMERIC_LITERAL",
-  "CONV_DEFINE"
+  "CONV_DEFINE",
+  "CONV_EXTERN_SHARED",
+  "CONV_KERNEL_LAUNCH"
 };
 
 const char *apiNames[NUM_API_TYPES] = {

hipify-clang/src/Statistics.h

@@ -119,6 +119,8 @@ enum ConvTypes {
   CONV_LITERAL,
   CONV_NUMERIC_LITERAL,
   CONV_DEFINE,
+  CONV_EXTERN_SHARED,
+  CONV_KERNEL_LAUNCH,
   CONV_LAST
 };
 constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;

include/hip/hcc_detail/device_functions.h

@@ -690,12 +690,7 @@ extern "C" uint64_t __clock_u64()  __HC__;
 __device__
 inline  __attribute((always_inline))
 long long int __clock64() {
-// ToDo: Unify HCC and HIP implementation.
+return (long long int)  __builtin_readcyclecounter();
-#if __HCC__
-    return (long long int) __clock_u64();
-#else
-    return (long long int) __builtin_amdgcn_s_memrealtime();
-#endif
 }
 
 __device__

include/hip/hcc_detail/hip_fp16.h

@@ -223,6 +223,7 @@ THE SOFTWARE.
                     Enable_if_t<
                         std::is_floating_point<T>{} &&
                         !std::is_same<T, double>{}>* = nullptr>
+                __host__ __device__
                 operator T() const { return data; }
             #endif
             __host__ __device__

include/hip/hcc_detail/hip_runtime.h

@@ -37,6 +37,7 @@ THE SOFTWARE.
 //#include <cstring>
 #if __cplusplus
 #include <cmath>
+#include <cstdint>
 #else
 #include <math.h>
 #include <string.h>
@@ -198,35 +199,93 @@ __device__ int __hip_move_dpp_N(int src);
 
 #if defined __HCC__
 
-template <
+namespace hip_impl {
-    typename std::common_type<decltype(hc_get_group_id), decltype(hc_get_group_size),
+  struct GroupId {
-                              decltype(hc_get_num_groups), decltype(hc_get_workitem_id)>::type f>
+    using R = decltype(hc_get_group_id(0));
-class Coordinates {
-    using R = decltype(f(0));
 
-    struct X {
+    __device__
-        __device__ operator R() const { return f(0); }
+    R operator()(std::uint32_t x) const noexcept { return hc_get_group_id(x); }
-        __device__ uint32_t operator=(R _) { return f(0); }
+  };
-    };
+  struct GroupSize {
-    struct Y {
+    using R = decltype(hc_get_group_size(0));
-        __device__ operator R() const { return f(1); }
-        __device__ uint32_t operator=(R _) { return f(1); }
-    };
-    struct Z {
-        __device__ operator R() const { return f(2); }
-        __device__ uint32_t operator=(R _) { return f(2); }
-    };
 
-   public:
+    __device__
-    static constexpr X x{};
+    R operator()(std::uint32_t x) const noexcept {
-    static constexpr Y y{};
+      return hc_get_group_size(x);
-    static constexpr Z z{};
+    }
+  };
+  struct NumGroups {
+    using R = decltype(hc_get_num_groups(0));
+
+    __device__
+    R operator()(std::uint32_t x) const noexcept {
+      return hc_get_num_groups(x);
+    }
+  };
+  struct WorkitemId {
+    using R = decltype(hc_get_workitem_id(0));
+
+    __device__
+    R operator()(std::uint32_t x) const noexcept {
+      return hc_get_workitem_id(x);
+    }
+  };
+} // Namespace hip_impl.
+
+template <typename F>
+struct Coordinates {
+  using R = decltype(F{}(0));
+
+  struct X { __device__ operator R() const noexcept { return F{}(0); } };
+  struct Y { __device__ operator R() const noexcept { return F{}(1); } };
+  struct Z { __device__ operator R() const noexcept { return F{}(2); } };
+
+  static constexpr X x{};
+  static constexpr Y y{};
+  static constexpr Z z{};
 };
 
-static constexpr Coordinates<hc_get_group_size> blockDim;
+inline
-static constexpr Coordinates<hc_get_group_id> blockIdx;
+__device__
-static constexpr Coordinates<hc_get_num_groups> gridDim;
+std::uint32_t operator*(Coordinates<hip_impl::NumGroups>::X,
-static constexpr Coordinates<hc_get_workitem_id> threadIdx;
+                        Coordinates<hip_impl::GroupSize>::X) noexcept {
+  return hc_get_grid_size(0);
+}
+inline
+__device__
+std::uint32_t operator*(Coordinates<hip_impl::GroupSize>::X,
+                        Coordinates<hip_impl::NumGroups>::X) noexcept {
+  return hc_get_grid_size(0);
+}
+inline
+__device__
+std::uint32_t operator*(Coordinates<hip_impl::NumGroups>::Y,
+                        Coordinates<hip_impl::GroupSize>::Y) noexcept {
+  return hc_get_grid_size(1);
+}
+inline
+__device__
+std::uint32_t operator*(Coordinates<hip_impl::GroupSize>::Y,
+                        Coordinates<hip_impl::NumGroups>::Y) noexcept {
+  return hc_get_grid_size(1);
+}
+inline
+__device__
+std::uint32_t operator*(Coordinates<hip_impl::NumGroups>::Z,
+                        Coordinates<hip_impl::GroupSize>::Z) noexcept {
+  return hc_get_grid_size(2);
+}
+inline
+__device__
+std::uint32_t operator*(Coordinates<hip_impl::GroupSize>::Z,
+                        Coordinates<hip_impl::NumGroups>::Z) noexcept {
+  return hc_get_grid_size(2);
+}
+
+static constexpr Coordinates<hip_impl::GroupSize> blockDim{};
+static constexpr Coordinates<hip_impl::GroupId> blockIdx{};
+static constexpr Coordinates<hip_impl::NumGroups> gridDim{};
+static constexpr Coordinates<hip_impl::WorkitemId> threadIdx{};
 
 #define hipThreadIdx_x (hc_get_workitem_id(0))
 #define hipThreadIdx_y (hc_get_workitem_id(1))

include/hip/hcc_detail/hip_runtime_api.h

@@ -1519,34 +1519,6 @@ hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind,
 } // Namespace hip_impl.
 #endif
 
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-/**
- * @brief C compliant kernel launch API
- *
- * @param [in] function_address - kernel function pointer.
- * @param [in] numBlocks - number of blocks
- * @param [in] dimBlocks - dimension of a block
- * @param [in] args - kernel arguments
- * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel.  The
- *  Kernel can access this with HIP_DYNAMIC_SHARED.
- * @param [in] stream - Stream where the kernel should be dispatched.  May be 0, in which case th
- *  default stream is used with associated synchronization rules.
- *
- * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
- *
- */
-
-hipError_t  hipLaunchKernel(const void* function_address,
-			    dim3 numBlocks, dim3 dimBlocks, void** args,
-				size_t sharedMemBytes, hipStream_t stream);
-
-#ifdef __cplusplus
-}
-#endif
-
 #if defined(__cplusplus)
 extern "C" {
 #endif
@@ -3055,6 +3027,65 @@ hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset);
 hipError_t hipLaunchByPtr(const void* func);
 
 
+/**
+ * @brief Push configuration of a kernel launch.
+ *
+ * @param [in] gridDim   grid dimension specified as multiple of blockDim.
+ * @param [in] blockDim  block dimensions specified in work-items
+ * @param [in] sharedMem Amount of dynamic shared memory to allocate for this kernel.  The
+ * kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] stream    Stream where the kernel should be dispatched.  May be 0, in which case the
+ * default stream is used with associated synchronization rules.
+ *
+ * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
+ *
+ */
+
+hipError_t __hipPushCallConfiguration(dim3 gridDim,
+                                      dim3 blockDim,
+                                      size_t sharedMem __dparm(0),
+                                      hipStream_t stream __dparm(0));
+
+/**
+ * @brief Pop configuration of a kernel launch.
+ *
+ * @param [out] gridDim   grid dimension specified as multiple of blockDim.
+ * @param [out] blockDim  block dimensions specified in work-items
+ * @param [out] sharedMem Amount of dynamic shared memory to allocate for this kernel.  The
+ * kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [out] stream    Stream where the kernel should be dispatched.  May be 0, in which case the
+ * default stream is used with associated synchronization rules.
+ *
+ * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
+ *
+ */
+hipError_t __hipPopCallConfiguration(dim3 *gridDim,
+                                     dim3 *blockDim,
+                                     size_t *sharedMem,
+                                     hipStream_t *stream);
+
+/**
+ * @brief C compliant kernel launch API
+ *
+ * @param [in] function_address - kernel stub function pointer.
+ * @param [in] numBlocks - number of blocks
+ * @param [in] dimBlocks - dimension of a block
+ * @param [in] args - kernel arguments
+ * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel.  The
+ *  Kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] stream - Stream where the kernel should be dispatched.  May be 0, in which case th
+ *  default stream is used with associated synchronization rules.
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
+ *
+ */
+
+hipError_t hipLaunchKernel(const void* function_address,
+                           dim3 numBlocks,
+                           dim3 dimBlocks,
+                           void** args,
+                           size_t sharedMemBytes __dparm(0),
+                           hipStream_t stream __dparm(0));
 
 /**
  * @}
@@ -3065,7 +3096,9 @@ hipError_t hipLaunchByPtr(const void* func);
 } /* extern "c" */
 #endif
 
+#if USE_PROF_API
 #include <hip/hcc_detail/hip_prof_str.h>
+#endif
 
 #ifdef __cplusplus
 extern "C" {

src/hip_module.cpp

@@ -132,20 +132,6 @@ extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device);
         return ihipLogStatus(hipStatus);                                                           \
     }
 
-hipError_t hipModuleUnload(hipModule_t hmod) {
-    HIP_INIT_API(hipModuleUnload, hmod);
-
-    // TODO - improve this synchronization so it is thread-safe.
-    // Currently we want for all inflight activity to complete, but don't prevent another
-    // thread from launching new kernels before we finish this operation.
-    ihipSynchronize(tls);
-
-    delete hmod;  // The ihipModule_t dtor will clean everything up.
-    hmod = nullptr;
-
-    return ihipLogStatus(hipSuccess);
-}
-
 hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t globalWorkSizeX,
                                   uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                   uint32_t localWorkSizeX, uint32_t localWorkSizeY,
@@ -821,6 +807,18 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag
     return HSA_STATUS_SUCCESS;
 }
 
+inline hsa_status_t remove_agent_global_variables(hsa_executable_t, hsa_agent_t agent,
+                                                  hsa_executable_symbol_t x, void* unused) {
+    hsa_symbol_kind_t t = {};
+    hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t);
+
+    if (t == HSA_SYMBOL_KIND_VARIABLE) {
+        hc::am_memtracker_remove(hip_impl::address(x));
+    }
+
+    return HSA_STATUS_SUCCESS;
+}
+
 hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname,
                                             hsa_agent_t* agent = nullptr) {
     using namespace hip_impl;
@@ -898,8 +896,28 @@ namespace hip_impl {
 
         return r;
     }
+    void remove_agent_globals_from_tracker(hsa_agent_t agent, hsa_executable_t executable) {
+        hsa_executable_iterate_agent_symbols(executable, agent, remove_agent_global_variables, NULL);
+    }
 } // Namespace hip_impl.
 
+hipError_t hipModuleUnload(hipModule_t hmod) {
+    HIP_INIT_API(hipModuleUnload, hmod);
+
+    // TODO - improve this synchronization so it is thread-safe.
+    // Currently we want for all inflight activity to complete, but don't prevent another
+    // thread from launching new kernels before we finish this operation.
+    ihipSynchronize(tls);
+
+    // deleting ihipModule_t does not remove agent globals from hc_am memtracker
+    hip_impl::remove_agent_globals_from_tracker(hip_impl::this_agent(), hip_impl::executable_for(hmod));
+
+    delete hmod;  // The ihipModule_t dtor will clean everything up.
+    hmod = nullptr;
+
+    return ihipLogStatus(hipSuccess);
+}
+
 hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t hmod, const char* name,
                                  hsa_agent_t *agent = nullptr) {
     using namespace hip_impl;

tests/src/deviceLib/hipTestHalf.cpp

@@ -30,6 +30,11 @@ THE SOFTWARE.
 
 #if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ || __HIP_ARCH_GFX908__
 
+__device__ void test_convert() {
+  __half x;
+  float y = (float)x;
+}
+
 __global__
 void __halfMath(bool* result, __half a) {
   result[0] = __heq(__hadd(a, __half{1}), __half{2});