diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index a71d847af7..5073dc0905 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -200,7 +200,10 @@ endif()
 if(HIP_PLATFORM STREQUAL "hcc")
     find_package(amd_comgr REQUIRED CONFIG
         PATHS
-            /opt/rocm/lib/cmake/amd_comgr
+            /opt/rocm/
+        PATH_SUFFIXES
+            cmake/amd_comgr
+            lib/cmake/amd_comgr
     )
     MESSAGE(STATUS "Code Object Manager found at ${amd_comgr_DIR}.")
 endif()
diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl
index 8bb7bdf31f..3e39a5a11d 100755
--- a/projects/hip/bin/hipify-perl
+++ b/projects/hip/bin/hipify-perl
@@ -1,6 +1,6 @@
 #!/usr/bin/perl  -w
 ##
-# Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+# 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
@@ -25,16 +25,16 @@ use Getopt::Long;
 
 my $warn_whitelist ="";
 GetOptions(
-      "print-stats" => \$print_stats                   # print the command-line, like a header.
-    , "count-conversions" => \$count_conversions       # count conversions.
-    , "quiet-warnings" => \$quiet_warnings           # don't print warnings on unknown CUDA functions.
+      "print-stats" => \$print_stats                     # print the command-line, like a header.
+    , "count-conversions" => \$count_conversions         # count conversions.
+    , "quiet-warnings" => \$quiet_warnings               # don't print warnings on unknown CUDA functions.
     , "warn-whitelist=s"=> \$warn_whitelist
-    , "no-translate-builtins" => \$no_translate_builtins      # don't translate math functions.
-    , "no-translate-textures" => \$no_translate_textures
-    , "no-output" => \$no_output                      # don't write any translated output to stdout.
-    , "inplace" => \$inplace                          # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file.
-														# If .prehip file exists, use that as input to hip.
-    , "n" => \$n                                      # combination of print_stats + no-output.
+    , "no-translate-builtins" => \$no_translate_builtins # don't translate math functions.
+    , "no-translate-textures" => \$no_translate_textures # don't translate texture functions.
+    , "no-output" => \$no_output                         # don't write any translated output to stdout.
+    , "inplace" => \$inplace                             # modify input file inplace, replacing input with hipified output, save backup in ".prehip" file.
+                                                         # If .prehip file exists, use that as input to hip.
+    , "n" => \$n                                         # combination of print_stats + no-output.
 );
 
 $print_stats = 1 if $n;
@@ -47,20 +47,17 @@ $no_output   = 1 if $n;
         ,"cudaStatus"
         ,"cudaDevice"
         ,"cudaDevice_t"
-
         ,"cudaIDs"
         ,"cudaGridDim"
         ,"cudaDimGrid"
         ,"cudaDimBlock"
         ,"cudaDeviceId"
         ,"cudaDevices",
-
         ,"cudaGradOutput",
         ,"cudaInput",
         ,"cudaOutput",
         ,"cudaGradInput",
         ,"cudaIndices",
-
         ,"cudaColorSpinorField"
         ,"cudaGaugeField"
         ,"cudaMom"
@@ -77,126 +74,96 @@ $no_output   = 1 if $n;
     );
 #print "WW=@warn_whitelist\n";
 
-
 # Allow users to add their own functions.
 push (@warn_whitelist, split(',',$warn_whitelist));
 
-
 #---
 #Stats tracking code:
 @statNames = ("dev", "mem", "kern", 'coord_func', "math_func", "special_func", "stream", "event", "err", "def", "tex", "extern_shared", "other");
 
-
 #---
 #Compute total of all individual counts:
 sub totalStats
 {
     my %count = %{ shift() };
-
     my $total = 0;
     foreach $key (keys %count) {
         $total += $count{$key};
     }
-
     return $total;
 };
 
 #---
 sub printStats
 {
-    my $label =     shift();
+    my $label     = shift();
     my @statNames = @{ shift() };
-    my %counts =    %{ shift() };
-    my $warnings =  shift();
+    my %counts    = %{ shift() };
+    my $warnings  = shift();
     my $loc       = shift();
-
-    my $total = totalStats(\%counts);
-
+    my $total     = totalStats(\%counts);
     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;
 }
 
-
 #---
-# Add adder stats to dest.  Used to add stats for current file to a running total for all files:
+# 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() };
-
     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 @statNames = @{ shift() };
-
     foreach $stat (@statNames) {
         $dest_ref->{$stat}  = 0;
     }
 }
 
-
-
-
+#---
 # count of transforms in all files:
 my %tt;
 clearStats(\%tt, \@statNames);
-
-
-
 $Twarnings = 0;
 $TlineCount = 0;
