From fa07d40cdd273831dde8ed5a8610d1a6e161c174 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 3 Apr 2019 11:54:50 +0530
Subject: [PATCH 1/9] [hit] Rename RUN -> TEST & RUN_NAMED -> TEST_NAMED

Change-Id: I75e24f15129973cee15fc9dac65d678bd2172074


[ROCm/hip commit: 9f2d1453fb495bb3cf18d6055db44e896e78db8e]
---
 projects/hip/tests/README.md                  | 26 +++++-----
 projects/hip/tests/hit/HIT.cmake              | 16 +++---
 projects/hip/tests/hit/parser                 | 52 +++++++++----------
 .../device/hipFuncDeviceSynchronize.cpp       |  2 +-
 .../Functional/device/hipFuncGetDevice.cpp    |  2 +-
 .../Functional/device/hipFuncSetDevice.cpp    |  2 +-
 .../device/hipFuncSetDeviceFlags.cpp          |  2 +-
 .../hip/tests/src/compiler/hipClassKernel.cpp |  2 +-
 .../hip/tests/src/context/hipCtx_simple.cpp   |  2 +-
 .../hip/tests/src/context/hipMemsetD8.cpp     |  8 +--
 .../src/deviceLib/hipAsynchronousStreams.cpp  |  2 +-
 .../tests/src/deviceLib/hipDeviceMalloc.cpp   |  2 +-
 .../tests/src/deviceLib/hipDeviceMemcpy.cpp   |  2 +-
 .../hipDoublePrecisionIntrinsics.cpp          |  2 +-
 .../hipDoublePrecisionMathDevice.cpp          |  2 +-
 .../deviceLib/hipDoublePrecisionMathHost.cpp  |  2 +-
 .../hip/tests/src/deviceLib/hipFloatMath.cpp  |  2 +-
 .../src/deviceLib/hipFloatMathPrecise.cpp     |  2 +-
 .../src/deviceLib/hipIntegerIntrinsics.cpp    |  2 +-
 .../tests/src/deviceLib/hipMathFunctions.cpp  |  2 +-
 .../src/deviceLib/hipSimpleAtomicsTest.cpp    |  2 +-
 .../hipSinglePrecisionIntrinsics.cpp          |  2 +-
 .../hipSinglePrecisionMathDevice.cpp          |  2 +-
 .../deviceLib/hipSinglePrecisionMathHost.cpp  |  2 +-
 .../hip/tests/src/deviceLib/hipStdComplex.cpp |  2 +-
 .../hip/tests/src/deviceLib/hipTestClock.cpp  |  2 +-
 .../hip/tests/src/deviceLib/hipTestDevice.cpp |  2 +-
 .../src/deviceLib/hipTestDeviceDouble.cpp     |  2 +-
 .../src/deviceLib/hipTestDeviceSymbol.cpp     |  2 +-
 .../src/deviceLib/hipTestDotFunctions.cpp     |  2 +-
 .../hip/tests/src/deviceLib/hipTestFMA.cpp    |  2 +-
 .../hip/tests/src/deviceLib/hipTestHalf.cpp   |  2 +-
 .../src/deviceLib/hipTestIncludeMath.cpp      |  2 +-
 .../tests/src/deviceLib/hipTestNativeHalf.cpp |  2 +-
 .../hip/tests/src/deviceLib/hipTestNew.cpp    |  2 +-
 .../tests/src/deviceLib/hipThreadFence.cpp    |  2 +-
 .../tests/src/deviceLib/hipVectorTypes.cpp    |  2 +-
 .../src/deviceLib/hipVectorTypesDevice.cpp    |  2 +-
 .../hip/tests/src/deviceLib/hip_anyall.cpp    |  2 +-
 .../hip/tests/src/deviceLib/hip_ballot.cpp    |  2 +-
 .../tests/src/deviceLib/hip_bitextract.cpp    |  2 +-
 .../hip/tests/src/deviceLib/hip_bitinsert.cpp |  2 +-
 projects/hip/tests/src/deviceLib/hip_brev.cpp |  2 +-
 projects/hip/tests/src/deviceLib/hip_clz.cpp  |  2 +-
 projects/hip/tests/src/deviceLib/hip_ffs.cpp  |  2 +-
 .../hip/tests/src/deviceLib/hip_mbcnt.cpp     |  2 +-
 projects/hip/tests/src/deviceLib/hip_popc.cpp |  2 +-
 .../hip/tests/src/deviceLib/hip_test_ldg.cpp  |  2 +-
 .../src/deviceLib/hip_threadfence_system.cpp  |  2 +-
 projects/hip/tests/src/deviceLib/hip_trig.cpp |  2 +-
 projects/hip/tests/src/hipC.c                 |  2 +-
 projects/hip/tests/src/hipEnvVarDriver.cpp    |  2 +-
 projects/hip/tests/src/hipHcc.cpp             |  2 +-
 .../hip/tests/src/kernel/hipDynamicShared.cpp |  2 +-
 .../tests/src/kernel/hipDynamicShared2.cpp    |  2 +-
 .../hip/tests/src/kernel/hipEmptyKernel.cpp   |  2 +-
 .../hip/tests/src/kernel/hipGridLaunch.cpp    |  2 +-
 .../src/kernel/hipLanguageExtensions.cpp      |  2 +-
 .../hip/tests/src/kernel/hipLaunchParm.cpp    |  2 +-
 .../tests/src/kernel/hipLaunchParmFunctor.cpp |  2 +-
 .../hip/tests/src/kernel/hipPrintfKernel.cpp  |  2 +-
 .../hip/tests/src/kernel/hipTestConstant.cpp  |  2 +-
 .../src/kernel/hipTestGlobalVariable.cpp      |  2 +-
 .../tests/src/kernel/hipTestMallocKernel.cpp  |  2 +-
 .../hip/tests/src/kernel/hipTestMemKernel.cpp |  2 +-
 .../hip/tests/src/kernel/inline_asm_vadd.cpp  |  2 +-
 .../tests/src/p2p/hipPeerToPeer_simple.cpp    |  6 +--
 .../src/runtimeApi/device/hipChooseDevice.cpp |  2 +-
 .../device/hipDeviceComputeCapability.cpp     |  2 +-
 .../device/hipDeviceGetByPCIBusId.cpp         |  2 +-
 .../runtimeApi/device/hipDeviceGetName.cpp    |  2 +-
 .../device/hipDeviceGetPCIBusId.cpp           |  2 +-
 .../device/hipDeviceSynchronize.cpp           |  2 +-
 .../runtimeApi/device/hipDeviceTotalMem.cpp   |  2 +-
 .../src/runtimeApi/device/hipGetDevice.cpp    |  2 +-
 .../device/hipGetDeviceAttribute.cpp          |  2 +-
 .../device/hipRuntimeGetVersion.cpp           |  2 +-
 .../runtimeApi/device/hipSetCachceConfig.cpp  |  2 +-
 .../src/runtimeApi/device/hipSetDevice.cpp    |  2 +-
 .../runtimeApi/device/hipSetDeviceFlags.cpp   |  2 +-
 .../runtimeApi/error/hipPeekAtLastError.cpp   |  2 +-
 .../src/runtimeApi/event/hipEventRecord.cpp   |  2 +-
 .../src/runtimeApi/event/record_event.cpp     |  2 +-
 .../tests/src/runtimeApi/memory/hipArray.cpp  |  2 +-
 .../src/runtimeApi/memory/hipHostGetFlags.cpp |  2 +-
 .../src/runtimeApi/memory/hipHostMalloc.cpp   |  2 +-
 .../src/runtimeApi/memory/hipHostRegister.cpp |  2 +-
 .../runtimeApi/memory/hipMallocManaged.cpp    |  2 +-
 .../runtimeApi/memory/hipMemPtrGetInfo.cpp    |  2 +-
 .../tests/src/runtimeApi/memory/hipMemcpy.cpp | 10 ++--
 .../src/runtimeApi/memory/hipMemcpyDtoD.cpp   |  2 +-
 .../runtimeApi/memory/hipMemcpyDtoDAsync.cpp  |  2 +-
 .../src/runtimeApi/memory/hipMemcpyPeer.cpp   |  2 +-
 .../runtimeApi/memory/hipMemcpyPeerAsync.cpp  |  2 +-
 .../runtimeApi/memory/hipMemcpy_simple.cpp    |  4 +-
 .../hipMemoryAllocateCoherentDriver.cpp       |  2 +-
 .../tests/src/runtimeApi/memory/hipMemset.cpp |  8 +--
 .../src/runtimeApi/memory/hipMemset2D.cpp     |  2 +-
 .../src/runtimeApi/memory/hipMemset3D.cpp     |  2 +-
 .../memory/hipRandomMemcpyAsync.cpp           |  2 +-
 .../runtimeApi/memory/hipTestMemcpyPin.cpp    |  2 +-
 .../runtimeApi/memory/p2p_copy_coherency.cpp  |  2 +-
 .../module/hipFuncGetAttributes.cpp           |  2 +-
 .../module/hipFuncSetCacheConfig.cpp          |  2 +-
 .../multiThread/hipMultiThreadDevice.cpp      |  6 +--
 .../multiThread/hipMultiThreadStreams1.cpp    |  2 +-
 .../multiThread/hipMultiThreadStreams2.cpp    |  2 +-
 .../src/runtimeApi/stream/hipNullStream.cpp   |  2 +-
 .../stream/hipStreamAddCallback.cpp           |  2 +-
 .../stream/hipStreamCreateWithPriority.cpp    |  2 +-
 .../runtimeApi/stream/hipStreamGetFlags.cpp   |  2 +-
 .../src/runtimeApi/stream/hipStreamL5.cpp     |  2 +-
 .../src/runtimeApi/stream/hipStreamSync2.cpp  |  2 +-
 .../runtimeApi/stream/hipStreamWaitEvent.cpp  |  2 +-
 .../synchronization/copy_coherency.cpp        |  2 +-
 .../hip/tests/src/surface/hipSurfaceObj2D.cpp |  2 +-
 .../src/texture/hipBindTexRef1DFetch.cpp      |  2 +-
 .../hip/tests/src/texture/hipGetChanDesc.cpp  |  2 +-
 .../src/texture/hipTextureObj1DFetch.cpp      |  2 +-
 .../hip/tests/src/texture/hipTextureObj2D.cpp |  2 +-
 .../hip/tests/src/texture/hipTextureRef2D.cpp |  2 +-
 121 files changed, 180 insertions(+), 180 deletions(-)

diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md
index a9638ba95f..77929340c0 100644
--- a/projects/hip/tests/README.md
+++ b/projects/hip/tests/README.md
@@ -31,17 +31,17 @@ The parser looks for a code block similar to the one below.
 ```
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * //Small copy
- * RUN: %t -N 10    --memsetval 0x42
+ * TEST: %t -N 10    --memsetval 0x42
  * // Oddball size
- * RUN: %t -N 10013 --memsetval 0x5a
+ * TEST: %t -N 10013 --memsetval 0x5a
  * // Big copy
- * RUN: %t -N 256M  --memsetval 0xa6
+ * TEST: %t -N 256M  --memsetval 0xa6
  * HIT_END
  */
 ```
-In the above, BUILD commands provide instructions on how to build the test case while RUN commands provide instructions on how to execute the test case.
+In the above, BUILD commands provide instructions on how to build the test case while TEST commands provide instructions on how to execute the test case.
 
 #### BUILD command
 
@@ -57,21 +57,21 @@ NVCC_OPTIONS: All options specified after this delimiter are passed to hipcc on
 EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms.
 
 
-#### RUN command
+#### TEST command
 
-The supported syntax for the RUN command is:
+The supported syntax for the TEST command is:
 ```
-RUN: %t <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
+TEST: %t <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
 ```
 %t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
-EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. Note that if the test has been excluded for a specific platform in the BUILD command, it is automatically excluded from the RUN command as well for the same platform.
+EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. Note that if the test has been excluded for a specific platform in the BUILD command, it is automatically excluded from the TEST command as well for the same platform.
 
 
-#### RUN_NAMED command
+#### TEST_NAMED command
 
-When using the RUN command, HIT will squash and append the arguments specified to the test executable name to generate the CMAKE test name. Sometimes we might want to specify a more descriptive name. The RUN_NAMED command is used for that. The supported syntax for the RUN_NAMED command is:
+When using the TEST command, HIT will squash and append the arguments specified to the test executable name to generate the CMAKE test name. Sometimes we might want to specify a more descriptive name. The TEST_NAMED command is used for that. The supported syntax for the TEST_NAMED command is:
 ```
-RUN: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
+TEST: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
 ```
 
 
@@ -111,7 +111,7 @@ Find the test and commandline that fail:
 
 (From the build directory, perhaps hip/build)
 grep -IR hipMemcpy-modes -IR ../tests/
-../tests/src/runtimeApi/memory/hipMemcpy.cpp: * RUN_NAMED: %t hipMemcpy-modes --tests 0x1
+../tests/src/runtimeApi/memory/hipMemcpy.cpp: * TEST_NAMED: %t hipMemcpy-modes --tests 0x1
 
 # Guidelines for adding new tests
 
diff --git a/projects/hip/tests/hit/HIT.cmake b/projects/hip/tests/hit/HIT.cmake
index 82e8508dcd..3233a7832a 100644
--- a/projects/hip/tests/hit/HIT.cmake
+++ b/projects/hip/tests/hit/HIT.cmake
@@ -55,8 +55,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
     endforeach()
 endmacro()
 
-# Helper macro to parse RUN instructions
-macro(PARSE_RUN_COMMAND _target _arguments _exclude_platforms)
+# Helper macro to parse TEST instructions
+macro(PARSE_TEST_COMMAND _target _arguments _exclude_platforms)
     set(${_target})
     set(${_arguments} " ")
     set(${_exclude_platforms})
@@ -78,8 +78,8 @@ macro(PARSE_RUN_COMMAND _target _arguments _exclude_platforms)
     endforeach()
 endmacro()
 
-# Helper macro to parse RUN_NAMED instructions
-macro(PARSE_RUN_NAMED_COMMAND _target _testname _arguments _exclude_platforms)
+# Helper macro to parse TEST_NAMED instructions
+macro(PARSE_TEST_NAMED_COMMAND _target _testname _arguments _exclude_platforms)
     set(${_target})
     set(${_arguments} " ")
     set(${_exclude_platforms})
@@ -170,7 +170,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
         endforeach()
 
         # Add tests
-        execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --runCMDs ${file}
+        execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --testCMDs ${file}
             OUTPUT_VARIABLE _contents
             ERROR_QUIET
             WORKING_DIRECTORY ${_dir}
@@ -178,7 +178,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
         string(REGEX REPLACE "\n" ";" _contents "${_contents}")
         foreach(_cmd ${_contents})
             string(REGEX REPLACE " " ";" _cmd "${_cmd}")
-            parse_run_command(_target _arguments _exclude_platforms ${_cmd})
+            parse_test_command(_target _arguments _exclude_platforms ${_cmd})
             string(REGEX REPLACE "/" "." target ${_label}/${_target})
             read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
             if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM} OR
@@ -189,7 +189,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
         endforeach()
 
         # Add named tests
-        execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --runNamedCMDs ${file}
+        execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --testNamedCMDs ${file}
             OUTPUT_VARIABLE _contents
             ERROR_QUIET
             WORKING_DIRECTORY ${_dir}
@@ -197,7 +197,7 @@ macro(HIT_ADD_FILES _dir _label _parent)
         string(REGEX REPLACE "\n" ";" _contents "${_contents}")
         foreach(_cmd ${_contents})
             string(REGEX REPLACE " " ";" _cmd "${_cmd}")
-            parse_run_named_command(_target _testname _arguments _exclude_platforms ${_cmd})
+            parse_test_named_command(_target _testname _arguments _exclude_platforms ${_cmd})
             string(REGEX REPLACE "/" "." target ${_label}/${_target})
             read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
             if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM} OR