-my %TwarningTags ;
-my %Tkernels ;
-
+my %TwarningTags;
+my %Tkernels;
 my $fileCount = @ARGV;
 my $fileName = "";
 
 while (@ARGV) {
     $fileName=shift (@ARGV);
-
-
     if ($inplace) {
         my $file_prehip = "$fileName" . ".prehip";
-		my $infile;
-		my $outfile;
-		if (-e $file_prehip) {
-			$infile  = $file_prehip;
-			$outfile = $fileName;
-		} else {
-			system ("cp $fileName $file_prehip");
-			$infile = $file_prehip;
-			$outfile = $fileName;
-		}
-
+        my $infile;
+        my $outfile;
+        if (-e $file_prehip) {
+            $infile  = $file_prehip;
+            $outfile = $fileName;
+        } else {
+            system ("cp $fileName $file_prehip");
+            $infile = $file_prehip;
+            $outfile = $fileName;
+        }
         open(INFILE,"<", $infile) or die "error: could not open $infile";
         open(OUTFILE,">", $outfile) or die "error: could not open $outfile";
         $OUTFILE = OUTFILE;
-
     } else {
         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:
-    #
-
     # count of transforms in this file, init to 0 here:
     my %ft;
     clearStats(\%ft, \@statNames);
@@ -206,44 +173,34 @@ while (@ARGV) {
     my $warningsCublas = 0;
     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.
     while (<INFILE>)
     {
-
         #--------
         # Compiler Defines
         # __CUDACC__ is set by NVCC to indicate it is treating the input file as CUDA code (as opposed to host)
         # Typically we want any code treated as CUDA code to be treated as accelerator code by Kalmar too
         # __HIPCC__ will set KALMARCC
         $ft{'def'} += s/\b__CUDACC__\b/__HIPCC__/g;
-
         # __CUDA_ARCH is often used to detect when a function or kernel is being compiled for the device.
         # Don't automaticall convert this - likely these will need special attention with HIP_ARCH_HAS_* macros
         #$ft{'def'} += s/\b__CUDA_ARCH__\b/__HIP_ARCH__/g;
-
-
-
         #--------
         #Includes:
         $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime\.h[>"]/$1<hip\/hip_runtime.h>/;
         $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_runtime_api\.h[>"]/$1<hip\/hip_runtime_api.h>/;
         $countIncludes += s/(\s*#\s*include\s+)[<"]cuda_fp16\.h[>"]/$1<hip\/hip_fp16.h>/;
-
-
         #--------
         # Error codes and return types:
         $ft{'err'} += s/\bcudaError_t\b/hipError_t/g;
         $ft{'err'} += s/\bcudaError\b/hipError_t/g;
         $ft{'err'} += s/\bcudaSuccess\b/hipSuccess/g;
-
         $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g;
         $ft{'err'} += s/\bcudaErrorMemoryAllocation\b/hipErrorMemoryAllocation/g;
         $ft{'err'} += s/\bcudaErrorMemoryFree\b/hipErrorMemoryFree/g;
         $ft{'err'} += s/\bcudaErrorUnknownSymbol\b/hipErrorUnknownSymbol/g;
-	$ft{'err'} += s/\bcudaErrorInvalidSymbol\b/hipErrorInvalidSymbol/g;
+        $ft{'err'} += s/\bcudaErrorInvalidSymbol\b/hipErrorInvalidSymbol/g;
         $ft{'err'} += s/\bcudaErrorOutOfResources\b/hipErrorOutOfResources/g;
         $ft{'err'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g;
         $ft{'err'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g;
@@ -253,15 +210,12 @@ while (@ARGV) {
         $ft{'err'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g;
         $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g;
         $ft{'err'} += s/\bcudaErrorPeerAccessAlreadyEnabled\b/hipErrorPeerAccessAlreadyEnabled/g;
-
-		# error APIs:
+        #--------
+        # error APIs:
         $ft{'err'} += s/\bcudaGetLastError\b/hipGetLastError/g;
         $ft{'err'} += s/\bcudaPeekAtLastError\b/hipPeekAtLastError/g;
         $ft{'err'} += s/\bcudaGetErrorName\b/hipGetErrorName/g;
         $ft{'err'} += s/\bcudaGetErrorString\b/hipGetErrorString/g;
-
-
-
         #--------
         # Memcpy
         $ft{'mem'} += s/\bcudaMemcpy\b/hipMemcpy/g;
@@ -272,48 +226,39 @@ while (@ARGV) {
         $ft{'mem'} += s/\bcudaMemcpyDefault\b/hipMemcpyDefault/g;
         $ft{'mem'} += s/\bcudaMemcpyToSymbol\s*\(\s*(\w+)\b/hipMemcpyToSymbol\(HIP_SYMBOL\($1\)/g;
         $ft{'mem'} += s/\bcudaMemcpyFromSymbol\s*\(\s*(.+?)\s*,\s*(.+?)\b/hipMemcpyFromSymbol\($1, HIP_SYMBOL\($2\)/g;
-
         $ft{'mem'} += s/\bcudaMemset\b/hipMemset/g;
         $ft{'mem'} += s/\bcudaMemsetAsync\b/hipMemsetAsync/g;
-
         $ft{'mem'} += s/\bcudaMemcpyAsync\b/hipMemcpyAsync/g;
-
         $ft{'mem'} += s/\bcudaMemGetInfo\b/hipMemGetInfo/g;
-
         $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g;
-
         $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g;
         $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttributes/g;
-
         $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g;
         $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g;
         $ft{'mem'} += s/\bcudaMemcpyToArray\b/hipMemcpyToArray/g;
-
         $ft{'mem'} += s/\bcudaGetSymbolAddress\s*\(\s*(.+?)\s*,\s*(.+?)\b/hipGetSymbolAddress\($1, HIP_SYMBOL\($2\)/g;
         $ft{'mem'} += s/\bcudaGetSymbolSize\s*\(\s*&(\w+)\s*,\s*(.+?)\b/hipGetSymbolSize(&$1, HIP_SYMBOL\($2\)/g;
-
         #--------
         # Memory management:
         $ft{'mem'} += s/\bcudaMalloc\b/hipMalloc/g;
-        $ft{'mem'} += s/\bcudaMallocHost\b/hipHostMalloc/g;  # note conversion to standard hipHost* naming convention
+        # note conversion to standard hipHost* naming convention
+        $ft{'mem'} += s/\bcudaMallocHost\b/hipHostMalloc/g;
         $ft{'mem'} += s/\bcudaFree\b/hipFree/g;
-        $ft{'mem'} += s/\bcudaFreeHost\b/hipHostFree/g;    # note conversion to standard hipHost* naming convention
-		$ft{'mem'} += s/\bcudaHostAlloc\b/hipHostMalloc/g;
-		$ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g;
-		$ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostMallocDefault/g;
-		$ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g;
-		$ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g;
-		$ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g;
-		$ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g;
+        # note conversion to standard hipHost* naming convention
+        $ft{'mem'} += s/\bcudaFreeHost\b/hipHostFree/g;
+        $ft{'mem'} += s/\bcudaHostAlloc\b/hipHostMalloc/g;
+        $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g;
+        $ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostMallocDefault/g;
+        $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g;
+        $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g;
+        $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g;
+        $ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g;
         $ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g;
         $ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g;
         $ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g;
-
         $ft{'mem'} += s/\bcudaMallocArray\b/hipMallocArray/g;
         $ft{'mem'} += s/\bcudaFreeArray\b/hipFreeArray/g;
         $ft{'mem'} += s/\bcudaMallocPitch\b/hipMallocPitch/g;
-
-
         #--------
         # Events
         $ft{'event'} += s/\bcudaEvent_t\b/hipEvent_t/g;
@@ -325,7 +270,6 @@ while (@ARGV) {
         $ft{'event'} += s/\bcudaEventSynchronize\b/hipEventSynchronize/g;
         $ft{'event'} += s/\bcudaEventDisableTiming\b/hipEventDisableTiming/g;
         $ft{'event'} += s/\bcudaEventQuery\b/hipEventQuery/g;
-
         #--------
         # Streams
         $ft{'stream'} += s/\bcudaStream_t\b/hipStream_t/g;
@@ -336,23 +280,22 @@ while (@ARGV) {
         $ft{'stream'} += s/\bcudaStreamSynchronize\b/hipStreamSynchronize/g;
         $ft{'stream'} += s/\bcudaStreamDefault\b/hipStreamDefault/g;
         $ft{'stream'} += s/\bcudaStreamNonBlocking\b/hipStreamNonBlocking/g;
-
-
         #--------
         # Other synchronization
         $ft{'dev'} += s/\bcudaDeviceSynchronize\b/hipDeviceSynchronize/g;
-        $ft{'dev'} += s/\bcudaThreadSynchronize\b/hipDeviceSynchronize/g;  # translate deprecated cudaThreadSynchronize
+        # translate deprecated cudaThreadSynchronize
+        $ft{'dev'} += s/\bcudaThreadSynchronize\b/hipDeviceSynchronize/g;
         $ft{'dev'} += s/\bcudaDeviceReset\b/hipDeviceReset/g;
-        $ft{'dev'} += s/\bcudaThreadExit\b/hipDeviceReset/g;               # translate deprecated cudaThreadExit
+        # translate deprecated cudaThreadExit
+        $ft{'dev'} += s/\bcudaThreadExit\b/hipDeviceReset/g;
         $ft{'dev'} += s/\bcudaSetDevice\b/hipSetDevice/g;
         $ft{'dev'} += s/\bcudaGetDevice\b/hipGetDevice/g;
-
         #--------
         # Device
         $ft{'dev'} += s/\bcudaDeviceProp\b/hipDeviceProp_t/g;
         $ft{'dev'} += s/\bcudaGetDeviceProperties\b/hipGetDeviceProperties/g;
         $ft{'dev'} += s/\bcudaDeviceGetPCIBusId\b/hipDeviceGetPCIBusId/g;
-
+        #--------
         # Attribute
         $ft{'err'} += s/\bcudaDevAttrMaxThreadsPerBlock\b/hipDeviceAttributeMaxThreadsPerBlock/g;
         $ft{'err'} += s/\bcudaDevAttrMaxBlockDimX\b/hipDeviceAttributeMaxBlockDimX/g;
@@ -378,9 +321,10 @@ while (@ARGV) {
         $ft{'err'} += s/\bcudaDevAttrMaxSharedMemoryPerMultiprocessor\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g;
         $ft{'err'} += s/\bcudaDevAttrMemoryClockRate\b/hipDeviceAttributeMemoryClockRate/g;
         $ft{'err'} += s/\bcudaDevAttrGlobalMemoryBusWidth\b/hipDeviceAttributeMemoryBusWidth/g;
+        #--------
         $ft{'dev'} += s/\bcudaDeviceAttr\b/hipDeviceAttribute_t/g;
         $ft{'dev'} += s/\bcudaDeviceGetAttribute\b/hipDeviceGetAttribute/g;
-
+        #--------
         # Cache config
         $ft{'dev'} += s/\bcudaDeviceSetCacheConfig\b/hipDeviceSetCacheConfig/g;
         $ft{'dev'} += s/\bcudaThreadSetCacheConfig\b/hipDeviceSetCacheConfig/g; # translate deprecated
@@ -393,10 +337,7 @@ while (@ARGV) {
         $ft{'dev'} += s/\bcudaFuncCachePreferEqual\b/hipFuncCachePreferEqual/g;
         # function
         $ft{'dev'} += s/\bcudaFuncSetCacheConfig\b/hipFuncSetCacheConfig/g;
-
-
         $ft{'dev'} += s/\bcudaDriverGetVersion\b/hipDriverGetVersion/g;
-
         #--------
         # Peer2Peer
         $ft{'dev'} += s/\bcudaDeviceCanAccessPeer\b/hipDeviceCanAccessPeer/g;
@@ -409,8 +350,7 @@ while (@ARGV) {
         $ft{'mem'} += s/\bcudaIpcGetMemHandle\b/hipIpcGetMemHandle/g;
         $ft{'mem'} += s/\bcudaIpcMemHandle_t\b/hipIpcMemHandle_t/g;
         $ft{'mem'} += s/\bcudaIpcMemLazyEnablePeerAccess\b/hipIpcMemLazyEnablePeerAccess/g;
-
-
+        #--------
         # Shared mem:
         $ft{'dev'} += s/\bcudaDeviceSetSharedMemConfig\b/hipDeviceSetSharedMemConfig/g;
         $ft{'dev'} += s/\bcudaThreadSetSharedMemConfig\b/hipDeviceSetSharedMemConfig/g; # translate deprecated
@@ -420,33 +360,25 @@ while (@ARGV) {
         $ft{'dev'} += s/\bcudaSharedMemBankSizeDefault\b/hipSharedMemBankSizeDefault/g;
         $ft{'dev'} += s/\bcudaSharedMemBankSizeFourByte\b/hipSharedMemBankSizeFourByte/g;
         $ft{'dev'} += s/\bcudaSharedMemBankSizeEightByte\b/hipSharedMemBankSizeEightByte/g;
-
         $ft{'dev'} += s/\bcudaGetDeviceCount\b/hipGetDeviceCount/g;
-
         #--------
         # Profiler
-        #$aOt += s/\bcudaProfilerInitialize\b/hipProfilerInitialize/g;  // see if these are called anywhere.
+        #$aOt += s/\bcudaProfilerInitialize\b/hipProfilerInitialize/g;
         $ft{'other'} += s/\bcudaProfilerStart\b/hipProfilerStart/g;
         $ft{'other'} += s/\bcudaProfilerStop\b/hipProfilerStop/g;
-
-
-
+        #--------
         $countKeywords += m/__global__/;
         $countKeywords += m/__shared__/;
-
         #--------
         # 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.
             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;
-
             # test patterns for the regular expression above:
             #'extern __shared__ double foo[];'
             #'extern __shared__ unsigned int foo[];'
@@ -458,47 +390,35 @@ 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[];'
-
             $ft{'extern_shared'} += $k;
         }
-
         #--------
         # CUDA Launch Syntax
         # Note these only work if launch is on a single line.
-
         # Handle the <<numBlocks, blockDim>>> syntax:
         {
             # match uses ? for <.*> which will be unitialized if this is not present in launch syntax.
             no warnings qw/uninitialized/;
-
             my $k = 0;
             my $kernelName;
-
             # Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax:
             $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, $6, /g;
             $kernelName = $1 if $k;
-
             # Handle the <<numBlocks, blockDim, sharedSize>>> syntax:
             $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, 0, /g;
             $kernelName = $1 if $k;
-
             # Handle the <<numBlocks, blockDim>>> syntax:
             $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), 0, 0, /g;
             $kernelName = $1 if $k;
-
             $ft{'kern'} += $k;
             if ($k) {
                 $Tkernels{$kernelName} ++;
             }
-         }
-
-
-
+        }
         unless ($no_translate_textures) {
             $ft{'tex'} += s/\bcudaChannelFormatDesc\b/hipChannelFormatDesc/g;
             $ft{'tex'} += s/\bcudaFilterModePoint\b/hipFilterModePoint/g;
             $ft{'tex'} += s/\bcudaReadModeElementType\b/hipReadModeElementType/g;
-
             $ft{'tex'} += s/\bcudaArray\b/hipArray/g;
             $ft{'tex'} += s/\bcudaCreateChannelDesc\b/hipCreateChannelDesc/g;
             $ft{'tex'} += s/\bcudaBindTexture\b/hipBindTexture/g;
@@ -508,23 +428,15 @@ while (@ARGV) {
             $ft{'tex'} += s/\bcudaAddressMode/hipAddressMode/g;
             $ft{'tex'} += s/\bcudaFilterMode/hipFilterMode/g;
         }
-
-
         if ($count_conversions) {
             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.
         # will almost certainly call one of the coordiante functions - could be fooled by clever macros but usually works:
         my $hasDeviceCode = $countKeywords + $ft{'coord_func'} + $ft{'math_func'} + $ft{'special_func'};
-
-
-
         unless ($quiet_warnings) {
             #print STDERR "Check WARNINGs\n";
             # copy into array of lines, process line-by-line to show warnings:
@@ -532,58 +444,41 @@ while (@ARGV) {
                 my @lines = split /\n/, $_;
                 my $tmp = $_;  # copies the whole file, could be a little smarter here...
                 my $line_num = 0;
-
                 foreach (@lines) {
-
                     $line_num ++;
-
                     # remove any whitelisted words:
                     foreach $w (@warn_whitelist) {
                         s/\b$w\b/ZAP/
                     }
-
-					my $tag ;
+                    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
                         $tag = (defined $1) ? $1 : "Launch";
-					} elsif (/(\bcublas[A-Z]\w+)/) {
+                    } elsif (/(\bcublas[A-Z]\w+)/) {
                         $warningsCublas++;
-						$tag = $1;
-					} elsif (/(\bcurand[A-Z]\w+)/) {
+                        $tag = $1;
+                    } elsif (/(\bcurand[A-Z]\w+)/) {
                         $warningsCurand++;
-						$tag = $1;
-					}
-
-					if (defined $tag) {
+                        $tag = $1;
+                    }
+                    if (defined $tag) {
                         $warnings++;
                         $warningTags{$tag}++;
                         print STDERR "  warning: $fileName:#$line_num : $_";
                         print STDERR "\n";
                     }
-
                     $s = warnUnsupportedSpecialFunctions($line_num);
                     $warnings += $s;
-
-
                 }
-
                 $_ = $tmp;
             }
         }
-
-
-
-
         #--------
         # Math libraries
         # To limit bogus translations, try to make sure we are in a kernel (ft{'builtin'} != 0):
         if (not $no_translate_builtins and ($hasDeviceCode > 0)) {
             $ft{'special_func'} += countSupportedSpecialFunctions();
         }
-
-
-
-
         #--------
         # Print it!
         # TODO - would like to move this code outside loop but it uses $_ which contains the whole file.
@@ -591,33 +486,23 @@ while (@ARGV) {
             my $apiCalls   = $ft{'err'} + $ft{'event'} + $ft{'mem'} + $ft{'stream'} + $ft{'dev'} + $ft{'def'} +  $ft{'tex'} +  $ft{'other'} + $ft{'math_func'};
             my $kernStuff  = $hasDeviceCode + $ft{'kern'};
             my $totalCalls = $apiCalls + $kernStuff;
-
             $is_dos = m/\r\n$/;
-
             if ($totalCalls and ($countIncludes == 0) and ($kernStuff != 0)) {
                 # 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.
+                # 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  "$_";
         }
-
         $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);
         print STDERR " in '$fileName'\n";
     }
-
-
     # Update totals for all files:
     addStats(\%tt, \%ft);
     $Twarnings += $warnings;
@@ -626,41 +511,31 @@ while (@ARGV) {
         $TwarningTags{$key} += $warningTags{$key};
     }
 }
-
 #-- Print total stats for all files processed:
 if ($print_stats and ($fileCount > 1)) {
     print STDERR "\n";
     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};
     }
-
     my $kernelCnt = keys %Tkernels;
     printf STDERR "  kernels (%d total) : ", $kernelCnt;
     foreach my $key (sort { $Tkernels{$b} <=> $Tkernels{$a} } keys %Tkernels) {
         printf STDERR "  %s(%d)", $key, $Tkernels{$key};
     }
     print STDERR "\n";
-
-
     print STDERR "\n";
 }
-
 if ($count_conversions) {
     foreach my $key (sort { $convertedTags{$b} <=> $convertedTags{$a} } keys %convertedTags) {
         printf STDERR "  %s %d\n", $key, $convertedTags{$key};
     }
 }
 
-
-
-
 sub countSupportedSpecialFunctions
 {
     my $m = 0;
-
     #supported special functions:
     foreach $func (
         # Synchronization:
@@ -669,53 +544,35 @@ sub countSupportedSpecialFunctions
     {
         # match math at the beginning of a word, but not if it already has a namespace qualifier ('::') :
         $m += m/[:]?[:]?\b($func)\b(\w*\()/g;
-
     }
-
     return $m;
 }
 
 sub warnUnsupportedSpecialFunctions
 {
     my $line_num =  shift;
-
     my $m = 0;
-
-
-
     foreach $func (
-        # memory fence:
-        "__threadfence_block",
-        "__threadfence",
-        "__threadfence_system",
-
         # Synchronization:
         "__syncthreads_count",
         "__syncthreads_and",
         "__syncthreads_or",
-
         # Read-only cache function:
         "__ldg",
-
         # Cross-lane and warp-vote instructions:
         #"__all",
         #"__any",
         #"__ballot",
-
         #"__popc",
         #"__clz",
-
         #"__shfl",
         #"__shfl_up",
         #"__shfl_down",
         #"__shfl_xor",
-
         "__prof_trigger",
-
         # too popular, and we can't tell if we are in device or host code.
         #"assert",
         #"printf",
-
         #"malloc",
         #"free",
         #"memset",
@@ -728,8 +585,6 @@ sub warnUnsupportedSpecialFunctions
             $m += $mt;
             print STDERR "  warning: $fileName:#$line_num : unsupported device function : $_\n";
         }
-
     }
-
     return $m;
 }
diff --git a/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md b/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md
index ba1eab75a9..43e7ca419e 100644
--- a/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md
+++ b/projects/hip/docs/markdown/CUDNN_API_supported_by_HIP.md
@@ -292,7 +292,7 @@
 |`cudnnSetTensor`                                           |`hipdnnSetTensor`                                |
 |`cudnnScaleTensor`                                         |`hipdnnScaleTensor`                              |
 |`cudnnCreateFilterDescriptor`                              |`hipdnnCreateFilterDescriptor`                   |
-|`cudnnSetFilter4dDescriptor`                               |                                                 |
+|`cudnnSetFilter4dDescriptor`                               |`hipdnnSetFilter4dDescriptor`                    |
 |`cudnnGetFilter4dDescriptor`                               |                                                 |
 |`cudnnSetFilterNdDescriptor`                               |`hipdnnSetFilterNdDescriptor`                    |
 |`cudnnGetFilterNdDescriptor`                               |`hipdnnGetFilterNdDescriptor`                    |
diff --git a/projects/hip/hip-config-clang.cmake.in b/projects/hip/hip-config-clang.cmake.in
index e4862f3d6c..9680b2ff8e 100644
--- a/projects/hip/hip-config-clang.cmake.in
+++ b/projects/hip/hip-config-clang.cmake.in
@@ -50,6 +50,7 @@ set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" )
 set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc")
 set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig")
 
+find_dependency(amd_comgr)
 include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" )
 
 set( hip_LIBRARIES hip::host hip::device)
diff --git a/projects/hip/hip-config-hcc.cmake.in b/projects/hip/hip-config-hcc.cmake.in
index 7e236ba0b4..c0ffc6e2af 100644
--- a/projects/hip/hip-config-hcc.cmake.in
+++ b/projects/hip/hip-config-hcc.cmake.in
@@ -51,6 +51,7 @@ set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc")
 set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig")
 
 find_dependency(hcc)
+find_dependency(amd_comgr)
 include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" )
 
 set( hip_LIBRARIES hip::host hip::device)
diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md
index e3e26af390..0daaace2f0 100644
--- a/projects/hip/hipify-clang/README.md
+++ b/projects/hip/hipify-clang/README.md
@@ -32,10 +32,12 @@
 ## <a name="dependencies"></a> Dependencies
 
 `hipify-clang` requires:
-1. **LLVM+CLANG** of at least version 3.8.0, the latest stable and recommended release: **6.0.1 on Windows**, and **8.0.0 on Linux**.
+1. **LLVM+CLANG** of at least version 3.8.0; the latest stable and recommended release: **6.0.1 on Windows**, and **8.0.0 on Linux**.
 
 2. **CUDA** at least version 7.0, the latest supported version is **9.0 on Windows**, and **10.0 on Linux**.
 
+If the target CUDA is 9.1, 9.2 or 10.0, to work on Windows you may apply patches*: [for LLVM 7.0.0](patches/patch_for_clang_7.0.0_bug_38811.zip), [for LLVM 7.0.1](patches/patch_for_clang_7.0.1_bug_38811.zip), [for LLVM 8.0.0](patches/patch_for_clang_8.0.0_bug_38811.zip).
+
 | **LLVM release version** | **CUDA latest supported version** | **Windows**  | **Linux** |
 |:------------------------:|:---------------------------------:|:------------:|:---------:|
 | 3.8.0                    | 7.5                               | +            | +         |
@@ -49,12 +51,14 @@
 | 5.0.2                    | 8.0                               | +            | +         |
 | 6.0.0                    | 9.0                               | +            | +         |
 | **6.0.1**                | **9.0**                           | + <br/> **LATEST STABLE RELEASE** | + |
-| 7.0.0                    | 9.2                               | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
-| 7.0.1                    | 9.2                               | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
+| 7.0.0                    | 9.2                               | - <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                    | 9.2                               | - <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                    | 9.2 (?)                           | - <br/> LLVM 7.1.0 <br/> is not yet released | - <br/> LLVM 7.1.0 <br/> is not yet released |
-| **8.0.0**                | **10.0**                          | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) | + <br/> **LATEST STABLE RELEASE** |
+| **8.0.0**                | **10.0**                          | - <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> | + <br/> **LATEST STABLE RELEASE** |
 | 8.0.1                    | 10.1 (?)                          | - <br/> LLVM 8.0.1 <br/> is not yet released | - <br/> LLVM 8.0.1 <br/> is not yet released |
 
+`*` 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
@@ -174,9 +178,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-8.0 - cudnn-9.0
+Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.5.0.56
 
-Ubuntu 16-18: LLVM 8.0.0, CUDA 8.0 - 10.0, cudnn-8.0 - cudnn-10.0
+Ubuntu 16-18: LLVM 8.0.0, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.5.0.56
 
 Build system for the above configurations:
 
@@ -299,11 +303,13 @@ Testing Time: 2.51s
 
 ### <a name="windows"></a >Windows
 
-On Windows the following configurations are tested:
+On Windows 10 the following configurations are tested:
 
-LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-8.0
+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-9.0
+LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.5.0.56
+
+LLVM 7.0.0 - 8.0.0 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.5.0.56
 
 Build system for the above configurations:
 
diff --git a/projects/hip/hipify-clang/patches/patch_for_clang_7.0.0_bug_38811.zip b/projects/hip/hipify-clang/patches/patch_for_clang_7.0.0_bug_38811.zip
new file mode 100644
index 0000000000..28205a2645
Binary files /dev/null and b/projects/hip/hipify-clang/patches/patch_for_clang_7.0.0_bug_38811.zip differ
diff --git a/projects/hip/hipify-clang/patches/patch_for_clang_7.0.1_bug_38811.zip b/projects/hip/hipify-clang/patches/patch_for_clang_7.0.1_bug_38811.zip
new file mode 100644
index 0000000000..fca5bf52e3
Binary files /dev/null and b/projects/hip/hipify-clang/patches/patch_for_clang_7.0.1_bug_38811.zip differ
diff --git a/projects/hip/hipify-clang/patches/patch_for_clang_8.0.0_bug_38811.zip b/projects/hip/hipify-clang/patches/patch_for_clang_8.0.0_bug_38811.zip
new file mode 100644
index 0000000000..be1cea6a03
Binary files /dev/null and b/projects/hip/hipify-clang/patches/patch_for_clang_8.0.0_bug_38811.zip differ
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp
index d99a201b8e..85e138a8fd 100644
--- a/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp
+++ b/projects/hip/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp
@@ -75,7 +75,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DNN_FUNCTION_MAP{
 
   // cuDNN Filter functions
   {"cudnnCreateFilterDescriptor",                         {"hipdnnCreateFilterDescriptor",                         "", CONV_LIB_FUNC, API_DNN}},
-  {"cudnnSetFilter4dDescriptor",                          {"hipdnnSetFilter4dDescriptor",                          "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
+  {"cudnnSetFilter4dDescriptor",                          {"hipdnnSetFilter4dDescriptor",                          "", CONV_LIB_FUNC, API_DNN}},
   {"cudnnGetFilter4dDescriptor",                          {"hipdnnGetFilter4dDescriptor",                          "", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
   {"cudnnSetFilterNdDescriptor",                          {"hipdnnSetFilterNdDescriptor",                          "", CONV_LIB_FUNC, API_DNN}},
   {"cudnnGetFilterNdDescriptor",                          {"hipdnnGetFilterNdDescriptor",                          "", CONV_LIB_FUNC, API_DNN}},
diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp
index 95c8ab7837..87dedc630e 100644
--- a/projects/hip/hipify-clang/src/LLVMCompat.cpp
+++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp
@@ -110,4 +110,12 @@ std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
 #endif
 }
 
+bool pragma_once_outside_header() {
+#if LLVM_VERSION_MAJOR < 4
+  return false;
+#else
+  return true;
+#endif
+}
+
 } // namespace llcompat
diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h
index 98ff66fb41..069aeb91a3 100644
--- a/projects/hip/hipify-clang/src/LLVMCompat.h
+++ b/projects/hip/hipify-clang/src/LLVMCompat.h
@@ -81,4 +81,6 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
 std::error_code real_path(const Twine &path, SmallVectorImpl<char> &output,
                           bool expand_tilde = false);
 
+bool pragma_once_outside_header();
+
 } // namespace llcompat
diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp
index 3b832deb8f..ae22f83ce0 100644
--- a/projects/hip/hipify-clang/src/main.cpp
+++ b/projects/hip/hipify-clang/src/main.cpp
@@ -222,6 +222,9 @@ int main(int argc, const char **argv) {
 #if defined(HIPIFY_CLANG_RES)
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
 #endif
+    if (llcompat::pragma_once_outside_header()) {
+      Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-Wno-pragma-once-outside-header", ct::ArgumentInsertPosition::BEGIN));
+    }
     if (!MacroNames.empty()) {
       for (std::string s : MacroNames) {
         Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::END));
diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h
index b88df2e168..044ca1e634 100644
--- a/projects/hip/include/hip/hcc_detail/device_functions.h
+++ b/projects/hip/include/hip/hcc_detail/device_functions.h
@@ -42,7 +42,7 @@ __device__ static inline unsigned int __popc(unsigned int input) {
     return __builtin_popcount(input);
 }
 __device__ static inline unsigned int __popcll(unsigned long long int input) {
-    return __builtin_popcountl(input);
+    return __builtin_popcountll(input);
 }
 
 __device__ static inline int __clz(int input) {
@@ -58,7 +58,7 @@ __device__ static inline unsigned int __ffs(unsigned int input) {
 }
 
 __device__ static inline unsigned int __ffsll(unsigned long long int input) {
-    return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1;
+    return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
 }
 
 __device__ static inline unsigned int __ffs(int input) {
@@ -66,7 +66,7 @@ __device__ static inline unsigned int __ffs(int input) {
 }
 
 __device__ static inline unsigned int __ffsll(long long int input) {
-    return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1;
+    return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
 }
 
 __device__ static inline unsigned int __brev(unsigned int input) {
diff --git a/projects/hip/include/hip/hcc_detail/hip_common.h b/projects/hip/include/hip/hcc_detail/hip_common.h
new file mode 100644
index 0000000000..2e2abace0c
--- /dev/null
+++ b/projects/hip/include/hip/hcc_detail/hip_common.h
@@ -0,0 +1,40 @@
+/*
+Copyright (c) 2019 - 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.
+*/
+
+#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMMON_H
+#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMMON_H
+
+#if defined(__HCC__)
+#define __HCC_OR_HIP_CLANG__ 1
+#define __HCC_ONLY__ 1
+#define __HIP_CLANG_ONLY__ 0
+#elif defined(__clang__) && defined(__HIP__)
+#define __HCC_OR_HIP_CLANG__ 1
+#define __HCC_ONLY__ 0
+#define __HIP_CLANG_ONLY__ 1
+#else
+#define __HCC_OR_HIP_CLANG__ 0
+#define __HCC_ONLY__ 0
+#define __HIP_CLANG_ONLY__ 0
+#endif
+
+#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMMON_H
diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h
index 849b7278d8..93ede207c2 100644
--- a/projects/hip/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h
@@ -21,6 +21,11 @@ THE SOFTWARE.
 */
 
 #pragma once
+#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
+#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
+
+#include <hip/hcc_detail/hip_common.h>
+
 #include "hip/hcc_detail/host_defines.h"
 #include <assert.h>
 #if defined(__cplusplus)
@@ -1643,3 +1648,5 @@ THE SOFTWARE.
 #elif defined(__GNUC__)
     #include "hip_fp16_gcc.h"
 #endif // !defined(__clang__) && defined(__GNUC__)
+
+#endif // HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
diff --git a/projects/hip/include/hip/hcc_detail/hip_ldg.h b/projects/hip/include/hip/hcc_detail/hip_ldg.h
index a5b80b0a1b..ab86955cd6 100644
--- a/projects/hip/include/hip/hcc_detail/hip_ldg.h
+++ b/projects/hip/include/hip/hcc_detail/hip_ldg.h
@@ -24,7 +24,7 @@ THE SOFTWARE.
 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H
 
 #if defined(__HCC_OR_HIP_CLANG__)
-#if __hcc_workweek__ >= 16164 || defined(__HIP_CLANG_ONLY__)
+#if __hcc_workweek__ >= 16164 || __HIP_CLANG_ONLY__
 #include "hip_vector_types.h"
 #include "host_defines.h"
 
@@ -96,7 +96,7 @@ __device__ inline static double __ldg(const double* ptr) { return ptr[0]; }
 
 __device__ inline static double2 __ldg(const double2* ptr) { return ptr[0]; }
 
-#endif  // __hcc_workweek__ || defined(__HIP_CLANG_ONLY__)
+#endif  // __hcc_workweek__ || __HIP_CLANG_ONLY__
 
 #endif  // defined(__HCC_OR_HIP_CLANG__)
 
diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h
index 212502ea5e..e6f03ab797 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h
@@ -29,19 +29,7 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
 
-#if defined(__HCC__)
-#define __HCC_OR_HIP_CLANG__ 1
-#define __HCC_ONLY__ 1
-#define __HIP_CLANG_ONLY__ 0
-#elif defined(__clang__) && defined(__HIP__)
-#define __HCC_OR_HIP_CLANG__ 1
-#define __HCC_ONLY__ 0
-#define __HIP_CLANG_ONLY__ 1
-#else
-#define __HCC_OR_HIP_CLANG__ 0
-#define __HCC_ONLY__ 0
-#define __HIP_CLANG_ONLY__ 0
-#endif
+#include <hip/hcc_detail/hip_common.h>
 
 //---
 // Top part of file can be compiled with any compiler
diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
index ccfe7aefc3..e8196af702 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
@@ -177,6 +177,9 @@ enum hipLimit_t {
     0x80000000  ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific
                 ///< allocation.
 
+#define hipMemAttachGlobal 0x0
+#define hipMemAttachHost 0x1
+
 #define hipDeviceMallocDefault 0x0
 #define hipDeviceMallocFinegrained 0x1  ///< Memory is allocated in fine grained region of device.
 
@@ -1103,6 +1106,17 @@ hipError_t hipMallocHost(void** ptr, size_t size);
  */
 hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
 
+/**
+ *  @brief Allocates memory that will be automatically managed by the Unified Memory system.
+ *
+ *  @param[out] ptr Pointer to the allocated managed memory
+ *  @param[in]  size Requested memory size
+ *  @param[in]  flags must be either hipMemAttachGlobal/hipMemAttachHost
+ *
+ *  @return #hipSuccess, #hipErrorMemoryAllocation
+ */
+hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0));
+
 /**
  *  @brief Allocate device accessible page locked host memory [Deprecated]
  *
@@ -2966,7 +2980,9 @@ hipError_t hipLaunchByPtr(const void* func);
 } /* extern "c" */
 #endif
 
+#ifdef __cplusplus
 #include <hip/hcc_detail/hip_prof_api.h>
+#endif
 
 #ifdef __cplusplus
 extern "C" {
diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h
index 34363689e4..e7ecede8c1 100644
--- a/projects/hip/include/hip/hip_runtime_api.h
+++ b/projects/hip/include/hip/hip_runtime_api.h
@@ -337,6 +337,12 @@ static inline hipError_t hipHostMalloc(T** ptr, size_t size,
                                        unsigned int flags = hipHostMallocDefault) {
     return hipHostMalloc((void**)ptr, size, flags);
 }
+
+template <class T>
+static inline hipError_t hipMallocManaged(T** devPtr, size_t size,
+                                       unsigned int flags = hipMemAttachGlobal) {
+    return hipMallocManaged((void**)devPtr, size, flags);
+}
 #endif
 
 #endif
diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
index 4d84cfbc51..2ebd4f8b7d 100644
--- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
@@ -111,6 +111,9 @@ typedef enum hipChannelFormatKind {
 #define hipHostMallocCoherent 0x0
 #define hipHostMallocNonCoherent 0x0
 
+#define hipMemAttachGlobal cudaMemAttachGlobal
+#define hipMemAttachHost cudaMemAttachHost
+
 #define hipHostRegisterDefault cudaHostRegisterDefault
 #define hipHostRegisterPortable cudaHostRegisterPortable
 #define hipHostRegisterMapped cudaHostRegisterMapped
@@ -420,6 +423,10 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla
     return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
 }
 
+inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int flags) {
+    return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags));
+}
+
 inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
                                         size_t width, size_t height,
                                         unsigned int flags __dparm(hipArrayDefault)) {
@@ -866,7 +873,7 @@ inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBl
     return hipCUDAErrorTohipError(cerror);
 }
 
-inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, void* ptr) {
+inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr) {
     struct cudaPointerAttributes cPA;
     hipError_t err = hipCUDAErrorTohipError(cudaPointerGetAttributes(&cPA, ptr));
     if (err == hipSuccess) {
diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp
index b6a1164a4a..b8632970d4 100644
--- a/projects/hip/src/hip_memory.cpp
+++ b/projects/hip/src/hip_memory.cpp
@@ -301,9 +301,7 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
     return ihipLogStatus(hip_status);
 }
 
-hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
-    HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
-    HIP_SET_DEVICE();
+hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
     hipError_t hip_status = hipSuccess;
 
     if (HIP_SYNC_HOST_ALLOC) {
@@ -368,6 +366,25 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
     if (HIP_SYNC_HOST_ALLOC) {
         hipDeviceSynchronize();
     }
+    return hip_status;
+}
+
+hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
+    HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
+    HIP_SET_DEVICE();
+    hipError_t hip_status = hipSuccess;
+    hip_status = ihipHostMalloc(ptr, sizeBytes, flags);
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) {
+    HIP_INIT_SPECIAL_API(hipMallocManaged, (TRACE_MEM), devPtr, size, flags);
+    HIP_SET_DEVICE();
+    hipError_t hip_status = hipSuccess;
+    if(flags != hipMemAttachGlobal)
+        hip_status = hipErrorInvalidValue;
+    else
+        hip_status = ihipHostMalloc(devPtr, size, hipHostMallocDefault);
     return ihipLogStatus(hip_status);
 }
 
diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp
index 05daca6055..e1a3b1ff65 100644
--- a/projects/hip/src/hip_module.cpp
+++ b/projects/hip/src/hip_module.cpp
@@ -474,7 +474,7 @@ hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const ch
 // Get kernel for the given hsa agent. Internal use only.
 hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod,
                                   const char* name, hsa_agent_t *agent) {
-    HIP_INIT_API(hipModuleGetFunctionEx, hfunc, hmod, name);
+    HIP_INIT_API(hipModuleGetFunctionEx, hfunc, hmod, name, agent);
     return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name, agent));
 }
 
diff --git a/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu
index 23880097ff..6aae90e403 100644
--- a/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu
+++ b/projects/hip/tests/hipify-clang/unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu
@@ -113,7 +113,7 @@ int main() {
   cudnnFilterDescriptor_t filt_desc;
   // CHECK: CUDNN_CALL(hipdnnCreateFilterDescriptor(&filt_desc));
   CUDNN_CALL(cudnnCreateFilterDescriptor(&filt_desc));
-  // CHECK-NOT: CUDNN_CALL(hipdnnSetFilter4dDescriptor(
+  // CHECK: CUDNN_CALL(hipdnnSetFilter4dDescriptor(
   CUDNN_CALL(cudnnSetFilter4dDescriptor(
         // CHECK: filt_desc, HIPDNN_DATA_FLOAT, HIPDNN_TENSOR_NCHW,
         filt_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW,
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp
new file mode 100644
index 0000000000..bf7662f6b7
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp
@@ -0,0 +1,52 @@
+#include <hip/hip_runtime.h>
+#include <math.h>
+#include "test_common.h"
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t -N 256M
+ * HIT_END
+ */
+
+__global__
+void add(int n, float *x, float *y)
+{
+    int index = blockIdx.x * blockDim.x + threadIdx.x;
+    int stride = blockDim.x * gridDim.x;
+    for (int i = index; i < n; i += stride)
+        y[i] = x[i] + y[i];
+}
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true);
+    int numElements = N;
+    bool testResult = true;
+    float *A, *B;
+
+    hipMallocManaged(&A, numElements*sizeof(float));
+    hipMallocManaged(&B, numElements*sizeof(float));
+
+    for (int i = 0; i < numElements; i++) {
+        A[i] = 1.0f;
+        B[i] = 2.0f;
+    }
+
+    int blockSize = 256;
+    int numBlocks = (numElements + blockSize - 1) / blockSize;
+    dim3 dimGrid(numBlocks, 1, 1);
+    dim3 dimBlock(blockSize, 1, 1);
+    hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, numElements, A, B);
+
+    hipDeviceSynchronize();
+
+    float maxError = 0.0f;
+    for (int i = 0; i < numElements; i++)
+        maxError = fmax(maxError, fabs(B[i]-3.0f));
+
+    hipFree(A);
+    hipFree(B);
+    if(maxError == 0.0f)
+        passed();
+    failed("Output Mismatch\n");
+}
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
index 2d1a591e73..13eca9c24d 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
@@ -58,8 +58,8 @@ int main() {
         HIPCHECK(hipDeviceSynchronize());
         HipTest::checkVectorADD(A_h, B_h, C_h, N);
 
-        HIPCHECK(hipStreamCreate(&s));
         HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipStreamCreate(&s));
         HIPCHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, Nbytes, s));
         HIPCHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, Nbytes, s));