diff --git a/projects/hip/tests/hit/parser b/projects/hip/tests/hit/parser
index f77bd524f3..cb2330026e 100755
--- a/projects/hip/tests/hit/parser
+++ b/projects/hip/tests/hit/parser
@@ -8,7 +8,7 @@ use File::Spec;
 sub parse_file {
     my $file = shift;
     (my $exe = $file) =~ s/\.[^.]+$//g;
-    my (@buildCMDs, @runCMDs, @runNamedCMDs, @cmakeCMDs);
+    my (@buildCMDs, @testCMDs, @testNamedCMDs, @cmakeCMDs);
     if (open (SOURCE, '<:encoding(UTF-8)', "$file")) {
         while (<SOURCE>) {
             my $line=$_;
@@ -20,21 +20,21 @@ sub parse_file {
                 $line =~ s/\R//g;               # Remove line endings
                 push @buildCMDs, $line;
             }
-            # Look for RUN instructions
-            if ($line =~ /^ \* RUN:/) {
-                $line =~ s/^ \* RUN: //g;       # Remove " * RUN: "
+            # Look for TEST instructions
+            if ($line =~ /^ \* TEST:/) {
+                $line =~ s/^ \* TEST: //g;      # Remove " * TEST: "
                 $line =~ s/%s/$file/g;          # Substitute %s -> filename
                 $line =~ s/%t/$exe/g;           # Subsitute %t -> targetname
                 $line =~ s/\R//g;               # Remove line endings
-                push @runCMDs, $line;
+                push @testCMDs, $line;
             }
-            # Look for RUN_NAMED instructions
-            if ($line =~ /^ \* RUN_NAMED:/) {
-                $line =~ s/^ \* RUN_NAMED: //g; # Remove " * RUN_NAMED: "
+            # Look for TEST_NAMED instructions
+            if ($line =~ /^ \* TEST_NAMED:/) {
+                $line =~ s/^ \* TEST_NAMED: //g;# Remove " * TEST_NAMED: "
                 $line =~ s/%s/$file/g;          # Substitute %s -> filename
                 $line =~ s/%t/$exe/g;           # Subsitute %t -> targetname
                 $line =~ s/\R//g;               # Remove line endings
-                push @runNamedCMDs, $line;
+                push @testNamedCMDs, $line;
             }
             # Look for CMAKECMD instructions
             if ($line =~ /^ \* CMAKECMD:/) {
@@ -46,7 +46,7 @@ sub parse_file {
         }
         close(SOURCE);
     }
-    return (\@buildCMDs, \@runCMDs, \@runNamedCMDs, \@cmakeCMDs);
+    return (\@buildCMDs, \@testCMDs, \@testNamedCMDs, \@cmakeCMDs);
 }
 
 # Exit if no arguments specified
@@ -58,21 +58,21 @@ if(scalar @ARGV == 0){
 # Parse command
 my @options = ();
 my $retBuildCMDs = 0;
-my $retRunCMDs = 0;
-my $retRunNamedCMDs = 0;
+my $retTestCMDs = 0;
+my $retTestNamedCMDs = 0;
 my $retCmakeCMDs = 0;
 foreach $arg (@ARGV) {
-    if ($retBuildCMDs or $retRunCMDs or $retRunNamedCMDs or $retCmakeCMDs) {
+    if ($retBuildCMDs or $retTestCMDs or $retTestNamedCMDs or $retCmakeCMDs) {
         push (@options, $arg);
     }
     if ($arg eq '--buildCMDs') {
         $retBuildCMDs = 1;
     }
-    if ($arg eq '--runCMDs') {
-        $retRunCMDs = 1;
+    if ($arg eq '--testCMDs') {
+        $retTestCMDs = 1;
     }
-    if ($arg eq '--runNamedCMDs') {
-        $retRunNamedCMDs = 1;
+    if ($arg eq '--testNamedCMDs') {
+        $retTestNamedCMDs = 1;
     }
     if ($arg eq '--cmakeCMDs') {
         $retCmakeCMDs = 1;
@@ -80,26 +80,26 @@ foreach $arg (@ARGV) {
 }
 
 # Atleast one command needs to be specified
-if (($retBuildCMDs eq 0) and ($retRunCMDs eq 0) and ($retRunNamedCMDs eq 0) and ($retCmakeCMDs eq 0)) {
-    die "Usage: $0 <--buildCMDs|--runCMDs|--runNamedCMDs|--cmakeCMDs> FILENAMEs\n";
+if (($retBuildCMDs eq 0) and ($retTestCMDs eq 0) and ($retTestNamedCMDs eq 0) and ($retCmakeCMDs eq 0)) {
+    die "Usage: $0 <--buildCMDs|--testCMDs|--testNamedCMDs|--cmakeCMDs> FILENAMEs\n";
 }
 
 # Iterate over input files
 foreach $file (@options) {
     # Convert absolute path to path relative to working directory
     my $relfile = File::Spec->abs2rel($file);
-    my ($buildCMDs, $runCMDs, $runNamedCMDs, $cmakeCMDs) = parse_file("$relfile");
+    my ($buildCMDs, $testCMDs, $testNamedCMDs, $cmakeCMDs) = parse_file("$relfile");
     if ($retBuildCMDs) {
         # print "BuildCMDs:\n";
         print "$_\n" for @$buildCMDs;
     }
-    if ($retRunCMDs) {
-        # print "RunCMDs:\n";
-        print "$_\n" for @$runCMDs;
+    if ($retTestCMDs) {
+        # print "TestCMDs:\n";
+        print "$_\n" for @$testCMDs;
     }
-    if ($retRunNamedCMDs) {
-        # print "RunNamedCMDs:\n";
-        print "$_\n" for @$runNamedCMDs;
+    if ($retTestNamedCMDs) {
+        # print "TestNamedCMDs:\n";
+        print "$_\n" for @$testNamedCMDs;
     }
     if ($retCmakeCMDs) {
         # print "CmakeCMDs:\n";
diff --git a/projects/hip/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp b/projects/hip/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp
index 30866da654..68af1f1736 100644
--- a/projects/hip/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp
+++ b/projects/hip/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/Functional/device/hipFuncGetDevice.cpp b/projects/hip/tests/src/Functional/device/hipFuncGetDevice.cpp
index ac2c844408..664378f233 100644
--- a/projects/hip/tests/src/Functional/device/hipFuncGetDevice.cpp
+++ b/projects/hip/tests/src/Functional/device/hipFuncGetDevice.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/Functional/device/hipFuncSetDevice.cpp b/projects/hip/tests/src/Functional/device/hipFuncSetDevice.cpp
index d65f6d1676..5ee086c8ac 100644
--- a/projects/hip/tests/src/Functional/device/hipFuncSetDevice.cpp
+++ b/projects/hip/tests/src/Functional/device/hipFuncSetDevice.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t EXCLUDE_HIP_PLATFORM
+ * TEST: %t EXCLUDE_HIP_PLATFORM
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/Functional/device/hipFuncSetDeviceFlags.cpp b/projects/hip/tests/src/Functional/device/hipFuncSetDeviceFlags.cpp
index f889011b9f..1de5c8c42f 100644
--- a/projects/hip/tests/src/Functional/device/hipFuncSetDeviceFlags.cpp
+++ b/projects/hip/tests/src/Functional/device/hipFuncSetDeviceFlags.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/compiler/hipClassKernel.cpp b/projects/hip/tests/src/compiler/hipClassKernel.cpp
index edc3808d83..8295c6c62f 100644
--- a/projects/hip/tests/src/compiler/hipClassKernel.cpp
+++ b/projects/hip/tests/src/compiler/hipClassKernel.cpp
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 #include "hipClassKernel.h"
diff --git a/projects/hip/tests/src/context/hipCtx_simple.cpp b/projects/hip/tests/src/context/hipCtx_simple.cpp
index 9d398ddf82..8f4e1a95bd 100644
--- a/projects/hip/tests/src/context/hipCtx_simple.cpp
+++ b/projects/hip/tests/src/context/hipCtx_simple.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/context/hipMemsetD8.cpp b/projects/hip/tests/src/context/hipMemsetD8.cpp
index 7ae0e06b8d..30766da776 100644
--- a/projects/hip/tests/src/context/hipMemsetD8.cpp
+++ b/projects/hip/tests/src/context/hipMemsetD8.cpp
@@ -24,13 +24,13 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * //Small copy
- * RUN: %t -N 10    --memsetval 0x42
+ * TEST: %t -N 10    --memsetval 0x42
  * // Oddball size
- * RUN: %t -N 10013 --memsetval 0x5a
+ * TEST: %t -N 10013 --memsetval 0x5a
  * // Big copy
- * RUN: %t -N 256M  --memsetval 0xa6
+ * TEST: %t -N 256M  --memsetval 0xa6
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp b/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp
index b5f6d25793..165c469941 100644
--- a/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp
+++ b/projects/hip/tests/src/deviceLib/hipAsynchronousStreams.cpp
@@ -20,7 +20,7 @@ THE SOFTWARE.
 /*
  * HIT_START
  * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipDeviceMalloc.cpp b/projects/hip/tests/src/deviceLib/hipDeviceMalloc.cpp
index 4ec10077c5..4af7614a5c 100644
--- a/projects/hip/tests/src/deviceLib/hipDeviceMalloc.cpp
+++ b/projects/hip/tests/src/deviceLib/hipDeviceMalloc.cpp
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 /* HIT_START
  * BUILD: %t %s NVCC_OPTIONS -std=c++11
- * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 #include "test_common.h"
diff --git a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp
index 6ebdebf6fc..7033dc2113 100644
--- a/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp
+++ b/projects/hip/tests/src/deviceLib/hipDeviceMemcpy.cpp
@@ -9,7 +9,7 @@
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
index 24a8320840..84f63f9f36 100644
--- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
+++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp
index 2620a3cbde..43691882a9 100644
--- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp
+++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
index 4bac6614f5..a3283622ff 100644
--- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
+++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp
index ab1deb477a..2f6557cab7 100644
--- a/projects/hip/tests/src/deviceLib/hipFloatMath.cpp
+++ b/projects/hip/tests/src/deviceLib/hipFloatMath.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp b/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp
index 6a0c91e890..4c23255d64 100644
--- a/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp
+++ b/projects/hip/tests/src/deviceLib/hipFloatMathPrecise.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipIntegerIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipIntegerIntrinsics.cpp
index b8fe1d27c4..42d218acc5 100644
--- a/projects/hip/tests/src/deviceLib/hipIntegerIntrinsics.cpp
+++ b/projects/hip/tests/src/deviceLib/hipIntegerIntrinsics.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp b/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp
index 6a3d5501e0..6553094603 100644
--- a/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp
+++ b/projects/hip/tests/src/deviceLib/hipMathFunctions.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
index ec38a5d327..760d65e555 100644
--- a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
+++ b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11 --gpu-architecture=sm_60
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
index 751ea88641..96ee8d2962 100644
--- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
+++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp
index 15c0e6c1e3..65ac046e01 100644
--- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp
+++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
index a830f45d40..1dd5737f86 100644
--- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
+++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipStdComplex.cpp b/projects/hip/tests/src/deviceLib/hipStdComplex.cpp
index 5d91c574ec..0a57889be6 100644
--- a/projects/hip/tests/src/deviceLib/hipStdComplex.cpp
+++ b/projects/hip/tests/src/deviceLib/hipStdComplex.cpp
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 /* HIT_START
  * BUILD: %t %s NVCC_OPTIONS -std=c++11
- * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 #include "test_common.h"
diff --git a/projects/hip/tests/src/deviceLib/hipTestClock.cpp b/projects/hip/tests/src/deviceLib/hipTestClock.cpp
index ee6dca8a42..b5bb5d9a05 100644
--- a/projects/hip/tests/src/deviceLib/hipTestClock.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestClock.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestDevice.cpp b/projects/hip/tests/src/deviceLib/hipTestDevice.cpp
index 2bb5163dc9..2cab5a91d8 100644
--- a/projects/hip/tests/src/deviceLib/hipTestDevice.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestDevice.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestDeviceDouble.cpp b/projects/hip/tests/src/deviceLib/hipTestDeviceDouble.cpp
index b29a9d2c20..d3dfdfc3b9 100644
--- a/projects/hip/tests/src/deviceLib/hipTestDeviceDouble.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestDeviceDouble.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * XXBUILD: %t %s ../test_common.cpp
- * XXRUN: %t
+ * XXTEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/projects/hip/tests/src/deviceLib/hipTestDeviceSymbol.cpp
index a17dd75a31..4c2dd279ea 100644
--- a/projects/hip/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestDeviceSymbol.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestDotFunctions.cpp b/projects/hip/tests/src/deviceLib/hipTestDotFunctions.cpp
index 51e005f7fb..247ee02026 100644
--- a/projects/hip/tests/src/deviceLib/hipTestDotFunctions.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestDotFunctions.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestFMA.cpp b/projects/hip/tests/src/deviceLib/hipTestFMA.cpp
index 2771fac585..0c22f33d9f 100644
--- a/projects/hip/tests/src/deviceLib/hipTestFMA.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestFMA.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp
index 125edd76eb..96447960c1 100644
--- a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp b/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp
index 92b93eda36..21141eb10a 100644
--- a/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp
index 10a3f33a27..17aa54a4cf 100644
--- a/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipTestNew.cpp b/projects/hip/tests/src/deviceLib/hipTestNew.cpp
index d644f8b483..bb96da4f7b 100644
--- a/projects/hip/tests/src/deviceLib/hipTestNew.cpp
+++ b/projects/hip/tests/src/deviceLib/hipTestNew.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipThreadFence.cpp b/projects/hip/tests/src/deviceLib/hipThreadFence.cpp
index 296d1519aa..f91bd7bca6 100644
--- a/projects/hip/tests/src/deviceLib/hipThreadFence.cpp
+++ b/projects/hip/tests/src/deviceLib/hipThreadFence.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp
index 44681e3b05..f05e6372d8 100644
--- a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp
+++ b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp
index a1db8de361..ba23931ee5 100644
--- a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp
+++ b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_anyall.cpp b/projects/hip/tests/src/deviceLib/hip_anyall.cpp
index 6815ce8cbf..c77a3cd25d 100644
--- a/projects/hip/tests/src/deviceLib/hip_anyall.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_anyall.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_ballot.cpp b/projects/hip/tests/src/deviceLib/hip_ballot.cpp
index 92cfabbd36..324047a6b1 100644
--- a/projects/hip/tests/src/deviceLib/hip_ballot.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_ballot.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_bitextract.cpp b/projects/hip/tests/src/deviceLib/hip_bitextract.cpp
index e791d66b34..5ee3974107 100644
--- a/projects/hip/tests/src/deviceLib/hip_bitextract.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_bitextract.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp b/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp
index bf1146b91a..d2c1006b5f 100644
--- a/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_bitinsert.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_brev.cpp b/projects/hip/tests/src/deviceLib/hip_brev.cpp
index cc48cca31f..363c9faac3 100644
--- a/projects/hip/tests/src/deviceLib/hip_brev.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_brev.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_clz.cpp b/projects/hip/tests/src/deviceLib/hip_clz.cpp
index 77ba980fe3..c9c6a5eeec 100644
--- a/projects/hip/tests/src/deviceLib/hip_clz.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_clz.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_ffs.cpp b/projects/hip/tests/src/deviceLib/hip_ffs.cpp
index a877b113f1..92193873b8 100644
--- a/projects/hip/tests/src/deviceLib/hip_ffs.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_ffs.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_mbcnt.cpp b/projects/hip/tests/src/deviceLib/hip_mbcnt.cpp
index 440e54c222..343d6069d4 100644
--- a/projects/hip/tests/src/deviceLib/hip_mbcnt.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_mbcnt.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_popc.cpp b/projects/hip/tests/src/deviceLib/hip_popc.cpp
index 643b1ca783..24b6f89304 100644
--- a/projects/hip/tests/src/deviceLib/hip_popc.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_popc.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_test_ldg.cpp b/projects/hip/tests/src/deviceLib/hip_test_ldg.cpp
index 3d623f9cc3..0f97c1f0ad 100644
--- a/projects/hip/tests/src/deviceLib/hip_test_ldg.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_test_ldg.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --gpu-architecture=sm_35
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp b/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp
index a090aa811b..24732573e6 100644
--- a/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/deviceLib/hip_trig.cpp b/projects/hip/tests/src/deviceLib/hip_trig.cpp
index 3519d597b3..b712636a69 100644
--- a/projects/hip/tests/src/deviceLib/hip_trig.cpp
+++ b/projects/hip/tests/src/deviceLib/hip_trig.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/hipC.c b/projects/hip/tests/src/hipC.c
index 1484adf379..b769dadd15 100644
--- a/projects/hip/tests/src/hipC.c
+++ b/projects/hip/tests/src/hipC.c
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s EXCLUDE_HIP_PLATFORM all
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/hipEnvVarDriver.cpp b/projects/hip/tests/src/hipEnvVarDriver.cpp
index e52aa46063..5063319ead 100644
--- a/projects/hip/tests/src/hipEnvVarDriver.cpp
+++ b/projects/hip/tests/src/hipEnvVarDriver.cpp
@@ -17,7 +17,7 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
 
 /* HIT_START
  * BUILD: %t %s test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/hipHcc.cpp b/projects/hip/tests/src/hipHcc.cpp
index d2a24b4d2b..e0a843d4f5 100644
--- a/projects/hip/tests/src/hipHcc.cpp
+++ b/projects/hip/tests/src/hipHcc.cpp
@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s
- * RUN: %t EXCLUDE_HIP_PLATFORM all
+ * TEST: %t EXCLUDE_HIP_PLATFORM all
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipDynamicShared.cpp b/projects/hip/tests/src/kernel/hipDynamicShared.cpp
index 8edd94b48c..d9a5942085 100644
--- a/projects/hip/tests/src/kernel/hipDynamicShared.cpp
+++ b/projects/hip/tests/src/kernel/hipDynamicShared.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM hcc
- * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipDynamicShared2.cpp b/projects/hip/tests/src/kernel/hipDynamicShared2.cpp
index 0e85c1a8c0..11f0640537 100644
--- a/projects/hip/tests/src/kernel/hipDynamicShared2.cpp
+++ b/projects/hip/tests/src/kernel/hipDynamicShared2.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM hcc
- * RUN: %t EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipEmptyKernel.cpp b/projects/hip/tests/src/kernel/hipEmptyKernel.cpp
index 3a42b33072..08ac5afa8d 100644
--- a/projects/hip/tests/src/kernel/hipEmptyKernel.cpp
+++ b/projects/hip/tests/src/kernel/hipEmptyKernel.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipGridLaunch.cpp b/projects/hip/tests/src/kernel/hipGridLaunch.cpp
index 4f7883c656..b06fe9776e 100644
--- a/projects/hip/tests/src/kernel/hipGridLaunch.cpp
+++ b/projects/hip/tests/src/kernel/hipGridLaunch.cpp
@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp b/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp
index f40014d399..0448ebf177 100644
--- a/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp
+++ b/projects/hip/tests/src/kernel/hipLanguageExtensions.cpp
@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipLaunchParm.cpp b/projects/hip/tests/src/kernel/hipLaunchParm.cpp
index aa9ccdf360..772ddc1032 100644
--- a/projects/hip/tests/src/kernel/hipLaunchParm.cpp
+++ b/projects/hip/tests/src/kernel/hipLaunchParm.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp b/projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp
index f5c30b07a4..e7cc713f28 100755
--- a/projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp
+++ b/projects/hip/tests/src/kernel/hipLaunchParmFunctor.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -O3
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipPrintfKernel.cpp b/projects/hip/tests/src/kernel/hipPrintfKernel.cpp
index 9b798de895..1d4fa5fe30 100644
--- a/projects/hip/tests/src/kernel/hipPrintfKernel.cpp
+++ b/projects/hip/tests/src/kernel/hipPrintfKernel.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipTestConstant.cpp b/projects/hip/tests/src/kernel/hipTestConstant.cpp
index b114f8105f..575e28966f 100644
--- a/projects/hip/tests/src/kernel/hipTestConstant.cpp
+++ b/projects/hip/tests/src/kernel/hipTestConstant.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipTestGlobalVariable.cpp b/projects/hip/tests/src/kernel/hipTestGlobalVariable.cpp
index 8ab8bef9c2..05b5b746e9 100644
--- a/projects/hip/tests/src/kernel/hipTestGlobalVariable.cpp
+++ b/projects/hip/tests/src/kernel/hipTestGlobalVariable.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipTestMallocKernel.cpp b/projects/hip/tests/src/kernel/hipTestMallocKernel.cpp
index d8084f8883..2a9ee9c6ad 100644
--- a/projects/hip/tests/src/kernel/hipTestMallocKernel.cpp
+++ b/projects/hip/tests/src/kernel/hipTestMallocKernel.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM all
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/hipTestMemKernel.cpp b/projects/hip/tests/src/kernel/hipTestMemKernel.cpp
index 8d928d3974..fa9b9fd5e9 100644
--- a/projects/hip/tests/src/kernel/hipTestMemKernel.cpp
+++ b/projects/hip/tests/src/kernel/hipTestMemKernel.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/kernel/inline_asm_vadd.cpp b/projects/hip/tests/src/kernel/inline_asm_vadd.cpp
index 1c5a77537b..32f1b54403 100644
--- a/projects/hip/tests/src/kernel/inline_asm_vadd.cpp
+++ b/projects/hip/tests/src/kernel/inline_asm_vadd.cpp
@@ -17,7 +17,7 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
 
 /* HIT_START
  * BUILD: %t %s
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp
index 3fac91bc43..205e3fcc04 100644
--- a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp
+++ b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp
@@ -24,9 +24,9 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t EXCLUDE_HIP_PLATFORM hcc
- * RUN: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc
- * RUN: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM hcc
+ * TEST: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc
+ * TEST: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipChooseDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipChooseDevice.cpp
index 8b27f4e2af..3868f9dc9e 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipChooseDevice.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipChooseDevice.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp
index cb5d035973..794da5fbe2 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetByPCIBusId.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetByPCIBusId.cpp
index ac476337fb..2e81415cbe 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetByPCIBusId.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetByPCIBusId.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
index 10c31f96b6..e86fa512ca 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp
index 56a8ceeb3c..7e6897e9f9 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
index 30866da654..68af1f1736 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp
index 6c3139eb6b..6fad8c1b9e 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
index ac2c844408..664378f233 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
index ed0b7cb715..ec82861013 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp b/projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp
index 33a605a2a1..9cdcac5e3f 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp
@@ -26,7 +26,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp
index 350a4bec7c..39f5fb7c04 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
index d65f6d1676..5ee086c8ac 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t EXCLUDE_HIP_PLATFORM
+ * TEST: %t EXCLUDE_HIP_PLATFORM
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp
index f889011b9f..1de5c8c42f 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp b/projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp
index 2cada53d20..a1923d2715 100644
--- a/projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp
+++ b/projects/hip/tests/src/runtimeApi/error/hipPeekAtLastError.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/event/hipEventRecord.cpp b/projects/hip/tests/src/runtimeApi/event/hipEventRecord.cpp
index bceb770fe0..a2a9beaf79 100644
--- a/projects/hip/tests/src/runtimeApi/event/hipEventRecord.cpp
+++ b/projects/hip/tests/src/runtimeApi/event/hipEventRecord.cpp
@@ -25,7 +25,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t --iterations 10
+ * TEST: %t --iterations 10
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/event/record_event.cpp b/projects/hip/tests/src/runtimeApi/event/record_event.cpp
index 3eb54e1735..8ba9ec3a27 100644
--- a/projects/hip/tests/src/runtimeApi/event/record_event.cpp
+++ b/projects/hip/tests/src/runtimeApi/event/record_event.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp b/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp
index 0c211eecfc..c022a40079 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipArray.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipHostGetFlags.cpp b/projects/hip/tests/src/runtimeApi/memory/hipHostGetFlags.cpp
index cc93273f8a..ebb323c223 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipHostGetFlags.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipHostGetFlags.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp b/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp
index 6a358c4fa6..3c902b3784 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp
@@ -22,7 +22,7 @@
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipHostRegister.cpp b/projects/hip/tests/src/runtimeApi/memory/hipHostRegister.cpp
index ff6ea963e9..7957a50142 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipHostRegister.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipHostRegister.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp
index bf7662f6b7..8f30195948 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMallocManaged.cpp
@@ -4,7 +4,7 @@
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t -N 256M
+ * TEST: %t -N 256M
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemPtrGetInfo.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemPtrGetInfo.cpp
index e470b1f325..2041ea41d1 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemPtrGetInfo.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemPtrGetInfo.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp
index d7acc9c177..c8ee170336 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy.cpp
@@ -22,11 +22,11 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN_NAMED: %t hipMemcpy-modes --tests 0x1
- * RUN_NAMED: %t hipMemcpy-size --tests 0x6
- * RUN_NAMED: %t hipMemcpy-dev-offsets --tests 0x10
- * RUN_NAMED: %t hipMemcpy-host-offsets --tests 0x20
- * RUN_NAMED: %t hipMemcpy-multithreaded --tests 0x8
+ * TEST_NAMED: %t hipMemcpy-modes --tests 0x1
+ * TEST_NAMED: %t hipMemcpy-size --tests 0x6
+ * TEST_NAMED: %t hipMemcpy-dev-offsets --tests 0x10
+ * TEST_NAMED: %t hipMemcpy-host-offsets --tests 0x20
+ * TEST_NAMED: %t hipMemcpy-multithreaded --tests 0x8
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp
index f546ae3ffd..b68cb3f708 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp
@@ -25,7 +25,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
index 13eca9c24d..ec3c075cdf 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
@@ -25,7 +25,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp
index ec7f0754d7..565cef91ee 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp
@@ -25,7 +25,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
index 57ec2038f5..d08ff2a7e8 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
@@ -25,7 +25,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp  EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp
index 265a5ced22..8ce02d6164 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy_simple.cpp
@@ -22,8 +22,8 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
- * RUN_NAMED: %t hipMemcpyAsync-simple --async
+ * TEST: %t
+ * TEST_NAMED: %t hipMemcpyAsync-simple --async
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp
index 8487d92877..f2f1e49c33 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemoryAllocateCoherentDriver.cpp
@@ -17,7 +17,7 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp
index abff987437..7edfe16b3c 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset.cpp
@@ -24,13 +24,13 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * //Small copy
- * RUN: %t -N 10    --memsetval 0x42 --memsetD32val 0x101
+ * TEST: %t -N 10    --memsetval 0x42 --memsetD32val 0x101
  * // Oddball size
- * RUN: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF
+ * TEST: %t -N 10013 --memsetval 0x5a --memsetD32val 0xDEADBEEF
  * // Big copy
- * RUN: %t -N 256M  --memsetval 0xa6 --memsetD32val 0xCAFEBABE
+ * TEST: %t -N 256M  --memsetval 0xa6 --memsetD32val 0xCAFEBABE
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp
index 51477e9eb2..956d26a323 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp
index 151a86a6c0..11bd656761 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipRandomMemcpyAsync.cpp b/projects/hip/tests/src/runtimeApi/memory/hipRandomMemcpyAsync.cpp
index bb25324e5c..98d07a316c 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipRandomMemcpyAsync.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipRandomMemcpyAsync.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/hipTestMemcpyPin.cpp b/projects/hip/tests/src/runtimeApi/memory/hipTestMemcpyPin.cpp
index 9cb1260b67..ab4b82b17d 100644
--- a/projects/hip/tests/src/runtimeApi/memory/hipTestMemcpyPin.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/hipTestMemcpyPin.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp
index 57810de02f..3db855c400 100644
--- a/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp
+++ b/projects/hip/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
- * RUN: %t EXCLUDE_HIP_PLATFORM all
+ * TEST: %t EXCLUDE_HIP_PLATFORM all
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp b/projects/hip/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp
index 6af87edb25..9ed9beb078 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/module/hipFuncSetCacheConfig.cpp b/projects/hip/tests/src/runtimeApi/module/hipFuncSetCacheConfig.cpp
index 350a4bec7c..39f5fb7c04 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipFuncSetCacheConfig.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipFuncSetCacheConfig.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp
index 249fbeb792..af48d5a367 100644
--- a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp
+++ b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadDevice.cpp
@@ -1,8 +1,8 @@
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN_NAMED: %t hipMultiThreadDevice-serial --tests 0x1
- * RUN_NAMED: %t hipMultiThreadDevice-pyramid --tests 0x4
- * RUN_NAMED: %t hipMultiThreadDevice-nearzero --tests 0x10
+ * TEST_NAMED: %t hipMultiThreadDevice-serial --tests 0x1
+ * TEST_NAMED: %t hipMultiThreadDevice-pyramid --tests 0x4
+ * TEST_NAMED: %t hipMultiThreadDevice-nearzero --tests 0x10
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams1.cpp b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams1.cpp
index 418a411613..bed7ba7b9f 100644
--- a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams1.cpp
+++ b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams1.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
index 7fadc60a05..2a0093be8e 100644
--- a/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
+++ b/projects/hip/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipNullStream.cpp b/projects/hip/tests/src/runtimeApi/stream/hipNullStream.cpp
index bf4699ddb6..16c9b1917b 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipNullStream.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipNullStream.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp
index 37e5cba920..d94d7d82d4 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp
index 582da67d85..ba2b140611 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp
index 19ae7c11b0..8d1129e567 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamL5.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamL5.cpp
index 31f8d128c0..32bce643b7 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipStreamL5.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamL5.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp
index cf25d0bd2b..c365dce70a 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamSync2.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp
index 8dbf37186e..4ddc7b98b0 100644
--- a/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp
+++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * ZZZBUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
- * ZZZRUN: %t
+ * ZZZTEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/runtimeApi/synchronization/copy_coherency.cpp b/projects/hip/tests/src/runtimeApi/synchronization/copy_coherency.cpp
index be813e1a7d..4dce05ef72 100644
--- a/projects/hip/tests/src/runtimeApi/synchronization/copy_coherency.cpp
+++ b/projects/hip/tests/src/runtimeApi/synchronization/copy_coherency.cpp
@@ -27,7 +27,7 @@ THE SOFTWARE.
 // TODO - add code object support here.
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp
index 6f9de9c094..b6c016ecc0 100644
--- a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp
+++ b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp
@@ -1,6 +1,6 @@
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 #include <stdio.h>
diff --git a/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp b/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp
index fa63efc1bb..546e9775e8 100644
--- a/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp
+++ b/projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp
@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/texture/hipGetChanDesc.cpp b/projects/hip/tests/src/texture/hipGetChanDesc.cpp
index 9034668ccf..bbdfd85392 100644
--- a/projects/hip/tests/src/texture/hipGetChanDesc.cpp
+++ b/projects/hip/tests/src/texture/hipGetChanDesc.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/texture/hipTextureObj1DFetch.cpp b/projects/hip/tests/src/texture/hipTextureObj1DFetch.cpp
index faf6541ce1..206173b954 100644
--- a/projects/hip/tests/src/texture/hipTextureObj1DFetch.cpp
+++ b/projects/hip/tests/src/texture/hipTextureObj1DFetch.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /*HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 
diff --git a/projects/hip/tests/src/texture/hipTextureObj2D.cpp b/projects/hip/tests/src/texture/hipTextureObj2D.cpp
index e214295989..148c966e80 100644
--- a/projects/hip/tests/src/texture/hipTextureObj2D.cpp
+++ b/projects/hip/tests/src/texture/hipTextureObj2D.cpp
@@ -1,6 +1,6 @@
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 #include <stdlib.h>
diff --git a/projects/hip/tests/src/texture/hipTextureRef2D.cpp b/projects/hip/tests/src/texture/hipTextureRef2D.cpp
index b912f789a7..2b19d1b5b3 100644
--- a/projects/hip/tests/src/texture/hipTextureRef2D.cpp
+++ b/projects/hip/tests/src/texture/hipTextureRef2D.cpp
@@ -1,6 +1,6 @@
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * RUN: %t
+ * TEST: %t
  * HIT_END
  */
 #include <stdlib.h>

From 4b274cc6f8883d89ca07df5eb61a4744b999d113 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Mon, 8 Apr 2019 16:13:07 +0530
Subject: [PATCH 2/9] [hit] Remove CUSTOM_CMD

Change-Id: Ia156fe6aab9cfcc11284823ea5131e33eaf962bc


[ROCm/hip commit: 622ea32964d22203fd055b9f397e16e7d05721cc]
---
 projects/hip/tests/README.md     | 12 ------------
 projects/hip/tests/hit/HIT.cmake | 14 --------------
 projects/hip/tests/hit/parser    | 27 ++++++---------------------
 3 files changed, 6 insertions(+), 47 deletions(-)

diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md
index 77929340c0..e0a7254963 100644
--- a/projects/hip/tests/README.md
+++ b/projects/hip/tests/README.md
@@ -75,18 +75,6 @@ TEST: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hc
 ```
 
 
-#### CMAKECMD command
-
-The supported syntax for the CMAKECMD command is:
-```
-CMAKECMD: <cmake_command> <options_to_cmake_command>
-```
-cmake_command: refers to any of the commands supported by ```cmake -E``` as specified in the [cmake documentation](https://cmake.org/cmake/help/latest/manual/cmake.1.html#command-line-tool-mode). Note that the commands are limited by the version of cmake the user is running.
-options_to_cmake_command: refers to the arguments supported by the specific cmake_command. The arguments are parsed by HIT to replace special markers. The markers supported by HIT are:
-%S: Refers to the source directory containing the current source file.
-%B: Refers to the build directory for the current cmake project i.e. CMAKE_CURRENT_BINARY_DIR.
-
-
 ### Running tests:
 ```
 ctest
diff --git a/projects/hip/tests/hit/HIT.cmake b/projects/hip/tests/hit/HIT.cmake
index 3233a7832a..2823d8f050 100644
--- a/projects/hip/tests/hit/HIT.cmake
+++ b/projects/hip/tests/hit/HIT.cmake
@@ -155,20 +155,6 @@ macro(HIT_ADD_FILES _dir _label _parent)
             endif()
         endforeach()
 
-        # Run cmake commands
-        execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --cmakeCMDs ${file}
-            OUTPUT_VARIABLE _contents
-            ERROR_QUIET
-            WORKING_DIRECTORY ${_dir}
-            OUTPUT_STRIP_TRAILING_WHITESPACE)
-        string(REGEX REPLACE "\n" ";" _contents "${_contents}")
-        string(REGEX REPLACE "%S" ${_dir} _contents "${_contents}")
-        string(REGEX REPLACE "%B" ${CMAKE_CURRENT_BINARY_DIR} _contents "${_contents}")
-        foreach(_cmd ${_contents})
-            string(REGEX REPLACE " " ";" _cmd "${_cmd}")
-            execute_process(COMMAND ${CMAKE_COMMAND} -E ${_cmd})
-        endforeach()
-
         # Add tests
         execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --testCMDs ${file}
             OUTPUT_VARIABLE _contents
diff --git a/projects/hip/tests/hit/parser b/projects/hip/tests/hit/parser
index cb2330026e..c74c0f9d01 100755
--- a/projects/hip/tests/hit/parser
+++ b/projects/hip/tests/hit/parser
@@ -8,7 +8,7 @@ use File::Spec;
 sub parse_file {
     my $file = shift;
     (my $exe = $file) =~ s/\.[^.]+$//g;
-    my (@buildCMDs, @testCMDs, @testNamedCMDs, @cmakeCMDs);
+    my (@buildCMDs, @testCMDs, @testNamedCMDs);
     if (open (SOURCE, '<:encoding(UTF-8)', "$file")) {
         while (<SOURCE>) {
             my $line=$_;
@@ -36,17 +36,10 @@ sub parse_file {
                 $line =~ s/\R//g;               # Remove line endings
                 push @testNamedCMDs, $line;
             }
-            # Look for CMAKECMD instructions
-            if ($line =~ /^ \* CMAKECMD:/) {
-                $line =~ s/^ \* CMAKECMD: //g;  # Remove " * CMAKECMD: "
-                # Substitute %S -> srcdir and %B -> builddir happens in cmake
-                $line =~ s/\R//g;               # Remove line endings
-                push @cmakeCMDs, $line;
-            }
         }
         close(SOURCE);
     }
-    return (\@buildCMDs, \@testCMDs, \@testNamedCMDs, \@cmakeCMDs);
+    return (\@buildCMDs, \@testCMDs, \@testNamedCMDs);
 }
 
 # Exit if no arguments specified
@@ -60,9 +53,8 @@ my @options = ();
 my $retBuildCMDs = 0;
 my $retTestCMDs = 0;
 my $retTestNamedCMDs = 0;
-my $retCmakeCMDs = 0;
 foreach $arg (@ARGV) {
-    if ($retBuildCMDs or $retTestCMDs or $retTestNamedCMDs or $retCmakeCMDs) {
+    if ($retBuildCMDs or $retTestCMDs or $retTestNamedCMDs) {
         push (@options, $arg);
     }
     if ($arg eq '--buildCMDs') {
@@ -74,21 +66,18 @@ foreach $arg (@ARGV) {
     if ($arg eq '--testNamedCMDs') {
         $retTestNamedCMDs = 1;
     }
-    if ($arg eq '--cmakeCMDs') {
-        $retCmakeCMDs = 1;
-    }
 }
 
 # Atleast one command needs to be specified
-if (($retBuildCMDs eq 0) and ($retTestCMDs eq 0) and ($retTestNamedCMDs eq 0) and ($retCmakeCMDs eq 0)) {
-    die "Usage: $0 <--buildCMDs|--testCMDs|--testNamedCMDs|--cmakeCMDs> FILENAMEs\n";
+if (($retBuildCMDs eq 0) and ($retTestCMDs eq 0) and ($retTestNamedCMDs eq 0)) {
+    die "Usage: $0 <--buildCMDs|--testCMDs|--testNamedCMDs> FILENAMEs\n";
 }
 
 # Iterate over input files
 foreach $file (@options) {
     # Convert absolute path to path relative to working directory
     my $relfile = File::Spec->abs2rel($file);
-    my ($buildCMDs, $testCMDs, $testNamedCMDs, $cmakeCMDs) = parse_file("$relfile");
+    my ($buildCMDs, $testCMDs, $testNamedCMDs) = parse_file("$relfile");
     if ($retBuildCMDs) {
         # print "BuildCMDs:\n";
         print "$_\n" for @$buildCMDs;
@@ -101,10 +90,6 @@ foreach $file (@options) {
         # print "TestNamedCMDs:\n";
         print "$_\n" for @$testNamedCMDs;
     }
-    if ($retCmakeCMDs) {
-        # print "CmakeCMDs:\n";
-        print "$_\n" for @$cmakeCMDs;
-    }
 }
 
 # vim: ts=4:sw=4:expandtab:smartindent

From ba620dfbe440c4e7858b3aa66dd1e57772ab046e Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 9 May 2019 08:41:05 +0530
Subject: [PATCH 3/9] [hit] Add support for BUILD_CMD

[ROCm/hip commit: 49a2d785d0d4192f81a279bd39a0f04af00dd3e1]
---
 projects/hip/tests/README.md     | 14 +++++++++
 projects/hip/tests/hit/HIT.cmake | 50 ++++++++++++++++++++++++++++++++
 projects/hip/tests/hit/parser    | 41 ++++++++++++++++++++------
 3 files changed, 96 insertions(+), 9 deletions(-)

diff --git a/projects/hip/tests/README.md b/projects/hip/tests/README.md
index e0a7254963..b1f70048a1 100644
--- a/projects/hip/tests/README.md
+++ b/projects/hip/tests/README.md
@@ -57,6 +57,20 @@ NVCC_OPTIONS: All options specified after this delimiter are passed to hipcc on
 EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms.
 
 
+#### BUILD_CMD command
+
+The supported syntax for the BUILD_CMD command is:
+```
+BUILD_CMD: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
+```
+%s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path).
+%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
+%hc: refers to hipcc pointed to by $CMAKE_INSTALL_PREFIX/bin/hipcc.
+%cc: refers to system c compiler pointed to by $CC.
+%S: refers to path to current source file.
+%T: refers to path to current build target.
+
+
 #### TEST command
 
 The supported syntax for the TEST command is:
diff --git a/projects/hip/tests/hit/HIT.cmake b/projects/hip/tests/hit/HIT.cmake
index 2823d8f050..f62ba65657 100644
--- a/projects/hip/tests/hit/HIT.cmake
+++ b/projects/hip/tests/hit/HIT.cmake
@@ -55,6 +55,29 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
     endforeach()
 endmacro()
 
+# Helper macro to parse CUSTOM BUILD instructions
+macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms)
+    set(${_target})
+    set(${_buildcmd} " ")
+    set(${_exclude_platforms})
+    set(_target_found FALSE)
+    set(_exclude_platforms_found FALSE)
+    foreach(arg ${ARGN})
+        if(NOT _target_found)
+            set(_target_found TRUE)
+            set(${_target} ${arg})
+        elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
+            set(_exclude_platforms_found TRUE)
+        else()
+            if(_exclude_platforms_found)
+                set(${_exclude_platforms} ${arg})
+            else()
+                list(APPEND ${_buildcmd} ${arg})
+            endif()
+        endif()
+    endforeach()
+endmacro()
+
 # Helper macro to parse TEST instructions
 macro(PARSE_TEST_COMMAND _target _arguments _exclude_platforms)
     set(${_target})
@@ -155,6 +178,33 @@ macro(HIT_ADD_FILES _dir _label _parent)
             endif()
         endforeach()
 
+        # Custom build commands
+        execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --customBuildCMDs ${file}
+            OUTPUT_VARIABLE _contents
+            ERROR_QUIET
+            WORKING_DIRECTORY ${_dir}
+            OUTPUT_STRIP_TRAILING_WHITESPACE)
+        string(REGEX REPLACE "\n" ";" _contents "${_contents}")
+        string(REGEX REPLACE "%hc" "${HIP_HIPCC_EXECUTABLE}" _contents "${_contents}")
+        string(REGEX REPLACE "%cc" "${CC}" _contents "${_contents}")
+        string(REGEX REPLACE "%S" ${_dir} _contents "${_contents}")
+        string(REGEX REPLACE "%T" ${_label} _contents "${_contents}")
+        foreach(_cmd ${_contents})
+            string(REGEX REPLACE " " ";" _cmd "${_cmd}")
+            parse_custombuild_command(_target _buildcmd _exclude_platforms ${_cmd})
+            string(REGEX REPLACE "/" "." target ${_label}/${_target})
+            insert_into_map("_exclude" "${target}" "${_exclude_platforms}")
+            if(_exclude_platforms STREQUAL "all" OR _exclude_platforms STREQUAL ${HIP_PLATFORM})
+            else()
+                string(REGEX REPLACE ";" " " _buildcmd "${_buildcmd}")
+                #string(CONCAT buildscript ${CMAKE_CURRENT_BINARY_DIR}/${target} ".sh")
+                #file(WRITE ${buildscript} ${_buildcmd})
+                #add_custom_target(${target} COMMAND ${buildscript})
+                add_custom_target(${target} COMMAND sh -c "${_buildcmd}")
+                add_dependencies(${_parent} ${target})
+            endif()
+        endforeach()
+
         # Add tests
         execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --testCMDs ${file}
             OUTPUT_VARIABLE _contents
diff --git a/projects/hip/tests/hit/parser b/projects/hip/tests/hit/parser
index c74c0f9d01..8381c78a3c 100755
--- a/projects/hip/tests/hit/parser
+++ b/projects/hip/tests/hit/parser
@@ -4,16 +4,21 @@ use 5.006; use v5.10.1;
 use File::Basename;
 use File::Spec;
 
+my $patBUILD = "^".quotemeta(" * BUILD:");
+my $patTEST = "^".quotemeta(" * TEST:");
+my $patTEST_NAMED = "^".quotemeta(" * TEST_NAMED:");
+my $patBUILD_CMD = "^".quotemeta(" * BUILD_CMD:");
+
 # Scan input file for HIT information
 sub parse_file {
     my $file = shift;
     (my $exe = $file) =~ s/\.[^.]+$//g;
-    my (@buildCMDs, @testCMDs, @testNamedCMDs);
+    my (@buildCMDs, @testCMDs, @testNamedCMDs, @customBuildCMDs);
     if (open (SOURCE, '<:encoding(UTF-8)', "$file")) {
         while (<SOURCE>) {
             my $line=$_;
             # Look for BUILD instructions
-            if ($line =~ /^ \* BUILD:/) {
+            if ($line =~ /$patBUILD/) {
                 $line =~ s/^ \* BUILD: //g;     # Remove " * BUILD: "
                 $line =~ s/%s/$file/g;          # Substitute %s -> filename
                 $line =~ s/%t/$exe/g;           # Substitute %t -> targetname
@@ -21,7 +26,7 @@ sub parse_file {
                 push @buildCMDs, $line;
             }
             # Look for TEST instructions
-            if ($line =~ /^ \* TEST:/) {
+            if ($line =~ /$patTEST/) {
                 $line =~ s/^ \* TEST: //g;      # Remove " * TEST: "
                 $line =~ s/%s/$file/g;          # Substitute %s -> filename
                 $line =~ s/%t/$exe/g;           # Subsitute %t -> targetname
@@ -29,17 +34,27 @@ sub parse_file {
                 push @testCMDs, $line;
             }
             # Look for TEST_NAMED instructions
-            if ($line =~ /^ \* TEST_NAMED:/) {
+            if ($line =~ /$patTEST_NAMED/) {
                 $line =~ s/^ \* TEST_NAMED: //g;# Remove " * TEST_NAMED: "
                 $line =~ s/%s/$file/g;          # Substitute %s -> filename
                 $line =~ s/%t/$exe/g;           # Subsitute %t -> targetname
                 $line =~ s/\R//g;               # Remove line endings
                 push @testNamedCMDs, $line;
             }
+            # Look for BUILD_CMD instructions
+            if ($line =~ /$patBUILD_CMD/) {
+                $line =~ s/^ \* BUILD_CMD: //g; # Remove " * BUILD_CMD: "
+                $line =~ s/%s/$file/g;          # Substitute %s -> filename
+                $line =~ s/%t/$exe/g;           # Substitute %t -> targetname
+                # Substitute %hc -> /path/to/hipcc and %cc -> /path/to/cc happens in cmake
+                # Substitute %S -> src dir and %T -> target build dir happens in cmake
+                $line =~ s/\R//g;               # Remove line endings
+                push @customBuildCMDs, $line;
+            }
         }
         close(SOURCE);
     }
-    return (\@buildCMDs, \@testCMDs, \@testNamedCMDs);
+    return (\@buildCMDs, \@testCMDs, \@testNamedCMDs, \@customBuildCMDs);
 }
 
 # Exit if no arguments specified
@@ -53,8 +68,9 @@ my @options = ();
 my $retBuildCMDs = 0;
 my $retTestCMDs = 0;
 my $retTestNamedCMDs = 0;
+my $retCustomBuildCMDs = 0;
 foreach $arg (@ARGV) {
-    if ($retBuildCMDs or $retTestCMDs or $retTestNamedCMDs) {
+    if ($retBuildCMDs or $retTestCMDs or $retTestNamedCMDs or $retCustomBuildCMDs) {
         push (@options, $arg);
     }
     if ($arg eq '--buildCMDs') {
@@ -66,18 +82,21 @@ foreach $arg (@ARGV) {
     if ($arg eq '--testNamedCMDs') {
         $retTestNamedCMDs = 1;
     }
+    if ($arg eq '--customBuildCMDs') {
+        $retCustomBuildCMDs = 1;
+    }
 }
 
 # Atleast one command needs to be specified
-if (($retBuildCMDs eq 0) and ($retTestCMDs eq 0) and ($retTestNamedCMDs eq 0)) {
-    die "Usage: $0 <--buildCMDs|--testCMDs|--testNamedCMDs> FILENAMEs\n";
+if (($retBuildCMDs eq 0) and ($retTestCMDs eq 0) and ($retTestNamedCMDs eq 0) and($retCustomBuildCMDs eq 0)) {
+    die "Usage: $0 <--buildCMDs|--testCMDs|--testNamedCMDs|--customBuildCMDs> FILENAMEs\n";
 }
 
 # Iterate over input files
 foreach $file (@options) {
     # Convert absolute path to path relative to working directory
     my $relfile = File::Spec->abs2rel($file);
-    my ($buildCMDs, $testCMDs, $testNamedCMDs) = parse_file("$relfile");
+    my ($buildCMDs, $testCMDs, $testNamedCMDs, $customBuildCMDs) = parse_file("$relfile");
     if ($retBuildCMDs) {
         # print "BuildCMDs:\n";
         print "$_\n" for @$buildCMDs;
@@ -90,6 +109,10 @@ foreach $file (@options) {
         # print "TestNamedCMDs:\n";
         print "$_\n" for @$testNamedCMDs;
     }
+    if ($retCustomBuildCMDs) {
+        # print "CustomBuildCMDs:\n";
+        print "$_\n" for @$customBuildCMDs;
+    }
 }
 
 # vim: ts=4:sw=4:expandtab:smartindent

From 09f8c228cfa32ea0ac6eda55498a5000e5a990d0 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 9 May 2019 08:42:00 +0530
Subject: [PATCH 4/9] [dtests] Add hipModule test to unit tests

Change-Id: I1dac38f8580265e2e9c82d88e4f070a2ff87f60b


[ROCm/hip commit: 79843f3b1287fb31dabb9c0e7c0c422c98303410]
---
 .../tests/src/runtimeApi/module/hipModule.cpp |  61 ++++++++++--------
 .../src/runtimeApi/module/vcpy_kernel.code    | Bin 18811 -> 0 bytes
 .../src/runtimeApi/module/vcpy_kernel.cpp     |   6 +-
 3 files changed, 38 insertions(+), 29 deletions(-)
 delete mode 100755 projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code

diff --git a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
index 22b4e7ff49..019645ea32 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
@@ -17,6 +17,13 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+/* HIT_START
+ * BUILD_CMD: vcpy_kernel.code %hc --genco %S/vcpy_kernel.cpp -o vcpy_kernel.code
+ * BUILD: %t %s ../../test_common.cpp
+ * TEST: %t
+ * HIT_END
+ */
+
 #include "hip/hip_runtime.h"
 #include "hip/hip_runtime_api.h"
 #include <iostream>
@@ -33,13 +40,15 @@ THE SOFTWARE.
 #define fileName "vcpy_kernel.code"
 #define kernel_name "hello_world"
 
-__global__ void Cpy(hipLaunchParm lp, float* Ad, float* Bd) {
-    int tx = threadIdx.x;
-    Bd[tx] = Ad[tx];
-}
+#define HIP_CHECK(status)                                                                          \
+    if (status != hipSuccess) {                                                                    \
+        std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl;            \
+        exit(0);                                                                                   \
+    }
 
 int main() {
-    float *A, *B, *Ad, *Bd;
+    float *A, *B;
+    hipDeviceptr_t Ad, Bd;
     A = new float[LEN];
     B = new float[LEN];
 
@@ -48,32 +57,38 @@ int main() {
         B[i] = 0.0f;
     }
 
+    HIPCHECK(hipInit(0));
+
+    hipDevice_t device;
+    hipCtx_t context;
+    HIPCHECK(hipDeviceGet(&device, 0));
+    HIPCHECK(hipCtxCreate(&context, 0, device));
+
     HIPCHECK(hipMalloc((void**)&Ad, SIZE));
     HIPCHECK(hipMalloc((void**)&Bd, SIZE));
-
     HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
     HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
+
     hipModule_t Module;
     hipFunction_t Function;
+    hipFunction_t f;
     HIPCHECK(hipModuleLoad(&Module, fileName));
     HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
-    hipFunction_t f;
-    HIPCHECK(hipModuleGetFunction(&f, Module, kernel_name));
-    assert(f == Function);
+
     hipStream_t stream;
     HIPCHECK(hipStreamCreate(&stream));
-    void* args[2] = {&Ad, &Bd};
 
-    std::vector<void*> argBuffer(5);
-    memcpy(&argBuffer[3], &Ad, sizeof(void*));
-    memcpy(&argBuffer[4], &Bd, sizeof(void*));
+    struct {
+        void* _Ad;
+        void* _Bd;
+    } args;
+    args._Ad = (void*) Ad;
+    args._Bd = (void*) Bd;
+    size_t size = sizeof(args);
 
-    size_t size = argBuffer.size() * sizeof(void*);
-
-    void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0], HIP_LAUNCH_PARAM_BUFFER_SIZE,
-                      &size, HIP_LAUNCH_PARAM_END};
-
-    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config);
+    void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
+                      HIP_LAUNCH_PARAM_END};
+    HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config));
 
     HIPCHECK(hipStreamDestroy(stream));
 
@@ -83,12 +98,6 @@ int main() {
         assert(A[i] == B[i]);
     }
 
-    std::vector<hipFunction_t> vec(1024 * 1024 * 64);
-    for (unsigned i = 0; i < 1024 * 1024 * 64; i++) {
-        hipFunction_t func;
-        hipModuleGetFunction(&func, Module, kernel_name);
-        vec[i] = func;
-    }
+	HIPCHECK(hipCtxDestroy(context));
     passed();
-    return 0;
 }
diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.code
deleted file mode 100755
index 4fd6472496c32a392baa767681f333d69a37ce29..0000000000000000000000000000000000000000
GIT binary patch
literal 0
HcmV?d00001

literal 18811
zcmeHPOKclO7@n+Mw++u0icpKtLQ9*X#92RLJCUfICZ#PTX$S?W<*`|N*RjDb)|*l&
zO-m8)BE(B@Kmu{%fDku&BNAs2aN*J;NEH%yD)qqs&(81cY&KO<f=ZU1DErU<&2#_n
zpV{%+GiMG=JgD^aFbXk<aP3wXDT;(kqap&L$O>!|TyJEXSpxJp$K3RyD<IVki=t|h
zn*@$9mRl(t+w7*0WTHP<ZUZI>-7nQWZaI>aj^(hg-HsJefl%aEHtUM6q;d-2|D+t<
z-^}lh$?vZyD)z})0h1JFVvj&M-kwOmR)KBt_OQbkeDYJFlO)@t#6DyDqm$#j7{(Fb
z0srU^h{xGRo)6S62Kd40$K1S!p#F^_@4&!7?<v)2n90fU$1S^4Yc)sIC)9Mms;N))
z^Mq-;)^sX|o0S%(>`9GeMzZ{G=gyt#p|LS_=CP@xhbIrH2WDoD%siP-tM_$O!%-`>
zx&<02^;(n1%qN?B`qe$@Ro9SNbq(27*D$o|8gi?yA;0Px3ahI@eM&vun@lEq(brZA
z=%d?lCJ^($v=~4p;9C*j<(5<hev0tjGX5UneKP(f;rGk<BH{fq{wv{08UK@TO~yAs
zzj!~>GQNZGtc<4#&&l{hgx}@j{h(u*CH!tz9~bxu!c`Zi{O1Y3SH`~}e5;G^67^gl
z{0<lI6ZnsW-!A9>o$x#5{8tFy?&4Wd&o)r;9|O)SxA7QbHwcN{!FfwG!_gbzDjpkL
zzX`5ec#JW0rd!~jrbPiuLNv|iH%thchxxMy($MX2j=UNEdC-qfrPBJri6i?*Cv<&E
zPi1rZ@l>j6)$1+&`IcQbQ@J_2X6kigzFDs7ZNqL%S02@Ly#l(^vQxvlUOsu!XtZ_5
zhHDFFxC1pFrXQ|MSD3C>Yi+$+)@!z2O_!}&U58oJ54So2sCBZP4uC85R?7~+>#ez<
z>~^^wklF4$XFCCSr`8NYgEKoeSYdf~!!!drQfU|gXk)$}fOh6f%TdglUDg_))!H(2
zJ~%VNLAg7I$*Um0h-74e$O4fCA`3(o_zzlu);v9YH3gsc2w0)7J?txX9?~}R71dHE
zv3%yhCifG1H!&WM2Q8(#>DM!&u*4!RNk3x>tbE`7<s7UK&jF?ac7~n*4pw>T9wo8R
zvoTR%Z2BUsAL9yOF@-Jk^dt&xo~M0L!NnkZ8S<~MKy+U#|7->RaV<_E;ag?@xA|8V
zvETcdvWOIN&XHS1S%41$@R4Huaf<hIts&z7s{F&pyoi6<mxAa2P>z50`Y#M7`Y`)I
zH2$wL{@H60|MK7mEqMInyEh7n?i(~0Sx?V@SpO~JT;^Q|C`x>Rv5P;mST8IQ!u34(
zcg8<{YqS*B6aQY`SuDQqvJ&6NwrzmdzTx7X*jX$7<?)Xz>7}rq_+M%K$Jify{9jfQ
z<LD2U<B5H-5ctJf@h`9cLNoryVf=Hxqw!DUIr#Y>QvQp;e=(Z>iF=g6{Cj8o_4xjW
zp8M==W#RH$@a|{Z26$%~&S$=t&9$2UL&N_^!2gF4|6U)F;~^jauK@q&BK~F0y;{im
ze-rqBBjVr7xEv2T|1SgoXCwY)&b?a5`9B5xzZmiFWn7Mjoc}Yx|3bvS%(+(!IsYF3
z|L;fqdl{GGA?N=z@c&Z8zs$K;3pxMq0RQLzKm2={jpC67A`3(oh%69UAhN*!!vgff
z{LjS~IJCb~Y}4%ilK&8H!h}GZa<_<$l=QvNCsdJ}ZRv`exyyB<Ij5d*w}PwLq?XiF
zwXbA!EK_YY`F`&dfT_He%M4`l`J85H1BG-YtqoPUUQQ3GdzoC#5U-}M>NxGr$l%~y
z&8g0plI2!oaK<pLR`W>PYL2xU?Rm!<<Rz1kMYW7ldbm<bo0*|}rJ$wDh0JiOY+5-j
zRT(xdOEYuXTu#&Ws@GO?y?&y>D;`Q&W+`P^rcuhK$^~;clh0<d1;Z-mhlVpHt&}NM
zhP6U!FPo?}pSA1}bz)-NwFP6z&I=94D1mfrS5_(3Y&jN7nlCgV3rkjwPSw>Mn<X9V
zq{B6%Q7f~gyB{A+(5%$v7@DGC!T!)S33wI&x`3d4DIx$mSAZHLzaae{h@fv#dGa@s
ze~B+5u^dL}_XOmJdjZ3}w^W||jpR3whJn(leUgKahQCUobn<7C<j+_K2Krvq)fmeQ
ztn22u0T_dj-SZ|~bXf%t0FGt5zl<r0j6%)|mHJEfPt80C>Ha=pk=Y_%mj{1dIHrM@
z*L+#2{C&Oxq>{va)ZM;&>Uee3Ri5@>7^QLqqfWLb4CSS9SNlmMqL1JlL-$Yo<G${o
z_AeZC%YW}{(5JeY0<Jut0?ZVbC&utD0y-T7ko?wLDh53yf0Xe!BR`XI@Mt2~lMq8s
zDS7}Ajqwo}gp;+F($I!M=Of^D{h;#^1_<Bn)}aamjuGA`<HrfVU&d|1`(^x1!hPo>
zd`7r%OltRA!hPo>{6M(xe1zW!_nnV$g>c{b2wUO)xE*}wBj7m(99N+r1UesKAK~}P
z_))@r=Oa8rxbJ*~vxNK3NBEd<;qY{x!q<fR&PQPS__3+c$-`sJbHafCD1hjM0NEnA
OlYiO%{_6EQA>ePX$&UR1

diff --git a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp
index c99e3df985..d7c23a8810 100644
--- a/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/vcpy_kernel.cpp
@@ -1,5 +1,5 @@
 /*
-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
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 #include "hip/hip_runtime.h"
 
-extern "C" __global__ void hello_world(hipLaunchParm lp, float* a, float* b) {
-    int tx = threadIdx.x;
+extern "C" __global__ void hello_world(float* a, float* b) {
+    int tx = hipThreadIdx_x;
     b[tx] = a[tx];
 }

From 8dd98b8c82f3d0dd3fa45459ac81e199ac26bee1 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 9 May 2019 18:03:42 +0530
Subject: [PATCH 5/9] [dtests] Add complex_loading_behavior test

Change-Id: Iadf135cb727a1a3761abef20336d652b159c7dcd


[ROCm/hip commit: eb637766b9451db7c8613d9c687b7e2959b8be71]
---
 .../dynamicLoading/complex_loading_behavior.cpp   | 15 ++++++---------
 1 file changed, 6 insertions(+), 9 deletions(-)

diff --git a/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp b/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp
index b8937e0582..ebf532a5c4 100644
--- a/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp
+++ b/projects/hip/tests/src/dynamicLoading/complex_loading_behavior.cpp
@@ -20,13 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
-/**
- * Build the test:
- * hipcc complex_loading_behavior.cpp -o libfoo.so -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY
- * hipcc complex_loading_behavior.cpp -o complex_loading_behavior -ldl
- *
- * Run the test:
- * ./complex_loading_behavior
+/* HIT_START
+ * BUILD_CMD: libfoo %hc %S/%s -o libfoo.so -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY
+ * BUILD_CMD: %t %hc %S/%s -o %T/%t -ldl
+ * TEST: %t
+ * HIT_END
  */
 
 #if !defined(TEST_SHARED_LIBRARY)
@@ -93,8 +91,7 @@ int launch_local_kernel() {
     free(B_h);
     free(C_h);
 
-    std::cout << "local launch succedded\n";
-
+    std::cout << "PASSED!\n";
     return 0;
 }
 

From 242aeaf133869d756e6a85f1c8b3145c5cc97316 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 9 May 2019 18:17:19 +0530
Subject: [PATCH 6/9] [dtests] Fix hipModule test for nvcc path

Change-Id: If918b87b848a825242e06b0d552a7be188a1c4b6


[ROCm/hip commit: 88abfde2f8f3a19d593b503fc76fa8a916dd3913]
---
 projects/hip/tests/src/runtimeApi/module/hipModule.cpp | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

diff --git a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
index 019645ea32..30dec3ddd8 100644
--- a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
+++ b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp
@@ -19,7 +19,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD_CMD: vcpy_kernel.code %hc --genco %S/vcpy_kernel.cpp -o vcpy_kernel.code
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
  * TEST: %t
  * HIT_END
  */
@@ -66,12 +66,11 @@ int main() {
 
     HIPCHECK(hipMalloc((void**)&Ad, SIZE));
     HIPCHECK(hipMalloc((void**)&Bd, SIZE));
-    HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
-    HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
+    HIPCHECK(hipMemcpyHtoD(Ad, A, SIZE));
+    HIPCHECK(hipMemcpyHtoD(Bd, B, SIZE));
 
     hipModule_t Module;
     hipFunction_t Function;
-    hipFunction_t f;
     HIPCHECK(hipModuleLoad(&Module, fileName));
     HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
 
@@ -92,7 +91,7 @@ int main() {
 
     HIPCHECK(hipStreamDestroy(stream));
 
-    HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
+    HIPCHECK(hipMemcpyDtoH(B, Bd, SIZE));
 
     for (uint32_t i = 0; i < LEN; i++) {
         assert(A[i] == B[i]);

From d0252dfa79cdc1127dda008aa386702fb6fecf16 Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Sun, 12 May 2019 09:54:03 -0400
Subject: [PATCH 7/9] migrate program_state logic from header into shared
 library (phase I) (#1077)

* Revert "Revert "Use COMgr to read Kernel Args Metadata (#1006)""

This reverts commit 62e96cb4cf36f43aa191eb92a3e096908eff41c9.

* Revert "Use COMgr to read Kernel Args Metadata (#1006)"

This reverts commit 882006555bcd1ce9dc7ead3ea827c13ceb596736.

* Revert "improve program state commentary"

This reverts commit fb2beb0c88265f372549450a3f0b30c6e08b03d2.

* Revert "load program state once per agent"

This reverts commit 21f5e142f5d7c6f6de86966f4399e0578848b2de.

* start moving function_names() into the hip shared lib

* start moving code_object_blobs to a new "state" object

* Consolidate various program state related static objects into a
single program_state object

* minor clean up

* move more stuffs from functional_grid_launch into program_state

* debug make_kernarg

* moving lookup for kernargs size_align into program_state

* clean up old code for kernarg size and alignment

* update hip_module to use newer api in program_state

* Create public member functions for program_state

* move most program state functions into shared library

* Pass the data buffer size to load_executable
Otherwise, it can't figure what the data size is
just from the char* (since the data is not really a string)

* turning free functions in program state into members of program_state_impl

* change the free function globals() into a member of program_state_impl

* replace the static mutex used for populating globals

* moving associate_code_object_symbols_with_host_allocation into
program_state_impl

* move load_code_object_and_freeze_executable into program_state_impl

* moving executables and functions_names into program_state_impl

* moving kernels() into program_state_impl

* moving functions() into program_state_impl

* move get_kernargs into program_state_impl

* moving kernel_descriptor into program_state_impl

* moving kernargs_size_align calculation into program_state_impl

* Changing the handle to program_state_impl to a pointer

* moving program_state_impl into a separate inline source file

* fixing/cleaning up some header file includes

* moving member function for kernargs_size_align into program_state.cpp

* moving Kernel_descriptor into program_state.inl

* add a new class to manage agent globals

* moving all agent globals processing functions into agent_globals_impl

* load program state once per agent

re-merging PR991 against other program state changes

* fix per-agent program state member initialization

* cache executables based on elf name, isa, and agent.

This avoids program state reloading executables after a shared library is dlopened.

re-merging PR1057 against other program state changes

* protect executables cache by a global mutex

* return ref to executables cache

* adapt PR#981 Make hipModuleGetGlobal be in HIP runtime


[ROCm/hip commit: f5eb91d53dd548affce17527ce350a59240b6a1c]
---
 projects/hip/CMakeLists.txt                   |   1 +
 .../hip/hcc_detail/functional_grid_launch.hpp |  65 +-
 .../include/hip/hcc_detail/hip_runtime_api.h  | 191 +++--
 .../include/hip/hcc_detail/program_state.hpp  | 625 +--------------
 projects/hip/src/hip_clang.cpp                |   4 +-
 projects/hip/src/hip_hcc.cpp                  |  10 +
 projects/hip/src/hip_module.cpp               |  21 +-
 projects/hip/src/program_state.cpp            |  63 ++
 projects/hip/src/program_state.inl            | 713 ++++++++++++++++++
 9 files changed, 961 insertions(+), 732 deletions(-)
 create mode 100644 projects/hip/src/program_state.cpp
 create mode 100644 projects/hip/src/program_state.inl

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 78d2bc7a30..30aa071515 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -232,6 +232,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
 
     set(SOURCE_FILES_RUNTIME
         src/code_object_bundle.cpp
+        src/program_state.cpp
         src/hip_clang.cpp
         src/hip_hcc.cpp
         src/hip_context.cpp
diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp
index 42326cbdc1..8082188489 100644
--- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp
+++ b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp
@@ -59,7 +59,7 @@ template <
     typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
 inline std::vector<std::uint8_t> make_kernarg(
     const std::tuple<Ts...>&,
-    const std::vector<std::pair<std::size_t, std::size_t>>&,
+    const kernargs_size_align&,
     std::vector<std::uint8_t> kernarg) {
     return kernarg;
 }
@@ -70,7 +70,7 @@ template <
     typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
 inline std::vector<std::uint8_t> make_kernarg(
     const std::tuple<Ts...>& formals,
-    const std::vector<std::pair<std::size_t, std::size_t>>& size_align,
+    const kernargs_size_align& size_align,
     std::vector<std::uint8_t> kernarg) {
     using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
 
@@ -86,13 +86,12 @@ inline std::vector<std::uint8_t> make_kernarg(
     #endif
 
     kernarg.resize(round_up_to_next_multiple_nonnegative(
-        kernarg.size(), size_align[n].second) + size_align[n].first);
+        kernarg.size(), size_align.alignment(n)) + size_align.size(n));
 
     std::memcpy(
-        kernarg.data() + kernarg.size() - size_align[n].first,
+        kernarg.data() + kernarg.size() - size_align.size(n),
         &std::get<n>(formals),
-        size_align[n].first);
-
+        size_align.size(n));
     return make_kernarg<n + 1>(formals, size_align, std::move(kernarg));
 }
 
@@ -104,45 +103,17 @@ inline std::vector<std::uint8_t> make_kernarg(
 
     if (sizeof...(Formals) == 0) return {};
 
-    auto it = function_names().find(reinterpret_cast<std::uintptr_t>(kernel));
-    if (it == function_names().cend()) {
-        hip_throw(std::runtime_error{"Undefined __global__ function."});
-    }
-
-    auto it1 = kernargs().find(it->second);
-    if (it1 == kernargs().end()) {
-        hip_throw(std::runtime_error{
-            "Missing metadata for __global__ function: " + it->second});
-    }
-
     std::tuple<Formals...> to_formals{std::move(actuals)};
     std::vector<std::uint8_t> kernarg;
     kernarg.reserve(sizeof(to_formals));
 
-    return make_kernarg<0>(to_formals, it1->second, std::move(kernarg));
+    auto& ps = hip_impl::get_program_state();
+    return make_kernarg<0>(to_formals, 
+                           ps.get_kernargs_size_align(
+                               reinterpret_cast<std::uintptr_t>(kernel)),
+                           std::move(kernarg));
 }
 
-inline
-std::string name(std::uintptr_t function_address)
-{
-    const auto it = function_names().find(function_address);
-
-    if (it == function_names().cend())  {
-        hip_throw(std::runtime_error{
-            "Invalid function passed to hipLaunchKernelGGL."});
-    }
-
-    return it->second;
-}
-
-inline
-std::string name(hsa_agent_t agent)
-{
-    char n[64]{};
-    hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, n);
-
-    return std::string{n};
-}
 
 hsa_agent_t target_agent(hipStream_t stream);
 
@@ -156,17 +127,10 @@ void hipLaunchKernelGGLImpl(
     hipStream_t stream,
     void** kernarg) {
 
-    auto agent = target_agent(stream);
-    auto it = functions(agent).find(function_address);
+    const auto& kd = hip_impl::get_program_state().kernel_descriptor(function_address, 
+                                                               target_agent(stream));
 
-    if (it == functions(agent).cend()) {
-        hip_throw(std::runtime_error{
-            "No device code available for function: " +
-            name(function_address) +
-            ", for agent: " + name(agent)});
-    }
-
-    hipModuleLaunchKernel(it->second, numBlocks.x, numBlocks.y, numBlocks.z,
+    hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z,
                           dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes,
                           stream, nullptr, kernarg);
 }
@@ -178,8 +142,7 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
                         std::uint32_t sharedMemBytes, hipStream_t stream,
                         Args... args) {
     hip_impl::hip_init();
-    auto kernarg = hip_impl::make_kernarg(
-        kernel, std::tuple<Args...>{std::move(args)...});
+    auto kernarg = hip_impl::make_kernarg(kernel, std::tuple<Args...>{std::move(args)...});
     std::size_t kernarg_size = kernarg.size();
 
     void* config[]{
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 b6eb21e0ef..d9da632577 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
@@ -67,6 +67,7 @@ THE SOFTWARE.
 #define HIP_LAUNCH_PARAM_END ((void*)0x03)
 
 #ifdef __cplusplus
+  #include <algorithm>
   #include <mutex>
   #include <string>
   #include <unordered_map>
@@ -2643,90 +2644,129 @@ std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
                                              hsa_executable_t executable);
 hsa_agent_t this_agent();
 
+
+class agent_globals_impl {
+private:
+    std::pair<
+        std::mutex,
+        std::unordered_map<
+            std::string, std::vector<Agent_global>>> globals_from_module;
+
+    std::unordered_map<
+        hsa_agent_t,
+        std::pair<
+            std::once_flag,
+            std::vector<Agent_global>>> globals_from_process;
+
+public:
+
+    hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+            hipModule_t hmod, const char* name) {
+        // the key of the map would the hash of code object associated with the
+        // hipModule_t instance
+        std::string key(hash_for(hmod));
+
+        if (globals_from_module.second.count(key) == 0) {
+            std::lock_guard<std::mutex> lck{globals_from_module.first};
+
+            if (globals_from_module.second.count(key) == 0) {
+                globals_from_module.second.emplace(
+                        key, read_agent_globals(this_agent(), executable_for(hmod)));
+            }
+        }
+
+        const auto it0 = globals_from_module.second.find(key);
+        if (it0 == globals_from_module.second.cend()) {
+            hip_throw(
+                    std::runtime_error{"agent_globals data structure corrupted."});
+        }
+
+        std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
+                it0->second.cend(), name);
+
+        // HACK for SWDEV-173477
+        //
+        // For code objects with global symbols of length 0, ROCR runtime would
+        // ignore them even though they exist in the symbol table. Therefore the
+        // result from read_agent_globals() can't be trusted entirely.
+        //
+        // As a workaround to tame applications which depend on the existence of
+        // global symbols with length 0, always return hipSuccess here.
+        //
+        // This behavior shall be reverted once ROCR runtime has been fixed to
+        // address SWDEV-173477
+
+        //return *dptr ? hipSuccess : hipErrorNotFound;
+        return hipSuccess;
+    }
+
+    hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+            const char* name) {
+
+        auto agent = this_agent();
+
+        std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
+            std::vector<Agent_global> tmp0;
+            for (auto&& executable : hip_impl::get_program_state().executables(aa)) {
+                auto tmp1 = read_agent_globals(aa, executable);
+                tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
+                            make_move_iterator(tmp1.end()));
+            }
+            globals_from_process[aa].second = move(move(tmp0));
+        }, agent);
+
+        const auto it = globals_from_process.find(agent);
+
+        if (it == globals_from_process.cend()) return hipErrorNotInitialized;
+
+        std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
+                it->second.second.cend(), name);
+
+        return *dptr ? hipSuccess : hipErrorNotFound;
+    }
+  
+};
+
+class agent_globals {
+public:
+    agent_globals() : impl(new agent_globals_impl()) { 
+        if (!impl) 
+            hip_throw(
+                std::runtime_error{"Error when constructing agent global data structures."});
+    }
+    ~agent_globals() { delete impl; }
+
+    hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+                                             hipModule_t hmod, const char* name) {
+        return impl->read_agent_global_from_module(dptr, bytes, hmod, name);
+    }
+
+    hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                                              const char* name) {
+        return impl->read_agent_global_from_process(dptr, bytes, name);
+    }
+
+private:
+    agent_globals_impl* impl;
+};
+
 inline
 __attribute__((visibility("hidden")))
-hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
-                                         hipModule_t hmod, const char* name) {
-    // the key of the map would the hash of code object associated with the
-    // hipModule_t instance
-    static std::unordered_map<
-        std::string, std::vector<Agent_global>> agent_globals;
-    std::string key(hash_for(hmod));
-
-    if (agent_globals.count(key) == 0) {
-        static std::mutex mtx;
-        std::lock_guard<std::mutex> lck{mtx};
-
-        if (agent_globals.count(key) == 0) {
-            agent_globals.emplace(
-                key, read_agent_globals(this_agent(), executable_for(hmod)));
-        }
-    }
-
-    const auto it0 = agent_globals.find(key);
-    if (it0 == agent_globals.cend()) {
-        hip_throw(
-            std::runtime_error{"agent_globals data structure corrupted."});
-    }
-
-    std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
-                                                      it0->second.cend(), name);
-
-    // HACK for SWDEV-173477
-    //
-    // For code objects with global symbols of length 0, ROCR runtime would
-    // ignore them even though they exist in the symbol table. Therefore the
-    // result from read_agent_globals() can't be trusted entirely.
-    //
-    // As a workaround to tame applications which depend on the existence of
-    // global symbols with length 0, always return hipSuccess here.
-    //
-    // This behavior shall be reverted once ROCR runtime has been fixed to
-    // address SWDEV-173477
-
-    //return *dptr ? hipSuccess : hipErrorNotFound;
-    return hipSuccess;
+agent_globals& get_agent_globals() {
+    static agent_globals ag;
+    return ag;
 }
 
+
+extern "C"
 inline
 __attribute__((visibility("hidden")))
 hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
                                           const char* name) {
-    static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
-        std::vector<Agent_global>>> globals;
-    static std::once_flag f;
-    auto agent = this_agent();
-
-    // Create placeholder for each agent in the map.
-    std::call_once(f, []() {
-        for (auto&& x : hip_impl::all_hsa_agents()) {
-            (void)globals[x];
-        }
-    });
-
-    if (globals.find(agent) == globals.cend()) {
-        hip_throw(std::runtime_error{"invalid agent"});
-    }
-
-    std::call_once(globals[agent].first, [](hsa_agent_t aa) {
-        std::vector<Agent_global> tmp0;
-        for (auto&& executable : executables(aa)) {
-            auto tmp1 = read_agent_globals(aa, executable);
-            tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
-                        make_move_iterator(tmp1.end()));
-        }
-        globals[aa].second = move(tmp0);
-    }, agent);
-
-    const auto it = globals.find(agent);
-
-    if (it == globals.cend()) return hipErrorNotInitialized;
-
-    std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
-                                                      it->second.second.cend(), name);
-
-    return *dptr ? hipSuccess : hipErrorNotFound;
+    return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
 }
+
+
 } // Namespace hip_impl.
 
 #if defined(__cplusplus)
@@ -2748,6 +2788,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
 #endif // __HIP_VDI__
 
 hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
+
 /**
  * @brief builds module from code object which resides in host memory. Image is pointer to that
  * location.
diff --git a/projects/hip/include/hip/hcc_detail/program_state.hpp b/projects/hip/include/hip/hcc_detail/program_state.hpp
index f05f41d3f5..ca8b6dcd88 100644
--- a/projects/hip/include/hip/hcc_detail/program_state.hpp
+++ b/projects/hip/include/hip/hcc_detail/program_state.hpp
@@ -22,38 +22,17 @@ THE SOFTWARE.
 
 #pragma once
 
-#include "code_object_bundle.hpp"
-#include "hsa_helpers.hpp"
-
-#if !defined(__cpp_exceptions)
-    #define try if (true)
-    #define catch(...) if (false)
-#endif
-#include "elfio/elfio.hpp"
-#if !defined(__cpp_exceptions)
-    #undef try
-    #undef catch
-#endif
-
 #include <hsa/amd_hsa_kernel_code.h>
 #include <hsa/hsa.h>
 #include <hsa/hsa_ext_amd.h>
 #include <hsa/hsa_ven_amd_loader.h>
 
-#include <link.h>
-
-#include <cassert>
 #include <cstddef>
 #include <cstdint>
 #include <cstdlib>
-#include <istream>
-#include <memory>
-#include <mutex>
-#include <sstream>
 #include <stdexcept>
 #include <string>
 #include <unordered_map>
-#include <utility>
 #include <vector>
 
 struct ihipModuleSymbol_t;
@@ -84,590 +63,48 @@ inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
 
 namespace hip_impl {
 
-std::vector<hsa_agent_t> all_hsa_agents();
-
-class Kernel_descriptor {
-    std::uint64_t kernel_object_{};
-    amd_kernel_code_t const* kernel_header_{nullptr};
-    std::string name_{};
-public:
-    Kernel_descriptor() = default;
-    Kernel_descriptor(std::uint64_t kernel_object, const std::string& name)
-        : kernel_object_{kernel_object}, name_{name}
-    {
-        bool supported{false};
-        std::uint16_t min_v{UINT16_MAX};
-        auto r = hsa_system_major_extension_supported(
-            HSA_EXTENSION_AMD_LOADER, 1, &min_v, &supported);
-
-        if (r != HSA_STATUS_SUCCESS || !supported) return;
-
-        hsa_ven_amd_loader_1_01_pfn_t tbl{};
-
-        r = hsa_system_get_major_extension_table(
-            HSA_EXTENSION_AMD_LOADER,
-            1,
-            sizeof(tbl),
-            reinterpret_cast<void*>(&tbl));
-
-        if (r != HSA_STATUS_SUCCESS) return;
-        if (!tbl.hsa_ven_amd_loader_query_host_address) return;
-
-        r = tbl.hsa_ven_amd_loader_query_host_address(
-            reinterpret_cast<void*>(kernel_object_),
-            reinterpret_cast<const void**>(&kernel_header_));
-
-        if (r != HSA_STATUS_SUCCESS) return;
-    }
-    Kernel_descriptor(const Kernel_descriptor&) = default;
-    Kernel_descriptor(Kernel_descriptor&&) = default;
-    ~Kernel_descriptor() = default;
-
-    Kernel_descriptor& operator=(const Kernel_descriptor&) = default;
-    Kernel_descriptor& operator=(Kernel_descriptor&&) = default;
-
-    operator hipFunction_t() const {  // TODO: this is awful and only meant for illustration.
-        return reinterpret_cast<hipFunction_t>(const_cast<Kernel_descriptor*>(this));
-    }
-};
-
-template<typename P>
-inline
-ELFIO::section* find_section_if(ELFIO::elfio& reader, P p) {
-    const auto it = std::find_if(
-        reader.sections.begin(), reader.sections.end(), std::move(p));
-
-    return it != reader.sections.end() ? *it : nullptr;
-}
-
-inline
-__attribute__((visibility("hidden")))
-const std::unordered_map<
-    hsa_isa_t, std::vector<std::vector<char>>>& code_object_blobs() {
-    static std::unordered_map<hsa_isa_t, std::vector<std::vector<char>>> r;
-    static std::once_flag f;
-
-    std::call_once(f, []() {
-        static std::vector<std::vector<char>> blobs{};
-
-        dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void*) {
-            ELFIO::elfio tmp;
-
-            const auto elf =
-                info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
-
-            if (!tmp.load(elf)) return 0;
-
-            const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
-                return x->get_name() == ".kernel";
-            });
-
-            if (!it) return 0;
-
-            blobs.emplace_back(it->get_data(), it->get_data() + it->get_size());
-
-            return 0;
-        }, nullptr);
-
-        for (auto&& multi_arch_blob : blobs) {
-            auto it = multi_arch_blob.begin();
-            while (it != multi_arch_blob.end()) {
-                Bundled_code_header tmp{it, multi_arch_blob.end()};
-
-                if (!valid(tmp)) break;
-
-                for (auto&& bundle : bundles(tmp)) {
-                    r[triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
-                }
-
-                it += tmp.bundled_code_size;
-            };
-        }
-    });
-
-    return r;
-}
-
-struct Symbol {
-    std::string name;
-    ELFIO::Elf64_Addr value = 0;
-    ELFIO::Elf_Xword size = 0;
-    ELFIO::Elf_Half sect_idx = 0;
-    std::uint8_t bind = 0;
-    std::uint8_t type = 0;
-    std::uint8_t other = 0;
-};
-
-inline
-Symbol read_symbol(const ELFIO::symbol_section_accessor& section,
-                   unsigned int idx) {
-    assert(idx < section.get_symbols_num());
-
-    Symbol r;
-    section.get_symbol(
-        idx, r.name, r.value, r.size, r.bind, r.type, r.sect_idx, r.other);
-
-    return r;
-}
-
-inline
-__attribute__((visibility("hidden")))
-const std::unordered_map<
-    std::string,
-    std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>& symbol_addresses() {
-    static std::unordered_map<
-        std::string, std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>> r;
-    static std::once_flag f;
-
-    std::call_once(f, []() {
-        dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void*) {
-            ELFIO::elfio tmp;
-            const auto elf =
-                info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
-
-            if (!tmp.load(elf)) return 0;
-
-            auto it = find_section_if(tmp, [](const ELFIO::section* x) {
-                return x->get_type() == SHT_SYMTAB;
-            });
-
-            if (!it) return 0;
-
-            const ELFIO::symbol_section_accessor symtab{tmp, it};
-
-            for (auto i = 0u; i != symtab.get_symbols_num(); ++i) {
-                auto s = read_symbol(symtab, i);
-
-                if (s.type != STT_OBJECT || s.sect_idx == SHN_UNDEF) continue;
-
-                const auto addr = s.value + info->dlpi_addr;
-                r.emplace(std::move(s.name), std::make_pair(addr, s.size));
-            }
-
-            return 0;
-        }, nullptr);
-    });
-
-    return r;
-}
-
-inline
-__attribute__((visibility("hidden")))
-std::unordered_map<std::string, void*>& globals() {
-    static std::unordered_map<std::string, void*> r;
-    static std::once_flag f;
-
-    std::call_once(f, []() { r.reserve(symbol_addresses().size()); });
-
-    return r;
-}
-
-inline
-std::vector<std::string> copy_names_of_undefined_symbols(
-    const ELFIO::symbol_section_accessor& section) {
-    std::vector<std::string> r;
-
-    for (auto i = 0u; i != section.get_symbols_num(); ++i) {
-        // TODO: this is boyscout code, caching the temporaries
-        //       may be of worth.
-        auto tmp = read_symbol(section, i);
-        if (tmp.sect_idx != SHN_UNDEF || tmp.name.empty()) continue;
-
-        r.push_back(std::move(tmp.name));
-    }
-
-    return r;
-}
-
 [[noreturn]]
 void hip_throw(const std::exception&);
 
-inline
-void associate_code_object_symbols_with_host_allocation(
-    const ELFIO::elfio& reader,
-    ELFIO::section* code_object_dynsym,
-    hsa_agent_t agent,
-    hsa_executable_t executable) {
-    if (!code_object_dynsym) return;
+class kernargs_size_align;
+class program_state_impl;
+class program_state {
+public:
+    program_state();
+    ~program_state();
 
-    const auto undefined_symbols = copy_names_of_undefined_symbols(
-        ELFIO::symbol_section_accessor{reader, code_object_dynsym});
+    hipFunction_t kernel_descriptor(std::uintptr_t,
+                                    hsa_agent_t);
+    
+    kernargs_size_align get_kernargs_size_align(std::uintptr_t);
+    hsa_executable_t load_executable(const char*, const size_t,
+                                     hsa_executable_t,
+                                     hsa_agent_t);
 
-    for (auto&& x : undefined_symbols) {
-        if (globals().find(x) != globals().cend()) return;
+    void* global_addr_by_name(const char* name);
 
-        const auto it1 = symbol_addresses().find(x);
+    // to fix later
+    const std::vector<hsa_executable_t>& executables(hsa_agent_t agent);
 
-        if (it1 == symbol_addresses().cend()) {
-            hip_throw(std::runtime_error{
-                "Global symbol: " + x + " is undefined."});
-        }
+    program_state(const program_state&) = delete;
 
-        static std::mutex mtx;
-        std::lock_guard<std::mutex> lck{mtx};
+private:
+    program_state_impl* impl;
+};
 
-        if (globals().find(x) != globals().cend()) return;
-
-        globals().emplace(x, (void*)(it1->second.first));
-        void* p = nullptr;
-        hsa_amd_memory_lock(
-            reinterpret_cast<void*>(it1->second.first),
-            it1->second.second,
-            nullptr,  // All agents.
-            0,
-            &p);
-
-        hsa_executable_agent_global_variable_define(
-            executable, agent, x.c_str(), p);
-    }
-}
-
-inline
-void load_code_object_and_freeze_executable(
-    const std::string& file, hsa_agent_t agent, hsa_executable_t executable) {
-    // TODO: the following sequence is inefficient, should be refactored
-    //       into a single load of the file and subsequent ELFIO
-    //       processing.
-    static const auto cor_deleter = [](hsa_code_object_reader_t* p) {
-        if (!p) return;
-
-        hsa_code_object_reader_destroy(*p);
-        delete p;
-    };
-
-    using RAII_code_reader =
-        std::unique_ptr<hsa_code_object_reader_t, decltype(cor_deleter)>;
-
-    if (file.empty()) return;
-
-    RAII_code_reader tmp{new hsa_code_object_reader_t, cor_deleter};
-    hsa_code_object_reader_create_from_memory(
-        file.data(), file.size(), tmp.get());
-
-    hsa_executable_load_agent_code_object(
-        executable, agent, *tmp, nullptr, nullptr);
-
-    hsa_executable_freeze(executable, nullptr);
-
-    static std::vector<RAII_code_reader> code_readers;
-    static std::mutex mtx;
-
-    std::lock_guard<std::mutex> lck{mtx};
-    code_readers.push_back(move(tmp));
-}
-
-inline
-hsa_executable_t load_executable(const std::string& file,
-                                 hsa_executable_t executable,
-                                 hsa_agent_t agent) {
-    ELFIO::elfio reader;
-    std::stringstream tmp{file};
-
-    if (!reader.load(tmp)) return hsa_executable_t{};
-
-    const auto code_object_dynsym = find_section_if(
-        reader, [](const ELFIO::section* x) {
-            return x->get_type() == SHT_DYNSYM;
-    });
-
-    associate_code_object_symbols_with_host_allocation(reader,
-                                                       code_object_dynsym,
-                                                       agent, executable);
-
-    load_code_object_and_freeze_executable(file, agent, executable);
-
-    return executable;
-}
+class kernargs_size_align {
+public:
+    std::size_t size(std::size_t n) const;
+    std::size_t alignment(std::size_t n) const;
+private:
+    const void* handle;
+    friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t);
+};
 
 inline
 __attribute__((visibility("hidden")))
-const std::vector<hsa_executable_t>& executables(hsa_agent_t agent) {
-    static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
-        std::vector<hsa_executable_t>>> r;
-    static std::once_flag f;
-
-    // Create placeholder for each agent in the map.
-    std::call_once(f, []() {
-        for (auto&& x : hip_impl::all_hsa_agents()) {
-            (void)r[x];
-        }
-    });
-
-    if (r.find(agent) == r.cend()) {
-        hip_throw(std::runtime_error{"invalid agent"});
-    }
-
-    std::call_once(r[agent].first, [](hsa_agent_t aa) {
-        hsa_agent_iterate_isas(aa, [](hsa_isa_t x, void* pa) {
-            const auto it = code_object_blobs().find(x);
-
-            if (it == code_object_blobs().cend()) return HSA_STATUS_SUCCESS;
-
-            hsa_agent_t a = *static_cast<hsa_agent_t*>(pa);
-
-            for (auto&& blob : it->second) {
-                hsa_executable_t tmp = {};
-
-                hsa_executable_create_alt(
-                    HSA_PROFILE_FULL,
-                    HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
-                    nullptr,
-                    &tmp);
-
-                // TODO: this is massively inefficient and only meant for
-                // illustration.
-                std::string blob_to_str{blob.cbegin(), blob.cend()};
-                tmp = load_executable(blob_to_str, tmp, a);
-
-                if (tmp.handle) r[a].second.push_back(tmp);
-            }
-
-            return HSA_STATUS_SUCCESS;
-        }, &aa);
-    }, agent);
-
-    return r[agent].second;
-}
-
-inline
-std::vector<std::pair<std::uintptr_t, std::string>> function_names_for(
-    const ELFIO::elfio& reader, ELFIO::section* symtab) {
-    std::vector<std::pair<std::uintptr_t, std::string>> r;
-    ELFIO::symbol_section_accessor symbols{reader, symtab};
-
-    for (auto i = 0u; i != symbols.get_symbols_num(); ++i) {
-        // TODO: this is boyscout code, caching the temporaries
-        //       may be of worth.
-        auto tmp = read_symbol(symbols, i);
-
-        if (tmp.type != STT_FUNC) continue;
-        if (tmp.type == SHN_UNDEF) continue;
-        if (tmp.name.empty()) continue;
-
-        r.emplace_back(tmp.value, tmp.name);
-    }
-
-    return r;
-}
-
-inline
-__attribute__((visibility("hidden")))
-const std::unordered_map<std::uintptr_t, std::string>& function_names() {
-    static std::unordered_map<std::uintptr_t, std::string> r;
-    static std::once_flag f;
-
-    std::call_once(f, []() {
-        dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void*) {
-            ELFIO::elfio tmp;
-            const auto elf =
-                info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
-
-            if (!tmp.load(elf)) return 0;
-
-            const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
-                return x->get_type() == SHT_SYMTAB;
-            });
-
-            if (!it) return 0;
-
-            auto names = function_names_for(tmp, it);
-            for (auto&& x : names) x.first += info->dlpi_addr;
-
-            r.insert(
-                std::make_move_iterator(names.begin()),
-                std::make_move_iterator(names.end()));
-
-            return 0;
-        }, nullptr);
-    });
-
-    return r;
-}
-
-inline
-__attribute__((visibility("hidden")))
-const std::unordered_map<
-    std::string, std::vector<hsa_executable_symbol_t>>& kernels(hsa_agent_t agent) {
-    static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
-        std::unordered_map<std::string, std::vector<hsa_executable_symbol_t>>>> r;
-    static std::once_flag f;
-
-    // Create placeholder for each agent in the map.
-    std::call_once(f, []() {
-        for (auto&& x : hip_impl::all_hsa_agents()) {
-            (void)r[x];
-        }
-    });
-
-    if (r.find(agent) == r.cend()) {
-        hip_throw(std::runtime_error{"invalid agent"});
-    }
-
-    std::call_once(r[agent].first, [](hsa_agent_t aa) {
-        static const auto copy_kernels = [](
-            hsa_executable_t, hsa_agent_t a, hsa_executable_symbol_t x, void*) {
-            if (type(x) == HSA_SYMBOL_KIND_KERNEL) r[a].second[name(x)].push_back(x);
-
-            return HSA_STATUS_SUCCESS;
-        };
-
-        for (auto&& executable : executables(aa)) {
-            hsa_executable_iterate_agent_symbols(
-                executable, aa, copy_kernels, nullptr);
-        }
-    }, agent);
-
-    return r[agent].second;
-}
-
-inline
-__attribute__((visibility("hidden")))
-const std::unordered_map<
-    std::uintptr_t,
-    Kernel_descriptor>& functions(hsa_agent_t agent) {
-    static std::unordered_map<hsa_agent_t, std::pair<std::once_flag,
-        std::unordered_map<std::uintptr_t, Kernel_descriptor>>> r;
-    static std::once_flag f;
-
-    // Create placeholder for each agent in the map.
-    std::call_once(f, []() {
-        for (auto&& x : hip_impl::all_hsa_agents()) {
-            (void)r[x];
-        }
-    });
-
-    if (r.find(agent) == r.cend()) {
-        hip_throw(std::runtime_error{"invalid agent"});
-    }
-
-    std::call_once(r[agent].first, [](hsa_agent_t aa) {
-        for (auto&& function : function_names()) {
-            const auto it = kernels(aa).find(function.second);
-
-            if (it == kernels(aa).cend()) continue;
-
-            for (auto&& kernel_symbol : it->second) {
-                r[aa].second.emplace(
-                    function.first,
-                    Kernel_descriptor{kernel_object(kernel_symbol), it->first});
-            }
-        }
-    }, agent);
-
-    return r[agent].second;
-}
-
-inline
-std::size_t parse_args(
-    const std::string& metadata,
-    std::size_t f,
-    std::size_t l,
-    std::vector<std::pair<std::size_t, std::size_t>>& size_align) {
-    if (f == l) return f;
-    if (!size_align.empty()) return l;
-
-    do {
-        static constexpr size_t size_sz{5};
-        f = metadata.find("Size:", f) + size_sz;
-
-        if (l <= f) return f;
-
-        auto size = std::strtoul(&metadata[f], nullptr, 10);
-
-        static constexpr size_t align_sz{6};
-        f = metadata.find("Align:", f) + align_sz;
-
-        char* l{};
-        auto align = std::strtoul(&metadata[f], &l, 10);
-
-        f += (l - &metadata[f]) + 1;
-
-        size_align.emplace_back(size, align);
-    } while (true);
-}
-
-inline
-void read_kernarg_metadata(
-    ELFIO::elfio& reader,
-    std::unordered_map<
-        std::string,
-        std::vector<std::pair<std::size_t, std::size_t>>>& kernargs) {
-    // TODO: this is inefficient.
-    auto it = find_section_if(reader, [](const ELFIO::section* x) {
-        return x->get_type() == SHT_NOTE;
-    });
-
-    if (!it) return;
-
-    const ELFIO::note_section_accessor acc{reader, it};
-    for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) {
-        ELFIO::Elf_Word type{};
-        std::string name{};
-        void* desc{};
-        ELFIO::Elf_Word desc_size{};
-
-        acc.get_note(i, type, name, desc, desc_size);
-
-        if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA.
-
-        std::string tmp{
-            static_cast<char*>(desc), static_cast<char*>(desc) + desc_size};
-
-        auto dx = tmp.find("Kernels:");
-
-        if (dx == std::string::npos) continue;
-
-        static constexpr decltype(tmp.size()) kernels_sz{8};
-        dx += kernels_sz;
-
-        do {
-            dx = tmp.find("Name:", dx);
-
-            if (dx == std::string::npos) break;
-
-            static constexpr decltype(tmp.size()) name_sz{5};
-            dx = tmp.find_first_not_of(" '", dx + name_sz);
-
-            auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx);
-            dx += fn.size();
-
-            auto dx1 = tmp.find("CodeProps", dx);
-            dx = tmp.find("Args:", dx);
-
-            if (dx1 < dx) {
-                dx = dx1;
-                continue;
-            }
-            if (dx == std::string::npos) break;
-
-            static constexpr decltype(tmp.size()) args_sz{5};
-            dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]);
-        } while (true);
-    }
-}
-
-inline
-__attribute__((visibility("hidden")))
-const std::unordered_map<
-    std::string, std::vector<std::pair<std::size_t, std::size_t>>>& kernargs() {
-    static std::unordered_map<
-        std::string, std::vector<std::pair<std::size_t, std::size_t>>> r;
-    static std::once_flag f;
-
-    std::call_once(f, []() {
-        for (auto&& isa_blobs : code_object_blobs()) {
-            for (auto&& blob : isa_blobs.second) {
-                std::stringstream tmp{std::string{blob.cbegin(), blob.cend()}};
-
-                ELFIO::elfio reader;
-
-                if (!reader.load(tmp)) continue;
-
-                read_kernarg_metadata(reader, r);
-            }
-        }
-    });
-
-    return r;
+program_state& get_program_state() {
+    static program_state ps;
+    return ps;
 }
 }  // Namespace hip_impl.
diff --git a/projects/hip/src/hip_clang.cpp b/projects/hip/src/hip_clang.cpp
index 82e181ca69..b8beb05400 100644
--- a/projects/hip/src/hip_clang.cpp
+++ b/projects/hip/src/hip_clang.cpp
@@ -85,7 +85,9 @@ __hipRegisterFatBinary(const void* data)
           reinterpret_cast<uintptr_t>(header) + desc->offset), desc->size};
       if (HIP_DUMP_CODE_OBJECT)
         __hipDumpCodeObject(image);
-      module->executable = hip_impl::load_executable(image, module->executable, agent);
+      module->executable = hip_impl::get_program_state().load_executable(image.data(), image.size(),
+                                                                         module->executable,
+                                                                         agent);
 
       if (module->executable.handle) {
         modules->at(deviceId) = module;
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 8011160cc1..2af83b65a8 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -2489,4 +2489,14 @@ namespace hip_impl {
             std::terminate();
         #endif
     }
+
+    std::mutex executables_cache_mutex;
+
+    std::vector<hsa_executable_t>& executables_cache(
+            std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
+        static std::unordered_map<std::string,
+            std::unordered_map<hsa_isa_t,
+                std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
+        return cache[elf][isa][agent];
+    }
 } // Namespace hip_impl.
diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp
index e1a3b1ff65..e029ec6e78 100644
--- a/projects/hip/src/hip_module.cpp
+++ b/projects/hip/src/hip_module.cpp
@@ -25,6 +25,7 @@ THE SOFTWARE.
 #include "hip/hcc_detail/hsa_helpers.hpp"
 #include "hip/hcc_detail/program_state.hpp"
 #include "hip_hcc_internal.h"
+#include "program_state.inl"
 #include "trace_helper.h"
 
 #include <hsa/amd_hsa_kernel_code.h>
@@ -289,7 +290,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
 
     if (!name) return hipErrorNotInitialized;
 
-    return hip_impl::read_agent_global_from_module(dptr, bytes, hmod, name);
+    return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
 }
 
 namespace hip_impl {
@@ -512,11 +513,8 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
     if (!func) return hipErrorInvalidDeviceFunction;
 
     auto agent = this_agent();
-    const auto it = functions(agent).find(reinterpret_cast<uintptr_t>(func));
-
-    if (it == functions(agent).cend()) return hipErrorInvalidDeviceFunction;
-
-    const auto header = static_cast<hipFunction_t>(it->second)->_header;
+    auto kd = get_program_state().kernel_descriptor(reinterpret_cast<uintptr_t>(func), agent);
+    const auto header = kd->_header;
 
     if (!header) throw runtime_error{"Ill-formed Kernel_descriptor."};
 
@@ -548,7 +546,8 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
 
     auto content = tmp.empty() ? read_elf_file_as_string(image) : tmp;
 
-    (*module)->executable = load_executable(content, (*module)->executable,
+    (*module)->executable = get_program_state().load_executable(
+                                            content.data(), content.size(), (*module)->executable,
                                             this_agent());
 
     // compute the hash of the code object
@@ -591,10 +590,10 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
     if (!texRef) return ihipLogStatus(hipErrorInvalidValue);
 
     if (!hmod || !name) return ihipLogStatus(hipErrorNotInitialized);
+    
+    auto addr = get_program_state().global_addr_by_name(name);
+    if (addr == nullptr) return ihipLogStatus(hipErrorInvalidValue);
 
-    const auto it = globals().find(name);
-    if (it == globals().end()) return ihipLogStatus(hipErrorInvalidValue);
-
-    *texRef = reinterpret_cast<textureReference*>(it->second);
+    *texRef = reinterpret_cast<textureReference*>(addr);
     return ihipLogStatus(hipSuccess);
 }
diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp
new file mode 100644
index 0000000000..6783f85b9b
--- /dev/null
+++ b/projects/hip/src/program_state.cpp
@@ -0,0 +1,63 @@
+#include "../include/hip/hcc_detail/program_state.hpp"
+
+#include <hsa/hsa.h>
+
+#include <cstdint>
+#include <stdexcept>
+#include <unordered_map>
+#include <vector>
+
+// contains implementation of program_state_impl
+#include "program_state.inl"
+
+namespace hip_impl {
+
+std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
+    return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
+}
+
+std::size_t kernargs_size_align::alignment(std::size_t n) const{
+    return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].second;
+}
+
+program_state::program_state() :
+    impl(new program_state_impl) {
+    if (!impl) hip_throw(std::runtime_error {
+                         "Unknown error when constructing program state."});
+}
+
+program_state::~program_state() {
+    delete(impl);
+}
+
+void* program_state::global_addr_by_name(const char* name) {
+    const auto it = impl->get_globals().find(name);
+    if (it == impl->get_globals().end())
+      return nullptr;
+    else
+      return it->second;
+}
+
+hsa_executable_t program_state::load_executable(const char* data,
+                                                const size_t data_size,
+                                                hsa_executable_t executable,
+                                                hsa_agent_t agent) {
+    return impl->load_executable(data, data_size, executable, agent);
+}
+
+const std::vector<hsa_executable_t>& program_state::executables(hsa_agent_t agent) {
+    return impl->get_executables(agent);
+}
+
+hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
+                                                          hsa_agent_t agent) {
+    auto& kd = impl->kernel_descriptor(function_address, agent);
+    return kd;
+}
+
+kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
+  kernargs_size_align t;
+  t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
+  return t;
+}
+};
diff --git a/projects/hip/src/program_state.inl b/projects/hip/src/program_state.inl
new file mode 100644
index 0000000000..9729da8115
--- /dev/null
+++ b/projects/hip/src/program_state.inl
@@ -0,0 +1,713 @@
+#include "../include/hip/hcc_detail/program_state.hpp"
+
+#include "../include/hip/hcc_detail/code_object_bundle.hpp"
+#include "../include/hip/hcc_detail/hsa_helpers.hpp"
+
+#if !defined(__cpp_exceptions)
+    #define try if (true)
+    #define catch(...) if (false)
+#endif
+#include "../include/hip/hcc_detail/elfio/elfio.hpp"
+#if !defined(__cpp_exceptions)
+    #undef try
+    #undef catch
+#endif
+
+#include <hsa/amd_hsa_kernel_code.h>
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+#include <hsa/hsa_ven_amd_loader.h>
+
+#include <link.h>
+
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <cstdio>
+#include <memory>
+#include <mutex>
+#include <string>
+#include <sstream>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+namespace hip_impl {
+
+[[noreturn]]
+void hip_throw(const std::exception&);
+
+std::vector<hsa_agent_t> all_hsa_agents();
+
+extern std::mutex executables_cache_mutex;
+
+std::vector<hsa_executable_t>& executables_cache(std::string, hsa_isa_t, hsa_agent_t);
+
+template<typename P>
+inline
+ELFIO::section* find_section_if(ELFIO::elfio& reader, P p) {
+    const auto it = std::find_if(
+        reader.sections.begin(), reader.sections.end(), std::move(p));
+
+    return it != reader.sections.end() ? *it : nullptr;
+}
+
+struct Symbol {
+    std::string name;
+    ELFIO::Elf64_Addr value = 0;
+    ELFIO::Elf_Xword size = 0;
+    ELFIO::Elf_Half sect_idx = 0;
+    std::uint8_t bind = 0;
+    std::uint8_t type = 0;
+    std::uint8_t other = 0;
+};
+
+class Kernel_descriptor {
+    std::uint64_t kernel_object_{};
+    amd_kernel_code_t const* kernel_header_{nullptr};
+    std::string name_{};
+public:
+    Kernel_descriptor() = default;
+    Kernel_descriptor(std::uint64_t kernel_object, const std::string& name)
+        : kernel_object_{kernel_object}, name_{name}
+    {
+        bool supported{false};
+        std::uint16_t min_v{UINT16_MAX};
+        auto r = hsa_system_major_extension_supported(
+            HSA_EXTENSION_AMD_LOADER, 1, &min_v, &supported);
+
+        if (r != HSA_STATUS_SUCCESS || !supported) return;
+
+        hsa_ven_amd_loader_1_01_pfn_t tbl{};
+
+        r = hsa_system_get_major_extension_table(
+            HSA_EXTENSION_AMD_LOADER,
+            1,
+            sizeof(tbl),
+            reinterpret_cast<void*>(&tbl));
+
+        if (r != HSA_STATUS_SUCCESS) return;
+        if (!tbl.hsa_ven_amd_loader_query_host_address) return;
+
+        r = tbl.hsa_ven_amd_loader_query_host_address(
+            reinterpret_cast<void*>(kernel_object_),
+            reinterpret_cast<const void**>(&kernel_header_));
+
+        if (r != HSA_STATUS_SUCCESS) return;
+    }
+    Kernel_descriptor(const Kernel_descriptor&) = default;
+    Kernel_descriptor(Kernel_descriptor&&) = default;
+    ~Kernel_descriptor() = default;
+
+    Kernel_descriptor& operator=(const Kernel_descriptor&) = default;
+    Kernel_descriptor& operator=(Kernel_descriptor&&) = default;
+
+    operator hipFunction_t() const {  // TODO: this is awful and only meant for illustration.
+        return reinterpret_cast<hipFunction_t>(const_cast<Kernel_descriptor*>(this));
+    }
+};
+
+class program_state_impl {
+
+public:
+
+    std::pair<
+        std::once_flag,
+        std::unordered_map<
+            std::string,
+                std::unordered_map<
+                    hsa_isa_t,
+                    std::vector<std::vector<char>>>>> code_object_blobs;
+
+    std::pair<
+        std::once_flag,
+        std::unordered_map<
+            std::string, 
+            std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>> symbol_addresses;
+
+    std::unordered_map<
+        hsa_agent_t,
+        std::pair<
+            std::once_flag,
+            std::vector<hsa_executable_t>>> executables;
+
+    std::unordered_map<
+        hsa_agent_t,
+        std::pair<
+            std::once_flag,
+            std::unordered_map<
+                std::string,
+                std::vector<hsa_executable_symbol_t>>>> kernels;
+
+    std::pair<
+        std::once_flag,
+        std::unordered_map<
+            std::string, std::vector<std::pair<std::size_t, std::size_t>>>> kernargs;
+
+    std::pair<
+        std::once_flag,
+        std::unordered_map<std::uintptr_t, std::string>> function_names;
+
+    std::unordered_map<
+        hsa_agent_t,
+        std::pair<
+            std::once_flag,
+            std::unordered_map<
+                std::uintptr_t,
+                Kernel_descriptor>>> functions;
+      
+    std::tuple<
+        std::once_flag,
+        std::mutex,
+        std::unordered_map<std::string, void*>> globals;
+
+    using RAII_code_reader =
+        std::unique_ptr<hsa_code_object_reader_t, 
+                        std::function<void(hsa_code_object_reader_t*)>>;
+    std::pair<
+        std::mutex,
+        std::vector<RAII_code_reader>> code_readers;
+
+    program_state_impl() {
+        // Create placeholder for each agent for the per-agent members.
+        for (auto&& x : hip_impl::all_hsa_agents()) {
+            (void)executables[x];
+            (void)kernels[x];
+            (void)functions[x];
+        }
+    }
+
+    const std::unordered_map<
+        std::string,
+            std::unordered_map<
+                hsa_isa_t,
+                std::vector<std::vector<char>>>>& get_code_object_blobs() {
+
+        std::call_once(code_object_blobs.first, [this]() {
+            dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* p) {
+                ELFIO::elfio tmp;
+
+                const auto elf =
+                    info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
+
+                if (!tmp.load(elf)) return 0;
+
+                const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
+                    return x->get_name() == ".kernel";
+                });
+
+                if (!it) return 0;
+
+                auto& impl = *static_cast<program_state_impl*>(p);
+
+                std::vector<char> multi_arch_blob(it->get_data(), it->get_data() + it->get_size());
+                auto blob_it = multi_arch_blob.begin();
+                while (blob_it != multi_arch_blob.end()) {
+                    Bundled_code_header tmp{blob_it, multi_arch_blob.end()};
+
+                    if (!valid(tmp)) break;
+
+                    for (auto&& bundle : bundles(tmp)) {
+                        impl.code_object_blobs.second[elf][triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
+                    }
+
+                    blob_it += tmp.bundled_code_size;
+                };
+
+                return 0;
+            }, this);
+        });
+
+        return code_object_blobs.second;
+    }
+
+    Symbol read_symbol(const ELFIO::symbol_section_accessor& section,
+                   unsigned int idx) {
+        assert(idx < section.get_symbols_num());
+
+        Symbol r;
+        section.get_symbol(
+            idx, r.name, r.value, r.size, r.bind, r.type, r.sect_idx, r.other);
+
+        return r;
+    }
+
+    const std::unordered_map<
+        std::string,
+        std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>& get_symbol_addresses() {
+
+        std::call_once(symbol_addresses.first, [this]() {
+            dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* psi_ptr) {
+
+                if (!psi_ptr)
+                    return 0;
+
+                program_state_impl* t = static_cast<program_state_impl*>(psi_ptr);
+
+                ELFIO::elfio tmp;
+                const auto elf =
+                    info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
+
+                if (!tmp.load(elf)) return 0;
+
+                auto it = find_section_if(tmp, [](const ELFIO::section* x) { 
+                    return x->get_type() == SHT_SYMTAB;
+                });
+
+                if (!it) return 0;
+
+                const ELFIO::symbol_section_accessor symtab{tmp, it};
+
+                for (auto i = 0u; i != symtab.get_symbols_num(); ++i) {
+                    auto s = t->read_symbol(symtab, i);
+
+                    if (s.type != STT_OBJECT || s.sect_idx == SHN_UNDEF) continue;
+
+                    const auto addr = s.value + info->dlpi_addr;
+                    t->symbol_addresses.second.emplace(std::move(s.name), std::make_pair(addr, s.size));
+                }
+
+                return 0;
+            }, this);
+        });
+
+        return symbol_addresses.second;
+    }
+
+    std::unordered_map<std::string, void*>& get_globals() {
+        std::call_once(std::get<0>(globals), [this]() { 
+            std::get<2>(globals).reserve(get_symbol_addresses().size()); 
+        });
+        return std::get<2>(globals);
+    }
+
+    std::mutex& get_globals_mutex() {
+        return std::get<1>(globals);
+    }
+
+    std::vector<std::string> copy_names_of_undefined_symbols(
+        const ELFIO::symbol_section_accessor& section) {
+        std::vector<std::string> r;
+
+        for (auto i = 0u; i != section.get_symbols_num(); ++i) {
+            // TODO: this is boyscout code, caching the temporaries
+            //       may be of worth.
+            auto tmp = read_symbol(section, i);
+            if (tmp.sect_idx != SHN_UNDEF || tmp.name.empty()) continue;
+
+            r.push_back(std::move(tmp.name));
+        }
+
+        return r;
+    }
+
+    void associate_code_object_symbols_with_host_allocation(
+        const ELFIO::elfio& reader,
+        ELFIO::section* code_object_dynsym,
+        hsa_agent_t agent,
+        hsa_executable_t executable) {
+        if (!code_object_dynsym) return;
+
+        const auto undefined_symbols = copy_names_of_undefined_symbols(
+            ELFIO::symbol_section_accessor{reader, code_object_dynsym});
+
+        auto& g = get_globals();
+        auto& g_mutex = get_globals_mutex();
+        for (auto&& x : undefined_symbols) {
+
+            if (g.find(x) != g.cend()) return;
+
+            const auto it1 = get_symbol_addresses().find(x);
+
+            if (it1 == get_symbol_addresses().cend()) {
+                hip_throw(std::runtime_error{
+                    "Global symbol: " + x + " is undefined."});
+            }
+
+            std::lock_guard<std::mutex> lck{g_mutex};
+
+            if (g.find(x) != g.cend()) return;
+
+            g.emplace(x, (void*)(it1->second.first));
+            void* p = nullptr;
+            hsa_amd_memory_lock(
+                reinterpret_cast<void*>(it1->second.first),
+                it1->second.second,
+                nullptr,  // All agents.
+                0,
+                &p);
+
+            hsa_executable_agent_global_variable_define(
+                executable, agent, x.c_str(), p);
+        }
+    }
+
+    void load_code_object_and_freeze_executable(
+        const std::string& file, hsa_agent_t agent, hsa_executable_t executable) {
+        // TODO: the following sequence is inefficient, should be refactored
+        //       into a single load of the file and subsequent ELFIO
+        //       processing.
+        if (file.empty()) return;
+
+        static const auto cor_deleter = [] (hsa_code_object_reader_t* p) {
+            if (!p) return;
+            hsa_code_object_reader_destroy(*p);
+            delete p;
+        };
+
+        RAII_code_reader tmp{new hsa_code_object_reader_t, cor_deleter};
+        hsa_code_object_reader_create_from_memory(
+            file.data(), file.size(), tmp.get());
+
+        hsa_executable_load_agent_code_object(
+            executable, agent, *tmp, nullptr, nullptr);
+
+        hsa_executable_freeze(executable, nullptr);
+
+        std::lock_guard<std::mutex> lck{code_readers.first};
+        code_readers.second.push_back(move(tmp));
+    }
+
+
+    const std::vector<hsa_executable_t>& get_executables(hsa_agent_t agent) {
+
+        if (executables.find(agent) == executables.cend()) {
+            hip_throw(std::runtime_error{"invalid agent"});
+        }
+
+        std::call_once(executables[agent].first, [this](hsa_agent_t aa) {
+            auto data = std::make_pair(this, &aa);
+            hsa_agent_iterate_isas(aa, [](hsa_isa_t x, void* d) {
+                auto& p = *static_cast<decltype(data)*>(d);
+                auto& impl = *(p.first);
+                for (const auto code_object_it : impl.get_code_object_blobs()) {
+                    const auto elf = code_object_it.first;
+                    const auto code_object_blobs = code_object_it.second;
+                    const auto it = code_object_blobs.find(x);
+
+                    if (it == code_object_blobs.cend()) continue;
+
+                    hsa_agent_t a = *static_cast<hsa_agent_t*>(p.second);
+
+                    std::lock_guard<std::mutex> lck{executables_cache_mutex};
+
+                    std::vector<hsa_executable_t>& current_exes =
+                            hip_impl::executables_cache(elf, x, a);
+                    // check the cache for already loaded executables
+                    if (current_exes.empty()) {
+                        // executables do not yet exist for this elf+isa+agent, create and cache them
+                        for (auto&& blob : it->second) {
+                            hsa_executable_t tmp = {};
+
+                            hsa_executable_create_alt(
+                                HSA_PROFILE_FULL,
+                                HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
+                                nullptr,
+                                &tmp);
+
+                            // TODO: this is massively inefficient and only meant for
+                            // illustration.
+                            tmp = impl.load_executable(blob.data(), blob.size(), tmp, a);
+
+                            if (tmp.handle) current_exes.push_back(tmp);
+                        }
+                    }
+                    // append cached executables to our agent's vector of executables
+                    impl.executables[a].second.insert(impl.executables[a].second.end(),
+                            current_exes.begin(), current_exes.end());
+                }
+                return HSA_STATUS_SUCCESS;
+            }, &data);
+        }, agent);
+
+        return executables[agent].second;
+    }
+    
+    hsa_executable_t load_executable(const char* data,
+                                     const size_t data_size,
+                                     hsa_executable_t executable,
+                                     hsa_agent_t agent) {
+        ELFIO::elfio reader;
+        std::string ts = std::string(data, data_size);
+        std::stringstream tmp{ts};
+
+        if (!reader.load(tmp)) return hsa_executable_t{};
+        const auto code_object_dynsym = find_section_if(
+            reader, [](const ELFIO::section* x) {
+                return x->get_type() == SHT_DYNSYM;
+        });
+
+        associate_code_object_symbols_with_host_allocation(reader,
+                                                           code_object_dynsym,
+                                                           agent, executable);
+
+        load_code_object_and_freeze_executable(ts, agent, executable);
+
+        return executable;
+    }
+
+    std::vector<std::pair<std::uintptr_t, std::string>> function_names_for(
+        const ELFIO::elfio& reader, ELFIO::section* symtab) {
+        std::vector<std::pair<std::uintptr_t, std::string>> r;
+        ELFIO::symbol_section_accessor symbols{reader, symtab};
+
+        for (auto i = 0u; i != symbols.get_symbols_num(); ++i) {
+            // TODO: this is boyscout code, caching the temporaries
+            //       may be of worth.
+            auto tmp = read_symbol(symbols, i);
+
+            if (tmp.type != STT_FUNC) continue;
+            if (tmp.type == SHN_UNDEF) continue;
+            if (tmp.name.empty()) continue;
+
+            r.emplace_back(tmp.value, tmp.name);
+        }
+
+        return r;
+    }
+
+    const std::unordered_map<std::uintptr_t, std::string>& get_function_names() {
+
+        std::call_once(function_names.first, [this]() {
+            dl_iterate_phdr([](dl_phdr_info* info, std::size_t, void* p) {
+                ELFIO::elfio tmp;
+                const auto elf =
+                    info->dlpi_addr ? info->dlpi_name : "/proc/self/exe";
+
+                if (!tmp.load(elf)) return 0;
+
+                const auto it = find_section_if(tmp, [](const ELFIO::section* x) {
+                    return x->get_type() == SHT_SYMTAB;
+                });
+
+                if (!it) return 0;
+
+                auto& impl = *static_cast<program_state_impl*>(p);
+       
+                auto names = impl.function_names_for(tmp, it);
+                for (auto&& x : names) x.first += info->dlpi_addr;
+
+                impl.function_names.second.insert(
+                    std::make_move_iterator(names.begin()),
+                    std::make_move_iterator(names.end()));
+
+                return 0;
+            }, this);
+        });
+
+        return function_names.second;
+    }
+
+    const std::unordered_map<
+        std::string, std::vector<hsa_executable_symbol_t>>& get_kernels(hsa_agent_t agent) {
+
+        if (kernels.find(agent) == kernels.cend()) {
+            hip_throw(std::runtime_error{"invalid agent"});
+        }
+
+        std::call_once(kernels[agent].first, [this](hsa_agent_t aa) {
+            static const auto copy_kernels = [](
+                hsa_executable_t, hsa_agent_t a, hsa_executable_symbol_t x, void* p) {
+                auto& impl = *static_cast<program_state_impl*>(p);
+                if (type(x) == HSA_SYMBOL_KIND_KERNEL) impl.kernels[a].second[hip_impl::name(x)].push_back(x);
+
+                return HSA_STATUS_SUCCESS;
+            };
+
+            for (auto&& executable : get_executables(aa)) {
+                hsa_executable_iterate_agent_symbols(
+                    executable, aa, copy_kernels, this);
+            }
+        }, agent);
+
+        return kernels[agent].second;
+    }
+
+    const std::unordered_map<
+        std::uintptr_t,
+        Kernel_descriptor>& get_functions(hsa_agent_t agent) {
+
+        if (functions.find(agent) == functions.cend()) {
+            hip_throw(std::runtime_error{"invalid agent"});
+        }
+
+        std::call_once(functions[agent].first, [this](hsa_agent_t aa) {
+            for (auto&& function : get_function_names()) {
+                const auto it = get_kernels(aa).find(function.second);
+
+                if (it == get_kernels(aa).cend()) continue;
+
+                for (auto&& kernel_symbol : it->second) {
+                    functions[aa].second.emplace(
+                        function.first,
+                        Kernel_descriptor{kernel_object(kernel_symbol), it->first});
+                }
+            }
+        }, agent);
+
+        return functions[agent].second;
+    }
+
+    std::size_t parse_args(
+            const std::string& metadata,
+            std::size_t f,
+            std::size_t l,
+            std::vector<std::pair<std::size_t, std::size_t>>& size_align) {
+        if (f == l) return f;
+        if (!size_align.empty()) return l;
+
+        do {
+            static constexpr size_t size_sz{5};
+            f = metadata.find("Size:", f) + size_sz;
+
+            if (l <= f) return f;
+
+            auto size = std::strtoul(&metadata[f], nullptr, 10);
+
+            static constexpr size_t align_sz{6};
+            f = metadata.find("Align:", f) + align_sz;
+
+            char* l{};
+            auto align = std::strtoul(&metadata[f], &l, 10);
+
+            f += (l - &metadata[f]) + 1;
+
+            size_align.emplace_back(size, align);
+        } while (true);
+    }
+
+    void read_kernarg_metadata(
+            ELFIO::elfio& reader,
+            std::unordered_map<
+            std::string,
+            std::vector<std::pair<std::size_t, std::size_t>>>& kernargs) {
+        // TODO: this is inefficient.
+        auto it = find_section_if(reader, [](const ELFIO::section* x) {
+                return x->get_type() == SHT_NOTE;
+                });
+
+        if (!it) return;
+
+        const ELFIO::note_section_accessor acc{reader, it};
+        for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) {
+            ELFIO::Elf_Word type{};
+            std::string name{};
+            void* desc{};
+            ELFIO::Elf_Word desc_size{};
+
+            acc.get_note(i, type, name, desc, desc_size);
+
+            if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA.
+
+            std::string tmp{
+                static_cast<char*>(desc), static_cast<char*>(desc) + desc_size};
+
+            auto dx = tmp.find("Kernels:");
+
+            if (dx == std::string::npos) continue;
+
+            static constexpr decltype(tmp.size()) kernels_sz{8};
+            dx += kernels_sz;
+
+            do {
+                dx = tmp.find("Name:", dx);
+
+                if (dx == std::string::npos) break;
+
+                static constexpr decltype(tmp.size()) name_sz{5};
+                dx = tmp.find_first_not_of(" '", dx + name_sz);
+
+                auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx);
+                dx += fn.size();
+
+                auto dx1 = tmp.find("CodeProps", dx);
+                dx = tmp.find("Args:", dx);
+
+                if (dx1 < dx) {
+                    dx = dx1;
+                    continue;
+                }
+                if (dx == std::string::npos) break;
+
+                static constexpr decltype(tmp.size()) args_sz{5};
+                dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]);
+            } while (true);
+        }
+    }
+
+    const std::unordered_map<std::string, 
+        std::vector<std::pair<std::size_t, std::size_t>>>& get_kernargs() {
+
+        std::call_once(kernargs.first, [this]() {
+            for (auto&& name_and_isa_blobs : get_code_object_blobs()) {
+                for (auto&& isa_blobs : name_and_isa_blobs.second) {
+                    for (auto&& blob : isa_blobs.second) {
+                        std::stringstream tmp{std::string{blob.cbegin(), blob.cend()}};
+
+                        ELFIO::elfio reader;
+
+                        if (!reader.load(tmp)) continue;
+
+                        read_kernarg_metadata(reader, kernargs.second);
+                    }
+                }
+            }
+        });
+
+        return kernargs.second;
+    }
+
+    std::string name(std::uintptr_t function_address)
+    {
+        const auto it = get_function_names().find(function_address);
+
+        if (it == get_function_names().cend())  {
+            hip_throw(std::runtime_error{
+                    "Invalid function passed to hipLaunchKernelGGL."});
+        }
+
+        return it->second;
+    }
+
+    std::string name(hsa_agent_t agent)
+    {
+        char n[64]{};
+        hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, n);
+
+        return std::string{n};
+    }
+
+    const Kernel_descriptor& kernel_descriptor(std::uintptr_t function_address,
+            hsa_agent_t agent) {
+
+        auto it0 = get_functions(agent).find(function_address);
+
+        if (it0 == get_functions(agent).cend()) {
+            hip_throw(std::runtime_error{
+                    "No device code available for function: " +
+                    std::string(name(function_address)) +
+                    ", for agent: " + name(agent)});
+        }
+
+        return it0->second;
+    }
+
+    const std::vector<std::pair<std::size_t, std::size_t>>& 
+        kernargs_size_align(std::uintptr_t kernel) {
+
+        auto it = get_function_names().find(kernel);
+        if (it == get_function_names().cend()) {
+            hip_throw(std::runtime_error{"Undefined __global__ function."});
+        }
+
+        auto it1 = get_kernargs().find(it->second);
+        if (it1 == get_kernargs().end()) {
+            hip_throw(std::runtime_error{
+                      "Missing metadata for __global__ function: " + it->second});
+        }
+
+        return it1->second;
+    }
+};  // class program_state_impl
+
+};

From 376e19308b4f307edc4e590351e764fcb4630668 Mon Sep 17 00:00:00 2001
From: Nick Curtis <arghdos@users.noreply.github.com>
Date: Sun, 12 May 2019 08:57:44 -0500
Subject: [PATCH 8/9] Markdown fixes & Whitespace cleanup for samples (#1096)

* Fix multiline code blocks in README's

* Whitespace cleanup


[ROCm/hip commit: 5257b54a3914737f1141f91d0d10e4f8fb69cb35]
---
 .../2_Cookbook/0_MatrixTranspose/Readme.md    | 48 ++++++++++---------
 .../2_Cookbook/10_inline_asm/Readme.md        |  6 +--
 .../samples/2_Cookbook/1_hipEvent/Readme.md   | 26 ++++++----
 .../2_Cookbook/3_shared_memory/Readme.md      |  4 +-
 .../hip/samples/2_Cookbook/4_shfl/Readme.md   | 22 +++++----
 .../hip/samples/2_Cookbook/5_2dshfl/Readme.md | 28 ++++++-----
 .../2_Cookbook/6_dynamic_shared/Readme.md     |  8 ++--
 .../samples/2_Cookbook/7_streams/Readme.md    | 18 ++++---
 .../hip/samples/2_Cookbook/9_unroll/Readme.md | 20 ++++----
 9 files changed, 103 insertions(+), 77 deletions(-)

diff --git a/projects/hip/samples/2_Cookbook/0_MatrixTranspose/Readme.md b/projects/hip/samples/2_Cookbook/0_MatrixTranspose/Readme.md
index 4a52b862a0..8f19613f87 100644
--- a/projects/hip/samples/2_Cookbook/0_MatrixTranspose/Readme.md
+++ b/projects/hip/samples/2_Cookbook/0_MatrixTranspose/Readme.md
@@ -7,30 +7,32 @@ This tutorial shows how to get write simple HIP application. We will write the s
 HIP is a C++ runtime API and kernel language that allows developers to create portable applications that can run on AMD and other GPU’s. Our goal was to rise above the lowest-common-denominator paths and deliver a solution that allows you, the developer, to use essential hardware features and maximize your application’s performance on GPU hardware.
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 Here is simple example showing how to write your first program in HIP.
-In order to use the HIP framework, we need to add the "hip_runtime.h" header file. SInce its c++ api you can add any header file you have been using earlier while writing your c/c++ program. For gpgpu programming, we have host(microprocessor) and the device(gpu). 
+In order to use the HIP framework, we need to add the "hip_runtime.h" header file. SInce its c++ api you can add any header file you have been using earlier while writing your c/c++ program. For gpgpu programming, we have host(microprocessor) and the device(gpu).
 
 ## Device-side code
 We will work on device side code first, Here is simple example showing a snippet of HIP device side code:
 
-`__global__ void matrixTranspose(float *out,                  `
-`                                float *in,                   `
-`                                const int width,             `
-`                                const int height)            `
-`{                                                            `
-`    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;  `
-`    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;  `
-`                                                             `
-`    out[y * width + x] = in[x * height + y];                 `
-`}                                                            `
+```
+__global__ void matrixTranspose(float *out,
+                                float *in,
+                                const int width,
+                                const int height)
+{
+    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+    out[y * width + x] = in[x * height + y];
+}
+```
 
 `__global__` keyword is the Function-Type Qualifiers, it is used with functions that are executed on device and are called/launched from the hosts.
 other function-type qualifiers are:
@@ -40,11 +42,11 @@ other function-type qualifiers are:
 `__host__` can combine with `__device__`, in which case the function compiles for both the host and device. These functions cannot use the HIP grid coordinate functions (for example, "hipThreadIdx_x", will talk about it latter). A possible workaround is to pass the necessary coordinate info as an argument to the function.
 `__host__` cannot combine with `__global__`.
 
-`__global__` functions are often referred to as *kernels*, and calling one is termed *launching the kernel*. 
+`__global__` functions are often referred to as *kernels*, and calling one is termed *launching the kernel*.
 
 Next keyword is `void`. HIP `__global__` functions must have a `void` return type. Global functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute.
 
-The kernel function begins with 
+The kernel function begins with
 `    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;`
 `    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;`
 here the keyword hipBlockIdx_x, hipBlockIdx_y and hipBlockIdx_z(not used here) are the built-in functions to identify the threads in a block. The keyword hipBlockDim_x, hipBlockDim_y and hipBlockDim_z(not used here) are to identify the dimensions of the block.
@@ -60,18 +62,20 @@ We allocated memory to the Matrix on host side by using malloc and initiallized
 here the first parameter is the destination pointer, second is the source pointer, third is the size of memory copy and the last specify the direction on memory copy(which is in this case froom host to device). While in order to transfer memory from device to host, use `hipMemcpyDeviceToHost` and for device to device memory copy use `hipMemcpyDeviceToDevice`.
 
 Now, we'll see how to launch the kernel.
-`  hipLaunchKernelGGL(matrixTranspose,                                                      `
-`                  dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),             `
-`                  dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),                          `
-`                  0, 0,                                                                    `
-`                  gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);                          `
+```
+  hipLaunchKernelGGL(matrixTranspose,
+                  dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
+                  dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                  0, 0,
+                  gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
+```
 
 HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel (this convention replaces the `Cuda <<< >>>` syntax). In HIP,
 - Kernels launch with the `"hipLaunchKernelGGL"` function
 - The first five parameters to hipLaunchKernelGGL are the following:
    - **symbol kernelName**: the name of the kernel to launch.  To support template kernels which contains "," use the HIP_KERNEL_NAME macro. In current application it's "matrixTranspose".
-   - **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. In MatrixTranspose sample, it's "dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y)". 
-   - **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block.In MatrixTranspose sample, it's "dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y)". 
+   - **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. In MatrixTranspose sample, it's "dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y)".
+   - **dim3 blockDim**: 3D-block dimensions specifying the number of threads in each block.In MatrixTranspose sample, it's "dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y)".
    - **size_t dynamicShared**: amount of additional shared memory to allocate when launching the kernel. In MatrixTranspose sample, it's '0'.
    - **hipStream_t**: stream where the kernel should execute. A value of 0 corresponds to the NULL stream.In MatrixTranspose sample, it's '0'.
 - Kernel arguments follow these first five parameters. Here, these are "gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT".
diff --git a/projects/hip/samples/2_Cookbook/10_inline_asm/Readme.md b/projects/hip/samples/2_Cookbook/10_inline_asm/Readme.md
index 7d0301bc74..a6cafac683 100644
--- a/projects/hip/samples/2_Cookbook/10_inline_asm/Readme.md
+++ b/projects/hip/samples/2_Cookbook/10_inline_asm/Readme.md
@@ -6,7 +6,7 @@ This tutorial is about how to use inline GCN asm in kernel. In this tutorial, we
 
 If you want to take advantage of the extra performance benefits of writing in assembly as well as take advantage of special GPU hardware features that were only available through assemby, then this tutorial is for you. In this tutorial we'll be explaining how to start writing inline asm in kernel.
 
-For more insight Please read the following blogs by Ben Sander 
+For more insight Please read the following blogs by Ben Sander
 [The Art of AMDGCN Assembly: How to Bend the Machine to Your Will](gpuopen.com/amdgcn-assembly)
 [AMD GCN Assembly: Cross-Lane Operations](http://gpuopen.com/amd-gcn-assembly-cross-lane-operations/)
 
@@ -15,13 +15,13 @@ For more information:
 [User Guide for AMDGPU Back-end](llvm.org/docs/AMDGPUUsage.html)
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 We will be using the Simple Matrix Transpose application from the our very first tutorial.
 
diff --git a/projects/hip/samples/2_Cookbook/1_hipEvent/Readme.md b/projects/hip/samples/2_Cookbook/1_hipEvent/Readme.md
index 7f276f628f..1af4efb06f 100644
--- a/projects/hip/samples/2_Cookbook/1_hipEvent/Readme.md
+++ b/projects/hip/samples/2_Cookbook/1_hipEvent/Readme.md
@@ -7,13 +7,13 @@ This tutorial is follow-up of the previous one where we learn how to write our f
 Memory transfer and kernel execution are the most important parameter in parallel computing (specially HPC and machine learning). Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore obtaining the memory transfer timing and kernel execution timing plays key role in application optimization.
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to get the performance score for memory transfer and kernel execution time.
 
@@ -21,12 +21,16 @@ We will be using the Simple Matrix Transpose application from the previous tutor
 
 We'll learn how to use the event management functionality of HIP runtime api. In the same sourcecode, we used for MatrixTranspose we will declare the following events as follows:
 
-`  hipEvent_t start, stop;`
+```
+   hipEvent_t start, stop;
+```
 
 We'll create the event with the help of following code:
 
-`  hipEventCreate(&start);`
-`  hipEventCreate(&stop);`
+```
+   hipEventCreate(&start);
+   hipEventCreate(&stop);
+```
 
 We'll use the "eventMs" variable to store the time taken value:
 `  float eventMs = 1.0f;`
@@ -41,11 +45,13 @@ Now, we'll have the operation for which we need to compute the time taken. For t
 `  hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);`
 
 and for kernel execution time we'll use `hipKernelLaunch`:
-`  hipLaunchKernelGGL(matrixTranspose,                                                    `
-`                  dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),        `
-`                  dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),                     `
-`                  0, 0,                                                               `
-`                  gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);                     `
+```
+hipLaunchKernelGGL(matrixTranspose,
+                   dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
+                   dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                   0, 0,
+                   gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
+```
 
 Now to mark the end of the eventRecord, we will again use the hipEventRecord by passing the stop event:
 `  hipEventRecord(stop, NULL);`
diff --git a/projects/hip/samples/2_Cookbook/3_shared_memory/Readme.md b/projects/hip/samples/2_Cookbook/3_shared_memory/Readme.md
index 8b9e102ec9..0cce65206d 100644
--- a/projects/hip/samples/2_Cookbook/3_shared_memory/Readme.md
+++ b/projects/hip/samples/2_Cookbook/3_shared_memory/Readme.md
@@ -7,13 +7,13 @@ Earlier we learned how to write our first hip program, in which we compute Matri
 As we mentioned earlier  that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use static shared memory and will explain the dynamic one latter.
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
 
diff --git a/projects/hip/samples/2_Cookbook/4_shfl/Readme.md b/projects/hip/samples/2_Cookbook/4_shfl/Readme.md
index 923d2f3837..b64899a6cb 100644
--- a/projects/hip/samples/2_Cookbook/4_shfl/Readme.md
+++ b/projects/hip/samples/2_Cookbook/4_shfl/Readme.md
@@ -5,23 +5,25 @@ In this tutorial, we'll explain how to use the warp shfl operations to improve t
 ## Introduction:
 
 Let's talk about Warp first. The kernel code is executed in groups of fixed number of threads known as Warp. For nvidia WarpSize is 32 while for AMD, 32 for Polaris architecture and 64 for rest. Threads in a warp are referred to as lanes and are numbered from 0 to warpSize -1. With the help of shfl ops, we can directly exchange values of variable between threads without using any memory ops within a warp. There are four types of shfl ops:
-`  int   __shfl      (int var,   int srcLane, int width=warpSize);              `
-`  float __shfl      (float var, int srcLane, int width=warpSize);              `
-`  int   __shfl_up   (int var,   unsigned int delta, int width=warpSize);       `
-`  float __shfl_up   (float var, unsigned int delta, int width=warpSize);       `
-`  int   __shfl_down (int var,   unsigned int delta, int width=warpSize);       `
-`  float __shfl_down (float var, unsigned int delta, int width=warpSize);       `
-`  int   __shfl_xor  (int var,   int laneMask, int width=warpSize)              `
-`  float __shfl_xor  (float var, int laneMask, int width=warpSize);             `
+```
+   int   __shfl      (int var,   int srcLane, int width=warpSize);
+   float __shfl      (float var, int srcLane, int width=warpSize);
+   int   __shfl_up   (int var,   unsigned int delta, int width=warpSize);
+   float __shfl_up   (float var, unsigned int delta, int width=warpSize);
+   int   __shfl_down (int var,   unsigned int delta, int width=warpSize);
+   float __shfl_down (float var, unsigned int delta, int width=warpSize);
+   int   __shfl_xor  (int var,   int laneMask, int width=warpSize)
+   float __shfl_xor  (float var, int laneMask, int width=warpSize);
+```
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
 
diff --git a/projects/hip/samples/2_Cookbook/5_2dshfl/Readme.md b/projects/hip/samples/2_Cookbook/5_2dshfl/Readme.md
index 8efff49d8e..d6ba29c2e6 100644
--- a/projects/hip/samples/2_Cookbook/5_2dshfl/Readme.md
+++ b/projects/hip/samples/2_Cookbook/5_2dshfl/Readme.md
@@ -5,31 +5,35 @@ This tutorial is follow-up of the previous tutorial, where we learned how to use
 ## Introduction:
 
 Let's talk about Warp first. The kernel code is executed in groups of fixed number of threads known as Warp. For nvidia WarpSize is 32 while for AMD, 32 for Polaris architecture and 64 for rest. Threads in a warp are referred to as lanes and are numbered from 0 to warpSize -1. With the help of shfl ops, we can directly exchange values of variable between threads without using any memory ops within a warp. There are four types of shfl ops:
-`  int   __shfl      (int var,   int srcLane, int width=warpSize);              `
-`  float __shfl      (float var, int srcLane, int width=warpSize);              `
-`  int   __shfl_up   (int var,   unsigned int delta, int width=warpSize);       `
-`  float __shfl_up   (float var, unsigned int delta, int width=warpSize);       `
-`  int   __shfl_down (int var,   unsigned int delta, int width=warpSize);       `
-`  float __shfl_down (float var, unsigned int delta, int width=warpSize);       `
-`  int   __shfl_xor  (int var,   int laneMask, int width=warpSize)              `
-`  float __shfl_xor  (float var, int laneMask, int width=warpSize);             `
+```
+   int   __shfl      (int var,   int srcLane, int width=warpSize);
+   float __shfl      (float var, int srcLane, int width=warpSize);
+   int   __shfl_up   (int var,   unsigned int delta, int width=warpSize);
+   float __shfl_up   (float var, unsigned int delta, int width=warpSize);
+   int   __shfl_down (int var,   unsigned int delta, int width=warpSize);
+   float __shfl_down (float var, unsigned int delta, int width=warpSize);
+   int   __shfl_xor  (int var,   int laneMask, int width=warpSize);
+   float __shfl_xor  (float var, int laneMask, int width=warpSize);
+```
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
 
 ## __shfl ops in 2D
 
 In the same sourcecode, we used for MatrixTranspose. We'll add the following:
-`  int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;                     `
-`  out[x*width + y] = __shfl(val,y*width + x);                                 `
+```
+	int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+	out[x*width + y] = __shfl(val,y*width + x);
+```
 
 With the help of this application, we can say that kernel code can be converted into  multi-dimensional threads with ease.
 
diff --git a/projects/hip/samples/2_Cookbook/6_dynamic_shared/Readme.md b/projects/hip/samples/2_Cookbook/6_dynamic_shared/Readme.md
index 9622ac85fe..2b9a44d0ba 100644
--- a/projects/hip/samples/2_Cookbook/6_dynamic_shared/Readme.md
+++ b/projects/hip/samples/2_Cookbook/6_dynamic_shared/Readme.md
@@ -7,13 +7,13 @@ Earlier we learned how to use static shared memory. In this tutorial, we'll expl
 As we mentioned earlier  that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use dynamic shared memory.
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to use shared memory.
 
@@ -25,11 +25,13 @@ Shared memory is way more faster than that of global and constant memory and acc
 here the first parameter is the data type while the second one is the variable name.
 
 The other important change is:
-`  hipLaunchKernelGGL(matrixTranspose,                                                   `
+```
+    hipLaunchKernelGGL(matrixTranspose,
                   dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
                   dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                   sizeof(float)*WIDTH*WIDTH, 0,
                   gpuTransposeMatrix , gpuMatrix, WIDTH);
+```
 here we replaced 4th parameter with amount of additional shared memory to allocate when launching the kernel.
 
 ## How to build and run:
diff --git a/projects/hip/samples/2_Cookbook/7_streams/Readme.md b/projects/hip/samples/2_Cookbook/7_streams/Readme.md
index 456deada1f..c05f9e7707 100644
--- a/projects/hip/samples/2_Cookbook/7_streams/Readme.md
+++ b/projects/hip/samples/2_Cookbook/7_streams/Readme.md
@@ -7,13 +7,13 @@ In all Earlier tutorial we used single stream, In this tutorial, we'll explain h
 The various instances of kernel to be executed on device in exact launch order defined by Host are called streams. We can launch multiple streams on a single device. We will learn how to learn two streams which can we scaled with ease.
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 We will be using the Simple Matrix Transpose application from the previous tutorial and modify it to learn how to launch multiple streams.
 
@@ -23,22 +23,28 @@ In this tutorial, we'll use both instances of shared memory (i.e., static and dy
 `  hipStream_t streams[num_streams];                                             `
 
 and create stream using `hipStreamCreate` as follows:
-`  for(int i=0;i<num_streams;i++)                                                `
-`  hipStreamCreate(&streams[i]);                                                 `
+```
+for(int i=0;i<num_streams;i++)
+    hipStreamCreate(&streams[i]);
+```
 
 and while kernel launch, we make the following changes in 5th parameter to hipLaunchKernelGGL(having 0 as the default stream value):
 
-`  hipLaunchKernelGGL(matrixTranspose_static_shared,                             `
+```
+ hipLaunchKernelGGL(matrixTranspose_static_shared,
                     dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
                     dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                     0, streams[0],
                     gpuTransposeMatrix[0], data[0], width);
+```
 
-`  hipLaunchKernelGGL(matrixTranspose_dynamic_shared,                            `
+```
+ hipLaunchKernelGGL(matrixTranspose_dynamic_shared,
                     dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
                     dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                     sizeof(float)*WIDTH*WIDTH, streams[1],
                     gpuTransposeMatrix[1], data[1], width);
+```
 
 here we replaced 4th parameter with amount of additional shared memory to allocate when launching the kernel.
 
diff --git a/projects/hip/samples/2_Cookbook/9_unroll/Readme.md b/projects/hip/samples/2_Cookbook/9_unroll/Readme.md
index 194eb9c7ff..9b464f1472 100644
--- a/projects/hip/samples/2_Cookbook/9_unroll/Readme.md
+++ b/projects/hip/samples/2_Cookbook/9_unroll/Readme.md
@@ -1,6 +1,6 @@
 ## Using Pragma unroll ###
 
-In this tutorial, we'll explain how to use #pragma unroll to improve the performance. 
+In this tutorial, we'll explain how to use #pragma unroll to improve the performance.
 
 ## Introduction:
 
@@ -8,24 +8,26 @@ Loop unrolling optimization hints can be specified with #pragma unroll and #prag
 Specifying #pragma unroll without a parameter directs the loop unroller to attempt to fully unroll the loop if the trip count is known at compile time and attempt to partially unroll the loop if the trip count is not known at compile time.
 
 ## Requirement:
-For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) 
+For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md)
 
 ## prerequiste knowledge:
 
 Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming.
 
-## Simple Matrix Transpose 
+## Simple Matrix Transpose
 
 For this tutorial we will be using MatrixTranspose with shfl operation i.e., our 4_shfl tutorial since it is the only examples where we used loops inside the kernel.
 
 In this tutorial, we'll use `#pragma unroll`. In the same sourcecode, we used for MatrixTranspose. We'll add it just before the for loop as following:
 
-`#pragma unroll                                                        `
-`    for(int i=0;i<width;i++)                                          `
-`    {                                                                 `
-`        for(int j=0;j<width;j++)                                      `
-`            out[i*width + j] = __shfl(val,j*width + i);               `
-`    }                                                                 `
+```
+#pragma unroll
+	for(int i=0;i<width;i++)
+	{
+		for(int j=0;j<width;j++)
+			out[i*width + j] = __shfl(val,j*width + i);
+	}
+```
 
 Specifying the optional parameter, #pragma unroll value, directs the unroller to unroll the loop value times. Be careful while using it.
 Specifying #pragma nounroll indicates that the loop should not be unroll. #pragma unroll 1 will show the same behaviour.

From c4567ad01a223df65e77aaaa1efc8cf5aff5b9b3 Mon Sep 17 00:00:00 2001
From: Rahul Garg <garg.rahul@gmail.com>
Date: Sun, 12 May 2019 23:18:26 -0700
Subject: [PATCH 9/9] Add fine grained host memory lock support (#1095)

* Add fine grained host memory lock support

* Fix default flag check


[ROCm/hip commit: aeeab1b23fe4f8b5a5e7c0893e97a9b4ccbc3327]
---
 .../hip/include/hip/hcc_detail/hip_runtime_api.h   |  2 +-
 projects/hip/src/hip_memory.cpp                    | 14 ++++++++++++--
 2 files changed, 13 insertions(+), 3 deletions(-)

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 d9da632577..3eb59e485e 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
@@ -191,7 +191,7 @@ enum hipLimit_t {
     0x2  ///< Map the allocation into the address space for the current device.  The device pointer
          ///< can be obtained with #hipHostGetDevicePointer.
 #define hipHostRegisterIoMemory 0x4  ///< Not supported.
-
+#define hipExtHostRegisterCoarseGrained 0x8  ///< Coarse Grained host memory lock
 
 #define hipDeviceScheduleAuto 0x0  ///< Automatically select between Spin and Yield
 #define hipDeviceScheduleSpin                                                                      \
diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp
index b8632970d4..0fad8ab890 100644
--- a/projects/hip/src/hip_memory.cpp
+++ b/projects/hip/src/hip_memory.cpp
@@ -957,15 +957,25 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
         }
         // TODO-test : multi-gpu access to registered host memory.
         if (ctx) {
-            if (flags == hipHostRegisterDefault || flags == hipHostRegisterPortable ||
-                flags == hipHostRegisterMapped) {
+            if ((flags == hipHostRegisterDefault) || (flags & hipHostRegisterPortable) ||
+                (flags & hipHostRegisterMapped) || (flags == hipExtHostRegisterCoarseGrained)) {
                 auto device = ctx->getWriteableDevice();
                 std::vector<hc::accelerator> vecAcc;
                 for (int i = 0; i < g_deviceCnt; i++) {
                     vecAcc.push_back(ihipGetDevice(i)->_acc);
                 }
+#if (__hcc_workweek__ >= 19183)
+                if(flags & hipExtHostRegisterCoarseGrained) {
+                    am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0],
+                                                    vecAcc.size());
+                } else {
+                    am_status = hc::am_memory_host_lock_with_flag(device->_acc, hostPtr, sizeBytes, &vecAcc[0],
+                                                    vecAcc.size());
+                }
+#else
                 am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0],
                                                     vecAcc.size());
+#endif
                 if ( am_status == AM_SUCCESS ) {
                      am_status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);