From 9014c89f049fc8cda9369d7a3bc2c14fef6f53b5 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Mon, 13 Nov 2017 17:20:07 +0000
Subject: [PATCH 1/9] Add hipify mappings for all CUDA headers that have HIP
 equivalents

I'm particularly running into issues with `device_types.h` in real
CUDA code...
---
 hipamd/hipify-clang/src/CUDA2HipMap.cpp | 13 ++++++++++---
 1 file changed, 10 insertions(+), 3 deletions(-)

diff --git a/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/hipamd/hipify-clang/src/CUDA2HipMap.cpp
index de6ddb2d74..b2e5251139 100644
--- a/hipamd/hipify-clang/src/CUDA2HipMap.cpp
+++ b/hipamd/hipify-clang/src/CUDA2HipMap.cpp
@@ -312,9 +312,16 @@ const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
 /// Maps cuda header names to hip header names.
 const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
     // CUDA includes
-    {"cuda.h",             {"hip/hip_runtime.h",     CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER}},
-    {"cuda_runtime.h",     {"hip/hip_runtime.h",     CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME}},
-    {"cuda_runtime_api.h", {"hip/hip_runtime_api.h", CONV_INCLUDE,             API_RUNTIME}},
+    {"cuda.h",               {"hip/hip_runtime.h",        CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER}},
+    {"cuda_runtime.h",       {"hip/hip_runtime.h",        CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME}},
+    {"cuda_runtime_api.h",   {"hip/hip_runtime_api.h",    CONV_INCLUDE,             API_RUNTIME}},
+    {"channel_descriptor.h", {"hip/channel_descriptor.h", CONV_INCLUDE,             API_RUNTIME}},
+    {"device_functions.h",   {"hip/device_functions.h",   CONV_INCLUDE,             API_RUNTIME}},
+    {"driver_types.h",       {"hip/driver_types.h",       CONV_INCLUDE,             API_RUNTIME}},
+    {"cuComplex.h",          {"hip/hip_complex.h",        CONV_INCLUDE,             API_RUNTIME}},
+    {"cuda_fp16.h",          {"hip/hip_fp16.h",           CONV_INCLUDE,             API_RUNTIME}},
+    {"cuda_texture_types.h", {"hip/hip_texture_types.h",  CONV_INCLUDE,             API_RUNTIME}},
+    {"vector_types.h",       {"hip/hip_vector_types.h",   CONV_INCLUDE,             API_RUNTIME}},
 
     // CUBLAS includes
     {"cublas.h",    {"hipblas.h", CONV_INCLUDE, API_BLAS}},

From e60a95d7ddc1fdf7a36fff7a9f3403758c26f4b9 Mon Sep 17 00:00:00 2001
From: Pierre <zins.pierre@gmail.com>
Date: Mon, 13 Nov 2017 16:13:19 -0500
Subject: [PATCH 2/9] Fix missing MARKER_END

Logging status of hipCtxSynchronize was missing
Test if hip profiling is active for MARKER_END in ihipPostLaunchKernel
Add MARKER_END after the completion of a kernel launched through
the "grid launch"
---
 hipamd/src/grid_launch.cpp | 3 +++
 hipamd/src/hip_context.cpp | 2 +-
 hipamd/src/hip_hcc.cpp     | 4 +++-
 3 files changed, 7 insertions(+), 2 deletions(-)

diff --git a/hipamd/src/grid_launch.cpp b/hipamd/src/grid_launch.cpp
index fd5c2a1573..9b3cf509c5 100644
--- a/hipamd/src/grid_launch.cpp
+++ b/hipamd/src/grid_launch.cpp
@@ -92,5 +92,8 @@ namespace hip_impl
 
         delete static_cast<L*>(locked_stream);
         locked_stream = nullptr;
+        if(HIP_PROFILE_API) {
+            MARKER_END();
+        }
     }
 }
diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp
index 69d75e7f31..d3d0691e55 100644
--- a/hipamd/src/hip_context.cpp
+++ b/hipamd/src/hip_context.cpp
@@ -269,7 +269,7 @@ hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig )
 hipError_t hipCtxSynchronize ( void )
 {
     HIP_INIT_API(1);
-    return ihipSynchronize(); //TODP Shall check validity of ctx?
+    return ihipLogStatus(ihipSynchronize()); //TODP Shall check validity of ctx?
 }
 
 hipError_t hipCtxGetFlags ( unsigned int* flags )
diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp
index 2d2d2745fe..2eff57f18b 100644
--- a/hipamd/src/hip_hcc.cpp
+++ b/hipamd/src/hip_hcc.cpp
@@ -1592,7 +1592,9 @@ void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launc
     tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
 
     stream->lockclose_postKernelCommand(kernelName, lp.av);
-    MARKER_END();
+    if(HIP_PROFILE_API) {
+        MARKER_END();
+    }
 }
 
 //=================================================================================================

From 9fa3ed4f6f7082143b91fb01170e1e69dd797ddc Mon Sep 17 00:00:00 2001
From: Phaneendr-kumar Lanka <planka@hcBLR-Titan.amd.com>
Date: Tue, 14 Nov 2017 16:37:59 +0530
Subject: [PATCH 3/9] [nvccTests]Enabled inline_asm_vadd on nvcc

---
 hipamd/tests/src/kernel/inline_asm_vadd.cpp | 8 +++++++-
 1 file changed, 7 insertions(+), 1 deletion(-)

diff --git a/hipamd/tests/src/kernel/inline_asm_vadd.cpp b/hipamd/tests/src/kernel/inline_asm_vadd.cpp
index 7a941d31af..7adca312a9 100644
--- a/hipamd/tests/src/kernel/inline_asm_vadd.cpp
+++ b/hipamd/tests/src/kernel/inline_asm_vadd.cpp
@@ -39,7 +39,13 @@ __global__ void vadd_asm(hipLaunchParm lp,
 {
     int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
 
-    asm volatile ("v_add_f32_e32 %0, %1, %2" : "=v" (out[i]) : "v"(in[i]),"v" (out[i]));
+#ifdef __HIP_PLATFORM_NVCC__
+        asm volatile("add.f32 %0,%1,%2;":"=f"(out[i]):"f"(in[i]),"f"(out[i]));
+#endif
+
+#ifdef __HIP_PLATFORM_HCC__
+   asm volatile ("v_add_f32_e32 %0, %1, %2" : "=v" (out[i]) : "v"(in[i]),"v" (out[i]));
+#endif
 }
 
 // CPU implementation of Vector Result

From 20947f80fba711bc051742a06960131d46653628 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Wed, 15 Nov 2017 18:52:59 +0530
Subject: [PATCH 4/9] Fixed half2 issue on gfx900

---
 hipamd/bin/hipcc                           |  1 +
 hipamd/tests/src/deviceLib/hipTestHalf.cpp | 33 +++++++++++++---------
 2 files changed, 21 insertions(+), 13 deletions(-)

diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc
index bbe079db41..b991bd40aa 100755
--- a/hipamd/bin/hipcc
+++ b/hipamd/bin/hipcc
@@ -439,6 +439,7 @@ if($HIP_PLATFORM eq "hcc"){
     if ($target_gfx900 eq 1) {
         $HIPLDFLAGS .= " --amdgpu-target=gfx900";
         $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX900__=1 ";
+        $ENV{HCC_EXTRA_LIBRARIES_GFX900}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
     }
 }
 
diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 94a3882eea..873dd23805 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -19,13 +19,14 @@ THE SOFTWARE.
 
 #include <iostream>
 #include <hip/hip_fp16.h>
-#include "hip/hip_runtime_api.h"
+#include "hip/hip_runtime.h"
+#include "test_common.h"
 
 #define LEN 64
 #define HALF_SIZE 64*sizeof(__half)
 #define HALF2_SIZE 64*sizeof(__half2)
 
-#if __HIP_ARCH_GFX803__ > 0
+#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX803__
 
 __global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
   int tx = hipThreadIdx_x;
@@ -61,15 +62,21 @@ __global__ void __half2Math(hipLaunchParm lp, __half2 *A, __half2 *B, __half2 *C
 #endif
 
 int main(){
-  __half *A, *B, *C;
-  hipMalloc(&A, HALF_SIZE);
-  hipMalloc(&B, HALF_SIZE);
-  hipMalloc(&C, HALF_SIZE);
-  hipLaunchKernel(__halfMath, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A, B, C);
-  __half2 *A2, *B2, *C2;
-  hipMalloc(&A, HALF2_SIZE);
-  hipMalloc(&B, HALF2_SIZE);
-  hipMalloc(&C, HALF2_SIZE);
-  hipLaunchKernel(__half2Math, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A2, B2, C2);
-
+    __half *A, *B, *C;
+    hipMalloc(&A, HALF_SIZE);
+    hipMalloc(&B, HALF_SIZE);
+    hipMalloc(&C, HALF_SIZE);
+    hipLaunchKernel(__halfMath, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A, B, C);
+    hipFree(A);
+    hipFree(B);
+    hipFree(C);
+    __half2 *A2, *B2, *C2;
+    hipMalloc(&A2, HALF2_SIZE);
+    hipMalloc(&B2, HALF2_SIZE);
+    hipMalloc(&C2, HALF2_SIZE);
+    hipLaunchKernel(__half2Math, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A2, B2, C2);
+    hipFree(A2);
+    hipFree(B2);
+    hipFree(C2);
+    passed();
 }

From cfc9f8746564a70628007ad90ab37d63f6e2f7d7 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Thu, 16 Nov 2017 09:34:52 +0530
Subject: [PATCH 5/9] Fixed test case for GFX900

---
 hipamd/tests/src/deviceLib/hipTestHalf.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 873dd23805..f56538f080 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -26,7 +26,7 @@ THE SOFTWARE.
 #define HALF_SIZE 64*sizeof(__half)
 #define HALF2_SIZE 64*sizeof(__half2)
 
-#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX803__
+#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__
 
 __global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
   int tx = hipThreadIdx_x;

From 0cc921f103c9cdbcebe03898d3bee9b5411e8e80 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Sun, 19 Nov 2017 01:54:12 +0000
Subject: [PATCH 6/9] This implements the trivial change needed to move back
 from the hip{Something}_{x, y, z} macros to the natural CUDA syntax of
 Something.{x, y, z}. This is contained in lines 384-404 in hip_runtime.h. All
 of the other changes have to do with changing unit tests to use this syntax.
 The macros are retained for backwards compatibility.

---
 hipamd/include/hip/hcc_detail/hip_runtime.h   | 21 ++++++
 hipamd/src/device_util.cpp                    | 12 ++--
 .../device/hipFuncDeviceSynchronize.cpp       |  2 +-
 hipamd/tests/src/deviceLib/hipComplex.cpp     |  2 +-
 .../tests/src/deviceLib/hipDeviceMemcpy.cpp   |  6 +-
 hipamd/tests/src/deviceLib/hipFloatMath.cpp   |  2 +-
 .../src/deviceLib/hipSimpleAtomicsTest.cpp    |  4 +-
 hipamd/tests/src/deviceLib/hipTestDevice.cpp  | 32 ++++-----
 .../src/deviceLib/hipTestDeviceDouble.cpp     | 28 ++++----
 .../src/deviceLib/hipTestDeviceSymbol.cpp     |  2 +-
 hipamd/tests/src/deviceLib/hipTestHalf.cpp    |  4 +-
 hipamd/tests/src/deviceLib/hipThreadFence.cpp |  2 +-
 hipamd/tests/src/deviceLib/hip_anyall.cpp     |  8 +--
 hipamd/tests/src/deviceLib/hip_ballot.cpp     |  8 +--
 hipamd/tests/src/deviceLib/hip_brev.cpp       |  4 +-
 hipamd/tests/src/deviceLib/hip_clz.cpp        |  4 +-
 hipamd/tests/src/deviceLib/hip_ffs.cpp        |  4 +-
 hipamd/tests/src/deviceLib/hip_popc.cpp       |  4 +-
 hipamd/tests/src/deviceLib/hip_test_ldg.cpp   |  4 +-
 .../src/deviceLib/hip_test_make_type.cpp      | 70 +++++++++----------
 hipamd/tests/src/deviceLib/hip_trig.cpp       |  2 +-
 hipamd/tests/src/experimental/xcompile/hHip.c |  2 +-
 .../src/experimental/xcompile/hipxxKer.cpp    |  2 +-
 .../src/experimental/xcompile/hxxHip.cpp      |  2 +-
 hipamd/tests/src/hipC.c                       |  2 +-
 hipamd/tests/src/hipC.cpp                     |  2 +-
 hipamd/tests/src/hipCKernel.c                 |  2 +-
 hipamd/tests/src/kernel/hipDynamicShared.cpp  |  4 +-
 hipamd/tests/src/kernel/hipDynamicShared2.cpp |  2 +-
 hipamd/tests/src/kernel/hipGridLaunch.cpp     |  7 +-
 .../src/kernel/hipLanguageExtensions.cpp      |  8 +--
 hipamd/tests/src/kernel/hipTestConstant.cpp   |  4 +-
 .../tests/src/kernel/hipTestMallocKernel.cpp  |  4 +-
 hipamd/tests/src/kernel/hipTestMemKernel.cpp  | 20 +++---
 hipamd/tests/src/kernel/inline_asm_vadd.cpp   |  4 +-
 hipamd/tests/src/kernel/inline_asm_vmac.cpp   |  4 +-
 hipamd/tests/src/kernel/launch_bounds.cpp     |  2 +-
 .../device/hipDeviceSynchronize.cpp           |  2 +-
 .../src/runtimeApi/memory/hipHostGetFlags.cpp |  2 +-
 .../src/runtimeApi/memory/hipHostMalloc.cpp   | 26 +++----
 .../src/runtimeApi/memory/hipHostRegister.cpp | 12 ++--
 .../src/runtimeApi/memory/hipMemcpyAsync.cpp  | 24 +++----
 .../memory/hipMemoryAllocateCoherent.cpp      |  2 +-
 .../runtimeApi/memory/p2p_copy_coherency.cpp  | 12 ++--
 .../tests/src/runtimeApi/module/hipModule.cpp |  2 +-
 .../src/runtimeApi/module/vcpy_kernel.cpp     |  2 +-
 .../multiThread/hipMultiThreadStreams2.cpp    |  6 +-
 .../runtimeApi/stream/hipAPIStreamDisable.cpp |  4 +-
 .../runtimeApi/stream/hipAPIStreamEnable.cpp  |  4 +-
 .../src/runtimeApi/stream/hipNullStream.cpp   |  4 +-
 .../tests/src/runtimeApi/stream/hipStream.h   |  8 +--
 .../synchronization/copy_coherency.cpp        | 40 +++++------
 .../synchronization/memcpyInt.device.cpp      |  6 +-
 hipamd/tests/src/stress/hipStressAsync.cpp    |  4 +-
 hipamd/tests/src/stress/hipStressChain.cpp    |  2 +-
 hipamd/tests/src/stress/hipStressKernel.cpp   |  2 +-
 hipamd/tests/src/stress/hipStressSync.cpp     |  4 +-
 hipamd/tests/src/test_common.h                | 46 ++++++------
 hipamd/tests/src/texture/hipTextureObj2D.cpp  |  4 +-
 hipamd/tests/src/texture/hipTextureRef2D.cpp  |  4 +-
 60 files changed, 270 insertions(+), 248 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h
index d3211ed3f5..924e774af0 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -381,6 +381,27 @@ __device__ void  __threadfence_system(void) ;
  * @}
  */
 
+template<typename std::common_type<
+    decltype(hc_get_group_id),
+    decltype(hc_get_group_size),
+    decltype(hc_get_num_groups),
+    decltype(hc_get_workitem_id)>::type f>
+class Coordinates {
+    using R = decltype(f(0));
+
+    struct X { __device__ operator R() const { return f(0); } };
+    struct Y { __device__ operator R() const { return f(1); } };
+    struct Z { __device__ operator R() const { return f(2); } };
+public:
+    static constexpr X x{};
+    static constexpr Y y{};
+    static constexpr Z z{};
+};
+
+static constexpr Coordinates<hc_get_group_size> blockDim;
+static constexpr Coordinates<hc_get_group_id> blockIdx;
+static constexpr Coordinates<hc_get_num_groups> gridDim;
+static constexpr Coordinates<hc_get_workitem_id> threadIdx;
 
 #define hipThreadIdx_x (hc_get_workitem_id(0))
 #define hipThreadIdx_y (hc_get_workitem_id(1))
diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp
index 6afc797ec6..367a4c1a4f 100644
--- a/hipamd/src/device_util.cpp
+++ b/hipamd/src/device_util.cpp
@@ -45,8 +45,8 @@ __device__ void *__hip_hc_malloc(size_t size)
     {
         return (void*)nullptr;
     }
-    uint32_t totalThreads = hipBlockDim_x * hipGridDim_x * hipBlockDim_y * hipGridDim_y * hipBlockDim_z * hipGridDim_z;
-    uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x;
+    uint32_t totalThreads = blockDim.x * gridDim.x * blockDim.y * gridDim.y * blockDim.z * gridDim.z;
+    uint32_t currentWorkItem = threadIdx.x + blockDim.x * blockIdx.x;
 
     uint32_t numHeapsPerWorkItem = NUM_PAGES / totalThreads;
     uint32_t heapSizePerWorkItem = SIZE_OF_HEAP / totalThreads;
@@ -932,7 +932,7 @@ __device__  unsigned long long int atomicMax(unsigned long long int* address,
 template<typename T>
 __device__ T atomicCAS_impl(T* address, T compare, T val)
 {
-  // the implementation assumes the atomic is lock-free and 
+  // the implementation assumes the atomic is lock-free and
   // has the same size as the non-atmoic equivalent type
   static_assert(sizeof(T) == sizeof(std::atomic<T>)
                 , "size mismatch between atomic and non-atomic types");
@@ -945,7 +945,7 @@ __device__ T atomicCAS_impl(T* address, T compare, T val)
 
   T expected = compare;
 
-  // hcc should generate a system scope atomic CAS 
+  // hcc should generate a system scope atomic CAS
   std::atomic_compare_exchange_weak_explicit(u.atomic_address
                                             , &expected, val
                                             , std::memory_order_acq_rel
@@ -1110,8 +1110,8 @@ __device__ void* __get_dynamicgroupbaseptr() {
   return hc::get_dynamic_group_segment_base_pointer();
 }
 
-__host__ void* __get_dynamicgroupbaseptr() { 
-  return nullptr; 
+__host__ void* __get_dynamicgroupbaseptr() {
+  return nullptr;
 }
 
 // Precise Math Functions
diff --git a/hipamd/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp b/hipamd/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp
index dac56bf709..c8c2e644c3 100644
--- a/hipamd/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp
+++ b/hipamd/tests/src/Functional/device/hipFuncDeviceSynchronize.cpp
@@ -34,7 +34,7 @@ THE SOFTWARE.
 #define NUM_STREAMS 2
 
 __global__ void Iter(hipLaunchParm lp, int *Ad, int num){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     // Kernel loop designed to execute very slowly... ... ...   so we can test timing-related behavior below
     if(tx == 0){
         for(int i = 0; i<num;i++){
diff --git a/hipamd/tests/src/deviceLib/hipComplex.cpp b/hipamd/tests/src/deviceLib/hipComplex.cpp
index 8a153b6bf0..c01b430339 100644
--- a/hipamd/tests/src/deviceLib/hipComplex.cpp
+++ b/hipamd/tests/src/deviceLib/hipComplex.cpp
@@ -27,7 +27,7 @@ THE SOFTWARE.
 #define SIZE 64<<2
 
 __global__  void getSqAbs(hipLaunchParm lp, float *A, float *B, float *C){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     C[tx] = hipCsqabsf(make_hipFloatComplex(A[tx], B[tx]));
 }
 
diff --git a/hipamd/tests/src/deviceLib/hipDeviceMemcpy.cpp b/hipamd/tests/src/deviceLib/hipDeviceMemcpy.cpp
index 527df9bab1..46656a434c 100644
--- a/hipamd/tests/src/deviceLib/hipDeviceMemcpy.cpp
+++ b/hipamd/tests/src/deviceLib/hipDeviceMemcpy.cpp
@@ -16,13 +16,13 @@
 
 __global__ void cpy(hipLaunchParm lp, uint32_t *Out, uint32_t *In)
 {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     memcpy(Out + tx, In + tx, sizeof(uint32_t));
 }
 
 __global__ void set(hipLaunchParm lp, uint32_t *ptr, uint8_t val, size_t size)
 {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     memset(ptr + tx, val, (sizeof(uint32_t)*(size/LEN)));
 }
 
@@ -58,6 +58,6 @@ int main()
             return 0;
         }
     }
- 
+
     passed();
 }
diff --git a/hipamd/tests/src/deviceLib/hipFloatMath.cpp b/hipamd/tests/src/deviceLib/hipFloatMath.cpp
index 7a96b5cd0d..9c97e50aa6 100644
--- a/hipamd/tests/src/deviceLib/hipFloatMath.cpp
+++ b/hipamd/tests/src/deviceLib/hipFloatMath.cpp
@@ -35,7 +35,7 @@ THE SOFTWARE.
 
 
 __global__ void floatMath(hipLaunchParm lp, float *In, float *Out) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   Out[tid] = __cosf(In[tid]);
   Out[tid] = __exp10f(Out[tid]);
   Out[tid] = __expf(Out[tid]);
diff --git a/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
index 74db5d0f73..5485e4afc9 100644
--- a/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
+++ b/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
@@ -217,7 +217,7 @@ int computeGold(int *gpuData, const int len)
 __global__ void testKernel(hipLaunchParm lp,int *g_odata)
 {
     // access thread id
-    const unsigned int tid = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
 
     // Test various atomic instructions
 
@@ -240,7 +240,7 @@ __global__ void testKernel(hipLaunchParm lp,int *g_odata)
 
     // Atomic increment (modulo 17+1)
     atomicInc((unsigned int *)&g_odata[5], 17);
- 
+
     // Atomic decrement
     atomicDec((unsigned int *)&g_odata[6], 137);
 
diff --git a/hipamd/tests/src/deviceLib/hipTestDevice.cpp b/hipamd/tests/src/deviceLib/hipTestDevice.cpp
index fa85940839..068866021b 100644
--- a/hipamd/tests/src/deviceLib/hipTestDevice.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestDevice.cpp
@@ -32,82 +32,82 @@ THE SOFTWARE.
 #define SIZE N*sizeof(float)
 
 __global__ void test_sincosf(hipLaunchParm lp, float* a, float* b, float *c){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     sincosf(a[tid], b+tid, c+tid);
 }
 
 __global__ void test_sincospif(hipLaunchParm lp, float* a, float* b, float *c){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     sincospif(a[tid], b+tid, c+tid);
 }
 
 __global__ void test_fdividef(hipLaunchParm lp, float *a, float* b, float *c){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     c[tid] = fdividef(a[tid], b[tid]);
 }
 
 __global__ void test_llrintf(hipLaunchParm lp, float *a, long long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = llrintf(a[tid]);
 }
 
 __global__ void test_lrintf(hipLaunchParm lp, float *a, long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = lrintf(a[tid]);
 }
 
 __global__ void test_rintf(hipLaunchParm lp, float *a, float *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = rintf(a[tid]);
 }
 
 __global__ void test_llroundf(hipLaunchParm lp, float *a, long long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = llroundf(a[tid]);
 }
 
 __global__ void test_lroundf(hipLaunchParm lp, float *a, long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = lroundf(a[tid]);
 }
 
 __global__ void test_rhypotf(hipLaunchParm lp, float *a, float* b, float *c){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     c[tid] = rhypotf(a[tid], b[tid]);
 }
 
 __global__ void test_norm3df(hipLaunchParm lp, float *a, float* b, float *c, float *d){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     d[tid] = norm3df(a[tid], b[tid], c[tid]);
 }
 
 __global__ void test_norm4df(hipLaunchParm lp, float *a, float* b, float *c, float *d, float *e){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     e[tid] = norm4df(a[tid], b[tid], c[tid], d[tid]);
 }
 
 __global__ void test_normf(hipLaunchParm lp, float *a, float *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = normf(N, a);
 }
 
 __global__ void test_rnorm3df(hipLaunchParm lp, float *a, float* b, float *c, float *d){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     d[tid] = rnorm3df(a[tid], b[tid], c[tid]);
 }
 
 __global__ void test_rnorm4df(hipLaunchParm lp, float *a, float* b, float *c, float *d, float *e){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     e[tid] = rnorm4df(a[tid], b[tid], c[tid], d[tid]);
 }
 
 __global__ void test_rnormf(hipLaunchParm lp, float *a, float *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = rnormf(N, a);
 }
 
 __global__ void test_erfinvf(hipLaunchParm lp, float *a, float *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = erff(erfinvf(a[tid]));
 }
 
diff --git a/hipamd/tests/src/deviceLib/hipTestDeviceDouble.cpp b/hipamd/tests/src/deviceLib/hipTestDeviceDouble.cpp
index 3b919d0cab..0b5d18fe5a 100644
--- a/hipamd/tests/src/deviceLib/hipTestDeviceDouble.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestDeviceDouble.cpp
@@ -32,72 +32,72 @@ THE SOFTWARE.
 #define SIZE N*sizeof(double)
 
 __global__ void test_sincos(hipLaunchParm lp, double* a, double* b, double *c){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     sincos(a[tid], b+tid, c+tid);
 }
 
 __global__ void test_sincospi(hipLaunchParm lp, double* a, double* b, double *c){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     sincospi(a[tid], b+tid, c+tid);
 }
 
 __global__ void test_llrint(hipLaunchParm lp, double *a, long long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = llrint(a[tid]);
 }
 
 __global__ void test_lrint(hipLaunchParm lp, double *a, long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = lrint(a[tid]);
 }
 
 __global__ void test_rint(hipLaunchParm lp, double *a, double *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = rint(a[tid]);
 }
 
 __global__ void test_llround(hipLaunchParm lp, double *a, long long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = llround(a[tid]);
 }
 
 __global__ void test_lround(hipLaunchParm lp, double *a, long int *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = lround(a[tid]);
 }
 
 __global__ void test_rhypot(hipLaunchParm lp, double *a, double* b, double *c){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     c[tid] = rhypot(a[tid], b[tid]);
 }
 
 __global__ void test_norm3d(hipLaunchParm lp, double *a, double* b, double *c, double *d){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     d[tid] = norm3d(a[tid], b[tid], c[tid]);
 }
 
 __global__ void test_norm4d(hipLaunchParm lp, double *a, double* b, double *c, double *d, double *e){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     e[tid] = norm4d(a[tid], b[tid], c[tid], d[tid]);
 }
 
 __global__ void test_rnorm3d(hipLaunchParm lp, double *a, double* b, double *c, double *d){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     d[tid] = rnorm3d(a[tid], b[tid], c[tid]);
 }
 
 __global__ void test_rnorm4d(hipLaunchParm lp, double *a, double* b, double *c, double *d, double *e){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     e[tid] = rnorm4d(a[tid], b[tid], c[tid], d[tid]);
 }
 
 __global__ void test_rnorm(hipLaunchParm lp, double *a, double *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = rnorm(N, a);
 }
 
 __global__ void test_erfinv(hipLaunchParm lp, double *a, double *b){
-    int tid = hipThreadIdx_x;
+    int tid = threadIdx.x;
     b[tid] = erf(erfinv(a[tid]));
 }
 
diff --git a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
index 140f3d97dc..be064f643c 100644
--- a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
@@ -36,7 +36,7 @@ __device__ int globalOut[NUM];
 
 __global__ void Assign(hipLaunchParm lp, int* Out)
 {
-    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
     Out[tid] = globalIn[tid];
     globalOut[tid] = globalIn[tid];
 }
diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 94a3882eea..842540795c 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -28,7 +28,7 @@ THE SOFTWARE.
 #if __HIP_ARCH_GFX803__ > 0
 
 __global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
-  int tx = hipThreadIdx_x;
+  int tx = threadIdx.x;
   __half a = A[tx];
   __half b = B[tx];
   __half c = C[tx];
@@ -44,7 +44,7 @@ __global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
 }
 
 __global__ void __half2Math(hipLaunchParm lp, __half2 *A, __half2 *B, __half2 *C) {
-  int tx = hipThreadIdx_x;
+  int tx = threadIdx.x;
   __half2 a = A[tx];
   __half2 b = B[tx];
   __half2 c = C[tx];
diff --git a/hipamd/tests/src/deviceLib/hipThreadFence.cpp b/hipamd/tests/src/deviceLib/hipThreadFence.cpp
index 1bd9fbe02c..2f73b68529 100644
--- a/hipamd/tests/src/deviceLib/hipThreadFence.cpp
+++ b/hipamd/tests/src/deviceLib/hipThreadFence.cpp
@@ -33,7 +33,7 @@ THE SOFTWARE.
 
 __global__ void vAdd(hipLaunchParm lp, float *In1, float *In2, float *In3, float *In4, float *Out)
 {
-    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
     In4[tid] = In1[tid] + In2[tid];
     __threadfence();
     In3[tid] = In3[tid] + In4[tid];
diff --git a/hipamd/tests/src/deviceLib/hip_anyall.cpp b/hipamd/tests/src/deviceLib/hip_anyall.cpp
index bba7915052..f0b314ce8d 100644
--- a/hipamd/tests/src/deviceLib/hip_anyall.cpp
+++ b/hipamd/tests/src/deviceLib/hip_anyall.cpp
@@ -37,9 +37,9 @@ __global__ void
 	warpvote(hipLaunchParm lp, int* device_any, int* device_all , int Num_Warps_per_Block, int pshift)
 {
 
-   int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
-   device_any[hipThreadIdx_x>>pshift] = __any(tid -77);
-   device_all[hipThreadIdx_x>>pshift] = __all(tid -77);
+   int tid = threadIdx.x + blockIdx.x * blockDim.x;
+   device_any[threadIdx.x>>pshift] = __any(tid -77);
+   device_all[threadIdx.x>>pshift] = __all(tid -77);
 }
 
 int main(int argc, char *argv[])
@@ -49,7 +49,7 @@ int main(int argc, char *argv[])
   warpSize = devProp.warpSize;
 
   int w = warpSize;
-  pshift = 0; 
+  pshift = 0;
   while (w >>= 1) ++pshift;
 
   printf ("warpSize=%d pshift=%d\n", warpSize, pshift);
diff --git a/hipamd/tests/src/deviceLib/hip_ballot.cpp b/hipamd/tests/src/deviceLib/hip_ballot.cpp
index 742c47a065..14b8f314a1 100644
--- a/hipamd/tests/src/deviceLib/hip_ballot.cpp
+++ b/hipamd/tests/src/deviceLib/hip_ballot.cpp
@@ -34,12 +34,12 @@ __global__ void
 	gpu_ballot(hipLaunchParm lp, unsigned int* device_ballot, int Num_Warps_per_Block,int pshift)
 {
 
-   int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
-   const unsigned int warp_num = hipThreadIdx_x >> pshift;
+   int tid = threadIdx.x + blockIdx.x * blockDim.x;
+   const unsigned int warp_num = threadIdx.x >> pshift;
 #ifdef __HIP_PLATFORM_HCC__
-   atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popcll(__ballot(tid - 245)));
+   atomicAdd(&device_ballot[warp_num+blockIdx.x*Num_Warps_per_Block],__popcll(__ballot(tid - 245)));
 #else
-	atomicAdd(&device_ballot[warp_num+hipBlockIdx_x*Num_Warps_per_Block],__popc(__ballot(tid - 245)));
+	atomicAdd(&device_ballot[warp_num+blockIdx.x*Num_Warps_per_Block],__popc(__ballot(tid - 245)));
 #endif
 
 }
diff --git a/hipamd/tests/src/deviceLib/hip_brev.cpp b/hipamd/tests/src/deviceLib/hip_brev.cpp
index 855a8bec47..c08c39dec9 100644
--- a/hipamd/tests/src/deviceLib/hip_brev.cpp
+++ b/hipamd/tests/src/deviceLib/hip_brev.cpp
@@ -72,8 +72,8 @@ HIP_kernel(hipLaunchParm lp,
              unsigned int* a, unsigned int* b,unsigned long long int* c, unsigned long long int* d, int width, int height)
   {
 
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
diff --git a/hipamd/tests/src/deviceLib/hip_clz.cpp b/hipamd/tests/src/deviceLib/hip_clz.cpp
index bdb31f3e8d..53fd611184 100644
--- a/hipamd/tests/src/deviceLib/hip_clz.cpp
+++ b/hipamd/tests/src/deviceLib/hip_clz.cpp
@@ -83,8 +83,8 @@ HIP_kernel(hipLaunchParm lp,
     unsigned int* a, unsigned int* b,unsigned int* c, unsigned long long int* d, int width, int height)
   {
 
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
diff --git a/hipamd/tests/src/deviceLib/hip_ffs.cpp b/hipamd/tests/src/deviceLib/hip_ffs.cpp
index c855ede060..49530bb298 100644
--- a/hipamd/tests/src/deviceLib/hip_ffs.cpp
+++ b/hipamd/tests/src/deviceLib/hip_ffs.cpp
@@ -73,8 +73,8 @@ HIP_kernel(hipLaunchParm lp,
 			    int width, int height)
   {
 
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
diff --git a/hipamd/tests/src/deviceLib/hip_popc.cpp b/hipamd/tests/src/deviceLib/hip_popc.cpp
index e503e55b42..19dafb4d43 100644
--- a/hipamd/tests/src/deviceLib/hip_popc.cpp
+++ b/hipamd/tests/src/deviceLib/hip_popc.cpp
@@ -64,8 +64,8 @@ HIP_kernel(hipLaunchParm lp,
              unsigned int* a, unsigned int* b,unsigned  int* c, unsigned long long int* d, int width, int height)
   {
 
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
diff --git a/hipamd/tests/src/deviceLib/hip_test_ldg.cpp b/hipamd/tests/src/deviceLib/hip_test_ldg.cpp
index 171ff1afd0..63d50e881e 100644
--- a/hipamd/tests/src/deviceLib/hip_test_ldg.cpp
+++ b/hipamd/tests/src/deviceLib/hip_test_ldg.cpp
@@ -57,8 +57,8 @@ vectoradd_float(hipLaunchParm lp,
              T* a, const T*  bm,  int width, int height)
 
   {
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
diff --git a/hipamd/tests/src/deviceLib/hip_test_make_type.cpp b/hipamd/tests/src/deviceLib/hip_test_make_type.cpp
index ef493ac923..6eba236e12 100644
--- a/hipamd/tests/src/deviceLib/hip_test_make_type.cpp
+++ b/hipamd/tests/src/deviceLib/hip_test_make_type.cpp
@@ -40,13 +40,13 @@ THE SOFTWARE.
 #define THREADS_PER_BLOCK_Z  1
 
 
-__global__ void 
+__global__ void
 vectoradd_char1(hipLaunchParm lp,
-             char1* a, const char1*  bm, const char1* cm, int width, int height) 
+             char1* a, const char1*  bm, const char1* cm, int width, int height)
 
   {
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
@@ -54,40 +54,40 @@ vectoradd_char1(hipLaunchParm lp,
       }
   }
 
-__global__ void 
+__global__ void
 vectoradd_char2(hipLaunchParm lp,
-             char2* a, const char2*  bm, const char2* cm, int width, int height) 
+             char2* a, const char2*  bm, const char2* cm, int width, int height)
 
   {
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
         a[i] = make_char2(bm[i].x, bm[i].y) + make_char2(cm[i].x, cm[i].y);
       }
-} 
+}
 
-__global__ void 
+__global__ void
 vectoradd_char3(hipLaunchParm lp,
-             char3* a, const char3*  bm, const char3* cm, int width, int height) 
+             char3* a, const char3*  bm, const char3* cm, int width, int height)
 
   {
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
         a[i] = make_char3(bm[i].x, bm[i].y, bm[i].z) + make_char3(cm[i].x, cm[i].y, cm[i].z);
       }
 }
-__global__ void 
+__global__ void
 vectoradd_char4(hipLaunchParm lp,
-             char4* a, const char4*  bm, const char4* cm, int width, int height) 
+             char4* a, const char4*  bm, const char4* cm, int width, int height)
 
   {
-      int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-      int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+      int x = blockDim.x * blockIdx.x + threadIdx.x;
+      int y = blockDim.y * blockIdx.y + threadIdx.y;
 
       int i = y * width + x;
       if ( i < (width * height)) {
@@ -99,7 +99,7 @@ vectoradd_char4(hipLaunchParm lp,
 #if 0
 __kernel__ void vectoradd_float(float* a, const float* b, const float* c, int width, int height) {
 
-  
+
   int x = blockDimX * blockIdx.x + threadIdx.x;
   int y = blockDimY * blockIdy.y + threadIdx.y;
 
@@ -128,21 +128,21 @@ bool dataTypesRun(){
   hostA = (T*)malloc(NUM * sizeof(T));
   hostB = (T*)malloc(NUM * sizeof(T));
   hostC = (T*)malloc(NUM * sizeof(T));
-  
+
   // initialize the input data
   for (i = 0; i < NUM; i++) {
     hostB[i] = (T)i;
     hostC[i] = (T)i;
   }
-  
+
   HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T)));
-  
+
   HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice));
   HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice));
 
-  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), 
+  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1),
                   dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
                   dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                   0, 0,
@@ -192,21 +192,21 @@ bool dataTypesRun(){
   hostA = (T*)malloc(NUM * sizeof(T));
   hostB = (T*)malloc(NUM * sizeof(T));
   hostC = (T*)malloc(NUM * sizeof(T));
-  
+
   // initialize the input data
   for (i = 0; i < NUM; i++) {
     hostB[i] = (T)i;
     hostC[i] = (T)i;
   }
-  
+
   HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T)));
-  
+
   HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice));
   HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice));
 
-  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), 
+  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1),
                   dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
                   dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                   0, 0,
@@ -256,21 +256,21 @@ bool dataTypesRun(){
   hostA = (T*)malloc(NUM * sizeof(T));
   hostB = (T*)malloc(NUM * sizeof(T));
   hostC = (T*)malloc(NUM * sizeof(T));
-  
+
   // initialize the input data
   for (i = 0; i < NUM; i++) {
     hostB[i] = (T)i;
     hostC[i] = (T)i;
   }
-  
+
   HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T)));
-  
+
   HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice));
   HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice));
 
-  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), 
+  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1),
                   dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
                   dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                   0, 0,
@@ -319,21 +319,21 @@ bool dataTypesRunChar4(){
   hostA = (T*)malloc(NUM * sizeof(T));
   hostB = (T*)malloc(NUM * sizeof(T));
   hostC = (T*)malloc(NUM * sizeof(T));
-  
+
   // initialize the input data
   for (i = 0; i < NUM; i++) {
     hostB[i] = (T)i;
     hostC[i] = (T)i;
   }
-  
+
   HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T)));
   HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T)));
-  
+
   HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice));
   HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice));
 
-  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), 
+  hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1),
                   dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
                   dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
                   0, 0,
@@ -368,7 +368,7 @@ bool dataTypesRunChar4(){
 }
 
 int main() {
-  
+
   hipDeviceProp_t devProp;
   hipGetDeviceProperties(&devProp, 0);
   cout << " System minor " << devProp.minor << endl;
diff --git a/hipamd/tests/src/deviceLib/hip_trig.cpp b/hipamd/tests/src/deviceLib/hip_trig.cpp
index 5ec28101f3..6ee8dc58ad 100644
--- a/hipamd/tests/src/deviceLib/hip_trig.cpp
+++ b/hipamd/tests/src/deviceLib/hip_trig.cpp
@@ -36,7 +36,7 @@ THE SOFTWARE.
 #define SIZE LEN<<2
 
 __global__ void kernel_trig(hipLaunchParm lp, float *In, float *sin_d, float *cos_d, float *tan_d, float *sin_pd, float *cos_pd){
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   sin_d[tid] = __sinf(In[tid]);
   cos_d[tid] = __cosf(In[tid]);
   tan_d[tid] = __tanf(In[tid]);
diff --git a/hipamd/tests/src/experimental/xcompile/hHip.c b/hipamd/tests/src/experimental/xcompile/hHip.c
index 2ac4ebc73e..17e7e9ecf6 100644
--- a/hipamd/tests/src/experimental/xcompile/hHip.c
+++ b/hipamd/tests/src/experimental/xcompile/hHip.c
@@ -29,7 +29,7 @@ THE SOFTWARE.
 
 __global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd, size_t len)
 {
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     if(tx < len)
     {
         Cd[tx] = Ad[tx] + Bd[tx];
diff --git a/hipamd/tests/src/experimental/xcompile/hipxxKer.cpp b/hipamd/tests/src/experimental/xcompile/hipxxKer.cpp
index d1bbed63cd..5dca6c1bca 100644
--- a/hipamd/tests/src/experimental/xcompile/hipxxKer.cpp
+++ b/hipamd/tests/src/experimental/xcompile/hipxxKer.cpp
@@ -30,7 +30,7 @@ THE SOFTWARE.
 
 __global__ void Kern(hipLaunchParm lp, float *A)
 {
-	int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+	int tx = threadIdx.x + blockIdx.x * blockDim.x;
 	A[tx] += 1.0f;
 }
 
diff --git a/hipamd/tests/src/experimental/xcompile/hxxHip.cpp b/hipamd/tests/src/experimental/xcompile/hxxHip.cpp
index 6a748d5c89..bca5d64afc 100644
--- a/hipamd/tests/src/experimental/xcompile/hxxHip.cpp
+++ b/hipamd/tests/src/experimental/xcompile/hxxHip.cpp
@@ -33,7 +33,7 @@ class memManager;
 template<typename T>
 __global__ void Add(hipLaunchParm lp, T* Ad, T* Bd, T* Cd, size_t Len)
 {
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     if(tx < Len)
     {
         Cd[tx] = Ad[tx] + Bd[tx];
diff --git a/hipamd/tests/src/hipC.c b/hipamd/tests/src/hipC.c
index 644df6c98f..efa03bb909 100644
--- a/hipamd/tests/src/hipC.c
+++ b/hipamd/tests/src/hipC.c
@@ -34,7 +34,7 @@ THE SOFTWARE.
 #define SIZE 1024*1024*sizeof(int)
 
 __global__ void Iter(hipLaunchParm lp, int *Ad){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     if(tx == 0){
         for(int i=0;i<ITER;i++){
             Ad[tx] += 1;
diff --git a/hipamd/tests/src/hipC.cpp b/hipamd/tests/src/hipC.cpp
index 8abb877808..6eac543ecf 100644
--- a/hipamd/tests/src/hipC.cpp
+++ b/hipamd/tests/src/hipC.cpp
@@ -29,7 +29,7 @@ THE SOFTWARE.
 #define SIZE 1024*1024*sizeof(int)
 
 __global__ void Iter(hipLaunchParm lp, int *Ad){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     if(tx == 0){
         for(int i=0;i<ITER;i++){
             Ad[tx] += 1;
diff --git a/hipamd/tests/src/hipCKernel.c b/hipamd/tests/src/hipCKernel.c
index 7a72cf84ca..891165f831 100644
--- a/hipamd/tests/src/hipCKernel.c
+++ b/hipamd/tests/src/hipCKernel.c
@@ -2,7 +2,7 @@
 #include "hip/hip_runtime_api.h"
 
 __global__ void Kernel(hipLaunchParm lp, float *Ad){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     Ad[tx] += Ad[tx-1];
 }
 
diff --git a/hipamd/tests/src/kernel/hipDynamicShared.cpp b/hipamd/tests/src/kernel/hipDynamicShared.cpp
index ba19fcaa0d..4c16c52884 100644
--- a/hipamd/tests/src/kernel/hipDynamicShared.cpp
+++ b/hipamd/tests/src/kernel/hipDynamicShared.cpp
@@ -40,8 +40,8 @@ __global__ void testExternSharedKernel(hipLaunchParm lp, const T* A_d, const T*
     T *sdata = reinterpret_cast<T *>(my_sdata);
 #endif
 
-    size_t gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t tid = hipThreadIdx_x;
+    size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t tid = threadIdx.x;
 
     // initialize dynamic shared memory
     if (tid < groupElements) {
diff --git a/hipamd/tests/src/kernel/hipDynamicShared2.cpp b/hipamd/tests/src/kernel/hipDynamicShared2.cpp
index 95e70a9956..4567ff6fc2 100644
--- a/hipamd/tests/src/kernel/hipDynamicShared2.cpp
+++ b/hipamd/tests/src/kernel/hipDynamicShared2.cpp
@@ -34,7 +34,7 @@ THE SOFTWARE.
 
 __global__ void vectorAdd(hipLaunchParm lp, float *Ad, float *Bd) {
   HIP_DYNAMIC_SHARED(float, sBd);
-  int tx = hipThreadIdx_x;
+  int tx = threadIdx.x;
   for(int i=0;i<LEN/64;i++) {
     sBd[tx + i * 64] = Ad[tx + i * 64] + 1.0f;
     Bd[tx + i * 64] = sBd[tx + i * 64];
diff --git a/hipamd/tests/src/kernel/hipGridLaunch.cpp b/hipamd/tests/src/kernel/hipGridLaunch.cpp
index 6cd724a070..f2cea64492 100644
--- a/hipamd/tests/src/kernel/hipGridLaunch.cpp
+++ b/hipamd/tests/src/kernel/hipGridLaunch.cpp
@@ -48,11 +48,12 @@ vectorADD2( hipLaunchParm lp,
             T *C_d,
             size_t N)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     for (size_t i=offset; i<N; i+=stride) {
-        C_d[i] = A_d[i] + B_d[i] ;
+        auto foo = __hiloint2double(A_d[i], B_d[i]);
+        C_d[i] = __double2loint(foo) + __double2hiint(foo);//A_d[i] + B_d[i] ;
     }
 }
 
diff --git a/hipamd/tests/src/kernel/hipLanguageExtensions.cpp b/hipamd/tests/src/kernel/hipLanguageExtensions.cpp
index 709fe3f1b8..5c5397675c 100644
--- a/hipamd/tests/src/kernel/hipLanguageExtensions.cpp
+++ b/hipamd/tests/src/kernel/hipLanguageExtensions.cpp
@@ -62,7 +62,7 @@ __global__ void MyKernel (const hipLaunchParm lp, const float *a, const float *b
 {
     //KERNELBEGIN;
 
-    unsigned gid = hipThreadIdx_x;
+    unsigned gid = threadIdx.x;
     if (gid < N) {
         c[gid] = a[gid] + PlusOne(b[gid]);
     }
@@ -96,7 +96,7 @@ vectorADD(const hipLaunchParm lp,
     int zuzu = deviceVar + 1;
 
 
-    int b = hipThreadIdx_x;
+    int b = threadIdx.x;
     int c;
 #ifdef NOT_YET
     int a = __shfl_up(x, 1);
@@ -119,8 +119,8 @@ vectorADD(const hipLaunchParm lp,
     __syncthreads();
 
 
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     for (size_t i=offset; i<N; i+=stride) {
 		C_d[i] = A_d[i] + B_d[i];
diff --git a/hipamd/tests/src/kernel/hipTestConstant.cpp b/hipamd/tests/src/kernel/hipTestConstant.cpp
index 9aeaf59c00..3922046a2e 100644
--- a/hipamd/tests/src/kernel/hipTestConstant.cpp
+++ b/hipamd/tests/src/kernel/hipTestConstant.cpp
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp 
+ * BUILD: %t %s ../test_common.cpp
  * RUN: %t
  * HIT_END
  */
@@ -38,7 +38,7 @@ __constant__ int Value[LEN];
 
 __global__ void Get(hipLaunchParm lp, int *Ad)
 {
-    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
     Ad[tid] = Value[tid];
 }
 
diff --git a/hipamd/tests/src/kernel/hipTestMallocKernel.cpp b/hipamd/tests/src/kernel/hipTestMallocKernel.cpp
index 9dd8b053a5..4aa0eb1cf7 100644
--- a/hipamd/tests/src/kernel/hipTestMallocKernel.cpp
+++ b/hipamd/tests/src/kernel/hipTestMallocKernel.cpp
@@ -33,12 +33,12 @@ THE SOFTWARE.
 #define SIZE NUM * 8
 
 __global__ void Alloc(hipLaunchParm lp, uint64_t *Ptr) {
-    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
     Ptr[tid] = (uint64_t)malloc(128);
 }
 
 __global__ void Free(hipLaunchParm lp, uint64_t *Ptr) {
-    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
     free((void*)Ptr[tid]);
 }
 
diff --git a/hipamd/tests/src/kernel/hipTestMemKernel.cpp b/hipamd/tests/src/kernel/hipTestMemKernel.cpp
index bac1905cd1..c44f275eed 100644
--- a/hipamd/tests/src/kernel/hipTestMemKernel.cpp
+++ b/hipamd/tests/src/kernel/hipTestMemKernel.cpp
@@ -35,52 +35,52 @@ THE SOFTWARE.
 #define LEN12 12 * 4
 
 __global__ void MemCpy8(hipLaunchParm lp, uint8_t *In, uint8_t *Out) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memcpy(Out + tid*8, In + tid*8, 8);
 }
 
 __global__ void MemCpy9(hipLaunchParm lp, uint8_t *In, uint8_t *Out) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memcpy(Out + tid*9, In + tid*9, 9);
 }
 
 __global__ void MemCpy10(hipLaunchParm lp, uint8_t *In, uint8_t *Out) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memcpy(Out + tid*10, In + tid*10, 10);
 }
 
 __global__ void MemCpy11(hipLaunchParm lp, uint8_t *In, uint8_t *Out) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memcpy(Out + tid*11, In + tid*11, 11);
 }
 
 __global__ void MemCpy12(hipLaunchParm lp, uint8_t *In, uint8_t *Out) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memcpy(Out + tid*12, In + tid*12, 12);
 }
 
 __global__ void MemSet8(hipLaunchParm lp, uint8_t *In) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memset(In + tid*8, 1, 8);
 }
 
 __global__ void MemSet9(hipLaunchParm lp, uint8_t *In) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memset(In + tid*9, 1, 9);
 }
 
 __global__ void MemSet10(hipLaunchParm lp, uint8_t *In) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memset(In + tid*10, 1, 10);
 }
 
 __global__ void MemSet11(hipLaunchParm lp, uint8_t *In) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memset(In + tid*11, 1, 11);
 }
 
 __global__ void MemSet12(hipLaunchParm lp, uint8_t *In) {
-  int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
   memset(In + tid*12, 1, 12);
 }
 
diff --git a/hipamd/tests/src/kernel/inline_asm_vadd.cpp b/hipamd/tests/src/kernel/inline_asm_vadd.cpp
index 7a941d31af..093fe1eb2d 100644
--- a/hipamd/tests/src/kernel/inline_asm_vadd.cpp
+++ b/hipamd/tests/src/kernel/inline_asm_vadd.cpp
@@ -37,7 +37,7 @@ __global__ void vadd_asm(hipLaunchParm lp,
                                 float *out,
                                 float *in)
 {
-    int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int i = blockDim.x * blockIdx.x + threadIdx.x;
 
     asm volatile ("v_add_f32_e32 %0, %1, %2" : "=v" (out[i]) : "v"(in[i]),"v" (out[i]));
 }
@@ -49,7 +49,7 @@ void addCPUReference(
 {
     for(unsigned int j=0; j < NUM; j++)
     {
-        
+
         output[j]= input[j] + output[j];
     }
 }
diff --git a/hipamd/tests/src/kernel/inline_asm_vmac.cpp b/hipamd/tests/src/kernel/inline_asm_vmac.cpp
index 1b6941c249..a78da01e74 100644
--- a/hipamd/tests/src/kernel/inline_asm_vmac.cpp
+++ b/hipamd/tests/src/kernel/inline_asm_vmac.cpp
@@ -35,7 +35,7 @@ __global__ void vmac_asm(hipLaunchParm lp,
                                 float *out,
                                 float *in)
 {
-    int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int i = blockDim.x * blockIdx.x + threadIdx.x;
 
     asm volatile ("v_mac_f32_e32 %0, %2, %3" : "=v" (out[i]) : "0"(out[i]), "v" (a), "v" (in[i]));
 }
@@ -47,7 +47,7 @@ void CPUReference(
 {
     for(unsigned int j=0; j < NUM; j++)
     {
-        
+
         output[j]= a*input[j] + output[j];
     }
 }
diff --git a/hipamd/tests/src/kernel/launch_bounds.cpp b/hipamd/tests/src/kernel/launch_bounds.cpp
index 3b1476fb11..269b88e34d 100644
--- a/hipamd/tests/src/kernel/launch_bounds.cpp
+++ b/hipamd/tests/src/kernel/launch_bounds.cpp
@@ -33,7 +33,7 @@ void
 __launch_bounds__(256, 2)
 myKern(hipLaunchParm lp, int *C, const int *A, int N, int xfactor)
 {
-    int tid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    int tid = (blockIdx.x * blockDim.x + threadIdx.x);
 
     if (tid < N) {
         C[tid] = A[tid];
diff --git a/hipamd/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp b/hipamd/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
index dac56bf709..c8c2e644c3 100644
--- a/hipamd/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
+++ b/hipamd/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
@@ -34,7 +34,7 @@ THE SOFTWARE.
 #define NUM_STREAMS 2
 
 __global__ void Iter(hipLaunchParm lp, int *Ad, int num){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     // Kernel loop designed to execute very slowly... ... ...   so we can test timing-related behavior below
     if(tx == 0){
         for(int i = 0; i<num;i++){
diff --git a/hipamd/tests/src/runtimeApi/memory/hipHostGetFlags.cpp b/hipamd/tests/src/runtimeApi/memory/hipHostGetFlags.cpp
index 9fad60aec8..f9359fb5da 100644
--- a/hipamd/tests/src/runtimeApi/memory/hipHostGetFlags.cpp
+++ b/hipamd/tests/src/runtimeApi/memory/hipHostGetFlags.cpp
@@ -33,7 +33,7 @@ THE SOFTWARE.
 #define SIZE LEN*sizeof(float)
 
 __global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd){
-int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+int tx = threadIdx.x + blockIdx.x * blockDim.x;
 Cd[tx] = Ad[tx] + Bd[tx];
 }
 
diff --git a/hipamd/tests/src/runtimeApi/memory/hipHostMalloc.cpp b/hipamd/tests/src/runtimeApi/memory/hipHostMalloc.cpp
index 607e2a9f63..4210ef0654 100644
--- a/hipamd/tests/src/runtimeApi/memory/hipHostMalloc.cpp
+++ b/hipamd/tests/src/runtimeApi/memory/hipHostMalloc.cpp
@@ -33,13 +33,13 @@
 #define SIZE LEN*sizeof(float)
 
 __global__ void Add(float *Ad, float *Bd, float *Cd){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     Cd[tx] = Ad[tx] + Bd[tx];
 }
 
 
 __global__ void Set(int *Ad, int val){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     Ad[tx] = val;
 }
 
@@ -52,13 +52,13 @@ std::vector<std::string> syncMsg = {"event", "stream", "device"};
 
 void CheckHostPointer(int numElements, int *ptr, unsigned eventFlags, int syncMethod, std::string msg)
 {
-    std::cerr << "test: CheckHostPointer " << msg 
+    std::cerr << "test: CheckHostPointer " << msg
               //<< " HIP_COHERENT_HOST_ALLOC=" << HIP_COHERENT_HOST_ALLOC
               //<< " HIP_EVENT_SYS_RELEASE=" <<   HIP_EVENT_SYS_RELEASE
-              << " eventFlags = " << std::hex << eventFlags 
-              << ((eventFlags & hipEventReleaseToDevice) ?  " hipEventReleaseToDevice" : "")  
-              << ((eventFlags & hipEventReleaseToSystem) ? " hipEventReleaseToSystem" : "") 
-              << " ptr=" << ptr 
+              << " eventFlags = " << std::hex << eventFlags
+              << ((eventFlags & hipEventReleaseToDevice) ?  " hipEventReleaseToDevice" : "")
+              << ((eventFlags & hipEventReleaseToSystem) ? " hipEventReleaseToSystem" : "")
+              << " ptr=" << ptr
               << " syncMethod=" << syncMsg[syncMethod] << "\n";
 
     hipStream_t s;
@@ -93,7 +93,7 @@ void CheckHostPointer(int numElements, int *ptr, unsigned eventFlags, int syncMe
         default:
             assert(0);
     };
-            
+
     for (int i=0; i<numElements; i++) {
         if (ptr[i] != expected) {
             printf ("mismatch at %d: %d != %d\n", i, ptr[i], expected);
@@ -153,7 +153,7 @@ int main(){
         size_t sizeBytes = numElements * sizeof (int);
 
 #ifdef __HIP_PLATFORM_HCC__
-        { 
+        {
             // Stimulate error condition:
             int *A = &numElements;
             HIPCHECK_API(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent|hipHostMallocNonCoherent), hipErrorInvalidValue);
@@ -174,7 +174,7 @@ int main(){
             // agent-scope releases don't provide host visibility, don't use them here:
         }
 
-        if (1) { 
+        if (1) {
             int *A = nullptr;
             HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent));
             const char *ptrType = "coherent";
@@ -189,14 +189,14 @@ int main(){
 
 
         // Check defaults:
-        if (1) { 
+        if (1) {
             int *A = nullptr;
             HIPCHECK(hipHostMalloc((void**)&A, sizeBytes));
             const char *ptrType = "default";
             CheckHostPointer(numElements, A, 0, SYNC_DEVICE,   ptrType);
             CheckHostPointer(numElements, A, 0, SYNC_STREAM,  ptrType);
             CheckHostPointer(numElements, A, 0, SYNC_EVENT,   ptrType);
-            
+
             CheckHostPointer(numElements, A, 0, SYNC_DEVICE,   ptrType);
             CheckHostPointer(numElements, A, 0, SYNC_STREAM,  ptrType);
             CheckHostPointer(numElements, A, 0, SYNC_EVENT,   ptrType);
@@ -206,7 +206,7 @@ int main(){
 
 
     }
-        
+
     passed();
 
 }
diff --git a/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp b/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
index 3376ee04f1..a65196ba89 100644
--- a/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
+++ b/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
@@ -29,13 +29,13 @@ THE SOFTWARE.
 #include<malloc.h>
 
 __global__ void Inc(hipLaunchParm lp, float *Ad){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     Ad[tx] = Ad[tx] + float(1);
 }
 
 
 template<typename T>
-void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internalRegister) 
+void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internalRegister)
 {
     A = A + offset;
     numElements -= offset;
@@ -56,7 +56,7 @@ void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internal
     HIPCHECK(hipMemset(Bd, 13.0f, sizeBytes));
 
 
-    // 
+    //
     HIPCHECK(hipMemcpy(Bd, A, sizeBytes,  hipMemcpyHostToDevice));
     HIPCHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost));
 
@@ -81,7 +81,7 @@ int main(int argc, char *argv[])
 
     const size_t size = N * sizeof(float);
 
-    if (p_tests & 0x1) { 
+    if (p_tests & 0x1) {
         float *A, **Ad;
         int num_devices;
         HIPCHECK(hipGetDeviceCount(&num_devices));
@@ -115,7 +115,7 @@ int main(int argc, char *argv[])
     }
 
 
-    if (p_tests & 0x6) { 
+    if (p_tests & 0x6) {
         // Sensitize HIP bug if device does not match where the memory was registered.
         HIPCHECK(hipSetDevice(0));
 
@@ -129,7 +129,7 @@ int main(int argc, char *argv[])
         HIPCHECK(hipMalloc(&Bd, size));
 
         // TODO - set to 128
-#define OFFSETS_TO_TRY 128 
+#define OFFSETS_TO_TRY 128
         assert (N>OFFSETS_TO_TRY);
 
         if (p_tests & 0x2) {
diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp
index 22bd30689a..5cd46c808a 100644
--- a/hipamd/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp
+++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpyAsync.cpp
@@ -59,7 +59,7 @@ struct HostTraits<Pinned>
     static const char *Name() { return "Pinned"; } ;
 
     static void *Alloc(size_t sizeBytes) {
-        void *p; 
+        void *p;
         HIPCHECK(hipHostMalloc((void**)&p, sizeBytes, hipHostMallocDefault));
         return p;
     };
@@ -67,11 +67,11 @@ struct HostTraits<Pinned>
 
 
 template<typename T>
-__global__ void 
+__global__ void
 addK (hipLaunchParm lp, T *A, T K, size_t numElements)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     for (size_t i=offset; i<numElements; i+=stride) {
         A[i] = A[i] + K;
@@ -85,7 +85,7 @@ addK (hipLaunchParm lp, T *A, T K, size_t numElements)
 //IN: numInflight : number of copies inflight at any time:
 //IN: numPongs = number of iterations to run (iteration)
 template<typename T, class AllocType>
-void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int numPongs, bool doHostSide) 
+void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int numPongs, bool doHostSide)
 {
     HIPASSERT(numElements % numInflight == 0); // Must be evenly divisible.
     size_t Nbytes = numElements*sizeof(T);
@@ -95,7 +95,7 @@ void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int
     unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements);
 
     printf ("-----------------------------------------------------------------------------------------------\n");
-    printf ("testing: %s<%s>  Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu eachCopyBytes=%zu\n", 
+    printf ("testing: %s<%s>  Nbytes=%zu (%6.1f MB) numPongs=%d numInflight=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
             __func__, HostTraits<AllocType>::Name(), Nbytes, (double)(Nbytes)/1024.0/1024.0, numPongs, numInflight, eachCopyElements, eachCopyBytes);
 
     T *A_h = NULL;
@@ -176,7 +176,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
     size_t eachCopyBytes = eachCopyElements * sizeof(T);
 
     printf ("-----------------------------------------------------------------------------------------------\n");
-    printf ("testing: %s  Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n", 
+    printf ("testing: %s  Nbytes=%zu (%6.1f MB) numCopies=%d eachCopyElements=%zu eachCopyBytes=%zu\n",
             __func__, Nbytes, (double)(Nbytes)/1024.0/1024.0, numCopies, eachCopyElements, eachCopyBytes);
 
     T *A_d;
@@ -194,7 +194,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
     //stream=0; // fixme TODO
 
 
-    for (int i=0; i<numCopies; i++) 
+    for (int i=0; i<numCopies; i++)
     {
         HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
         HIPCHECK(hipMemcpyAsync(&A_d[i*eachCopyElements], &A_h1[i*eachCopyElements], eachCopyBytes, hipMemcpyHostToDevice, stream));
@@ -204,7 +204,7 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
         HIPCHECK(hipDeviceSynchronize());
     }
 
-    for (int i=0; i<numCopies; i++) 
+    for (int i=0; i<numCopies; i++)
     {
         HIPASSERT(A_d + i*eachCopyElements < A_d + Nbytes);
         HIPCHECK(hipMemcpyAsync(&A_h2[i*eachCopyElements], &A_d[i*eachCopyElements], eachCopyBytes, hipMemcpyDeviceToHost, stream));
@@ -252,7 +252,7 @@ void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemc
 
 
     hipStream_t *stream = (hipStream_t*)malloc(sizeof(hipStream_t) * nStreams);
-    if (useNullStream) { 
+    if (useNullStream) {
         nStreams = 1;
         stream[0] = NULL;
     } else  {
@@ -262,7 +262,7 @@ void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemc
     }
 
 
-    size_t workLeft = N; 
+    size_t workLeft = N;
     size_t workPerStream = N / nStreams;
     for (int i = 0; i < nStreams; ++i) {
         size_t work = (workLeft < workPerStream) ? workLeft : workPerStream;
@@ -287,7 +287,7 @@ void test_chunkedAsyncExample(int nStreams, bool useNullStream, bool useSyncMemc
         } else {
             HIPCHECK ( hipMemcpyAsync(&C_h[offset], &C_d[offset], workBytes, hipMemcpyDeviceToHost, stream[i]));
         }
-    } 
+    }
 
 
     HIPCHECK (hipDeviceSynchronize());
diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocateCoherent.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocateCoherent.cpp
index 6042f538b3..667d4b404b 100644
--- a/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocateCoherent.cpp
+++ b/hipamd/tests/src/runtimeApi/memory/hipMemoryAllocateCoherent.cpp
@@ -31,7 +31,7 @@ THE SOFTWARE.
 
 __global__ void Kernel(hipLaunchParm lp,volatile float* hostRes)
 {
-    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tid = threadIdx.x + blockIdx.x * blockDim.x;
     hostRes[tid] = tid + 1;
     __threadfence_system();
     // expecting that the data is getting flushed to host here!
diff --git a/hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp b/hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp
index cad5065e78..81450f1fba 100644
--- a/hipamd/tests/src/runtimeApi/memory/p2p_copy_coherency.cpp
+++ b/hipamd/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 
+ * RUN: %t EXCLUDE_HIP_PLATFORM all
  * HIT_END
  */
 
@@ -63,8 +63,8 @@ int enablePeers(int dev0, int dev1)
 __global__ void
 memsetIntKernel(int * ptr, const int val, size_t numElements)
 {
-    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    int stride = hipBlockDim_x * hipGridDim_x ;
+    int gid = (blockIdx.x * blockDim.x + threadIdx.x);
+    int stride = blockDim.x * gridDim.x ;
     for (size_t i= gid; i< numElements; i+=stride){
        ptr[i] = val;
     }
@@ -73,15 +73,15 @@ memsetIntKernel(int * ptr, const int val, size_t numElements)
 __global__ void
 memcpyIntKernel(const int * src, int* dst, size_t numElements)
 {
-    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    int stride = hipBlockDim_x * hipGridDim_x ;
+    int gid = (blockIdx.x * blockDim.x + threadIdx.x);
+    int stride = blockDim.x * gridDim.x ;
     for (size_t i= gid; i< numElements; i+=stride){
        dst[i] = src[i];
     }
 };
 
 
-// CHeck arrays in reverse order, to more easily detect cases where 
+// CHeck arrays in reverse order, to more easily detect cases where
 // the copy is "partially" done.
 void checkReverse(const int *ptr, int numElements, int expected) {
     for (int i=numElements-1; i>=0; i--) {
diff --git a/hipamd/tests/src/runtimeApi/module/hipModule.cpp b/hipamd/tests/src/runtimeApi/module/hipModule.cpp
index 1b7b62cff2..f2c2137738 100644
--- a/hipamd/tests/src/runtimeApi/module/hipModule.cpp
+++ b/hipamd/tests/src/runtimeApi/module/hipModule.cpp
@@ -34,7 +34,7 @@ THE SOFTWARE.
 #define kernel_name "hello_world"
 
 __global__ void Cpy(hipLaunchParm lp, float *Ad, float* Bd){
-  int tx = hipThreadIdx_x;
+  int tx = threadIdx.x;
   Bd[tx] = Ad[tx];
 }
 
diff --git a/hipamd/tests/src/runtimeApi/module/vcpy_kernel.cpp b/hipamd/tests/src/runtimeApi/module/vcpy_kernel.cpp
index 0375eee342..7ee1ad333b 100644
--- a/hipamd/tests/src/runtimeApi/module/vcpy_kernel.cpp
+++ b/hipamd/tests/src/runtimeApi/module/vcpy_kernel.cpp
@@ -24,7 +24,7 @@ THE SOFTWARE.
 
 extern "C" __global__ void hello_world(hipLaunchParm lp, float *a, float *b)
 {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     b[tx] = a[tx];
 }
 
diff --git a/hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp b/hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
index 43a3e9bdea..3727901645 100644
--- a/hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
+++ b/hipamd/tests/src/runtimeApi/multiThread/hipMultiThreadStreams2.cpp
@@ -35,7 +35,7 @@ THE SOFTWARE.
 
 template<typename T>
 __global__ void Inc(hipLaunchParm lp, T *Array){
-int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+int tx = threadIdx.x + blockIdx.x * blockDim.x;
 Array[tx] = Array[tx] + T(1);
 }
 
@@ -116,7 +116,7 @@ int main(int argc, char **argv)
 	}
 
 	const size_t size = N * sizeof(float);
-    
+
     for (int i=0; i< iterations; i++) {
 
 	std::thread t1(run1, size, stream[0]);
@@ -126,7 +126,7 @@ int main(int argc, char **argv)
 //	std::cout<<"T1"<<std::endl;
 	t2.join();
 //	std::cout<<"T2"<<std::endl;
-	t3.join(); 
+	t3.join();
     }
 	passed();
 }
diff --git a/hipamd/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp b/hipamd/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp
index 4e343121ed..66b93a164f 100644
--- a/hipamd/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp
+++ b/hipamd/tests/src/runtimeApi/stream/hipAPIStreamDisable.cpp
@@ -29,7 +29,7 @@ THE SOFTWARE.
 const int NN = 1 << 21;
 
 __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	if(tid < 1){
 		for(int i=0;i<n;i++){
 			x[i] = sqrt(powf(3.14159,i));
@@ -39,7 +39,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
 }
 
 __global__ void nKernel(hipLaunchParm lp, float *y){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	y[tid] = y[tid] + 1.0f;
 }
 
diff --git a/hipamd/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp b/hipamd/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp
index 63cde4563e..e91b62231f 100644
--- a/hipamd/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp
+++ b/hipamd/tests/src/runtimeApi/stream/hipAPIStreamEnable.cpp
@@ -31,7 +31,7 @@ THE SOFTWARE.
 const int NN = 1 << 21;
 
 __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	if(tid < 1){
 		for(int i=0;i<n;i++){
 			x[i] = sqrt(powf(3.14159,i));
@@ -41,7 +41,7 @@ __global__ void kernel(hipLaunchParm lp, float *x, float *y, int n){
 }
 
 __global__ void nKernel(hipLaunchParm lp, float *y){
-	int tid = hipThreadIdx_x;
+	int tid = threadIdx.x;
 	y[tid] = y[tid] + 1.0f;
 }
 
diff --git a/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp b/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp
index af5547e3fb..8e46a69273 100644
--- a/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp
+++ b/hipamd/tests/src/runtimeApi/stream/hipNullStream.cpp
@@ -41,8 +41,8 @@ vectorADDRepeat(hipLaunchParm lp,
             size_t NELEM,
             int repeat)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     for (int j=1; j<=repeat;j++) {
         for (size_t i=offset; i<NELEM; i+=stride) {
diff --git a/hipamd/tests/src/runtimeApi/stream/hipStream.h b/hipamd/tests/src/runtimeApi/stream/hipStream.h
index 6468667703..0ce06bbc3f 100644
--- a/hipamd/tests/src/runtimeApi/stream/hipStream.h
+++ b/hipamd/tests/src/runtimeApi/stream/hipStream.h
@@ -73,7 +73,7 @@ void D2H(T *Dst, T *Src, size_t size){
 
 template<typename T>
 __global__ void Inc(hipLaunchParm lp, T *In){
-int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+int tx = threadIdx.x + blockIdx.x * blockDim.x;
 In[tx] = In[tx] + 1;
 }
 
@@ -94,12 +94,12 @@ void initArrays(T **Ad, T **Ah,
 }
 
 template<typename T>
-void initArrays(T **Ad, size_t N, 
-		bool deviceMemory = false, 
+void initArrays(T **Ad, size_t N,
+		bool deviceMemory = false,
 		bool usePinnedHost = false){
 	size_t NBytes = N * sizeof(T);
 	if(deviceMemory){
-		HIPCHECK( hipMalloc(Ad, NBytes)); 
+		HIPCHECK( hipMalloc(Ad, NBytes));
 	}else{
 		if(usePinnedHost){
 			HIPCHECK(hipHostMalloc((void**)Ad, NBytes, hipHostMallocDefault));
diff --git a/hipamd/tests/src/runtimeApi/synchronization/copy_coherency.cpp b/hipamd/tests/src/runtimeApi/synchronization/copy_coherency.cpp
index 2e514e1e3a..b2a66f61e2 100644
--- a/hipamd/tests/src/runtimeApi/synchronization/copy_coherency.cpp
+++ b/hipamd/tests/src/runtimeApi/synchronization/copy_coherency.cpp
@@ -27,14 +27,14 @@ THE SOFTWARE.
 // TODO - add code object support here.
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
- * RUN: %t 
+ * RUN: %t
  * HIT_END
  */
 
 
 // Test cache management (fences) and synchronization between kernel and copy commands.
-// Exhaustively tests 3 command types (copy, kernel, module kernel), 
-// many sync types (see SyncType), followed by another command, across a sweep 
+// Exhaustively tests 3 command types (copy, kernel, module kernel),
+// many sync types (see SyncType), followed by another command, across a sweep
 // of data sizes designed to stress various levels of the memory hierarchy.
 
 #include "hip/hip_runtime.h"
@@ -102,8 +102,8 @@ MemcpyFunction g_moduleMemcpy("memcpyInt.hsaco", "memcpyIntKernel");
 __global__ void
 memsetIntKernel(int * ptr, const int val, size_t numElements)
 {
-    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    int stride = hipBlockDim_x * hipGridDim_x ;
+    int gid = (blockIdx.x * blockDim.x + threadIdx.x);
+    int stride = blockDim.x * gridDim.x ;
     for (size_t i= gid; i< numElements; i+=stride){
        ptr[i] = val;
     }
@@ -112,15 +112,15 @@ memsetIntKernel(int * ptr, const int val, size_t numElements)
 __global__ void
 memcpyIntKernel(int *dst, const int * src, size_t numElements)
 {
-    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    int stride = hipBlockDim_x * hipGridDim_x ;
+    int gid = (blockIdx.x * blockDim.x + threadIdx.x);
+    int stride = blockDim.x * gridDim.x ;
     for (size_t i= gid; i< numElements; i+=stride){
        dst[i] = src[i];
     }
 };
 
 
-// CHeck arrays in reverse order, to more easily detect cases where 
+// CHeck arrays in reverse order, to more easily detect cases where
 // the copy is "partially" done.
 void checkReverse(const int *ptr, int numElements, int expected) {
     int mismatchCnt = 0;
@@ -157,7 +157,7 @@ const char * CmdTypeStr(CmdType c)
         ENUM_CASE_STR(MODULE_KERNEL);
         default: return "UNKNOWN";
     };
-} 
+}
 
 
 enum SyncType {
@@ -223,16 +223,16 @@ void resetInputs( int * Ad, int * Bd, int *Cd, int *Ch, size_t numElements, int
 
 // Intended to test proper synchronization and cache flushing between CMDA and CMDB.
 // CMD are of type CmdType. All command copy memory, using either hipMemcpyAsync or kernel implementations.
-// CmdA copies from Ad to Bd, 
+// CmdA copies from Ad to Bd,
 // Some form of synchronization is applied.
 // Then cmdB copies from Bd to Cd.
 //
 // Cd is then copied to host Ch using a memory copy.
 //
 // Correct result at the end is that Ch contains the contents originally in Ad (integer 0x42)
-void runTestImpl(CmdType cmdAType, SyncType syncType, CmdType cmdBType, 
+void runTestImpl(CmdType cmdAType, SyncType syncType, CmdType cmdBType,
             hipStream_t stream1, hipStream_t stream2, int numElements,
-             int * Ad, int * Bd, int *Cd, int *Ch, 
+             int * Ad, int * Bd, int *Cd, int *Ch,
              int expected)
 {
     hipEvent_t e;
@@ -241,14 +241,14 @@ void runTestImpl(CmdType cmdAType, SyncType syncType, CmdType cmdBType,
     resetInputs(Ad, Bd, Cd, Ch, numElements, expected);
 
     const size_t sizeElements = numElements * sizeof(int);
-    fprintf (stderr, "test: runTest with %zu bytes (%6.2f MB) cmdA=%s; sync=%s; cmdB=%s\n", 
+    fprintf (stderr, "test: runTest with %zu bytes (%6.2f MB) cmdA=%s; sync=%s; cmdB=%s\n",
             sizeElements, (double) (sizeElements/1024.0), CmdTypeStr(cmdAType), SyncTypeStr(syncType), CmdTypeStr(cmdBType));
 
     if (SKIP_MODULE_KERNEL && ((cmdAType == MODULE_KERNEL) || (cmdBType == MODULE_KERNEL))) {
         fprintf (stderr, "warn: skipping since test infra does not yet support modules\n");
         return;
     }
-            
+
 
     // Step A:
     runCmd(cmdAType, Bd, Ad, stream1, numElements);
@@ -334,7 +334,7 @@ void testWrapper(size_t numElements)
     fprintf (stderr, "test: init complete, start running tests\n");
 
 
-    runTestImpl(COPY, EVENT_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+    runTestImpl(COPY, EVENT_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
 
     for (int cmdA=0; cmdA<MAX_CmdType; cmdA++) {
         for (int cmdB=0; cmdB<MAX_CmdType; cmdB++) {
@@ -347,7 +347,7 @@ void testWrapper(size_t numElements)
                     //case STREAM_QUERY:
                     case STREAM_SYNC:
                     case DEVICE_SYNC:
-                        runTestImpl(CmdType(cmdA), SyncType(syncMode), CmdType(cmdB), stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+                        runTestImpl(CmdType(cmdA), SyncType(syncMode), CmdType(cmdB), stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
                         break;
                     default:
                         break;
@@ -358,11 +358,11 @@ void testWrapper(size_t numElements)
     }
 
 #if 0
-    runTestImpl(COPY, STREAM_SYNC, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
-    runTestImpl(COPY, STREAM_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
-    runTestImpl(COPY, STREAM_WAIT_EVENT, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+    runTestImpl(COPY, STREAM_SYNC, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
+    runTestImpl(COPY, STREAM_SYNC, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
+    runTestImpl(COPY, STREAM_WAIT_EVENT, MODULE_KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
 
-    runTestImpl(COPY, STREAM_WAIT_EVENT, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected); 
+    runTestImpl(COPY, STREAM_WAIT_EVENT, KERNEL, stream1, stream2, numElements, Ad, Bd, Cd, Ch, expected);
 #endif
 
     HIPCHECK(hipFree(Ad));
diff --git a/hipamd/tests/src/runtimeApi/synchronization/memcpyInt.device.cpp b/hipamd/tests/src/runtimeApi/synchronization/memcpyInt.device.cpp
index d86f36a1f4..2916d51bf9 100644
--- a/hipamd/tests/src/runtimeApi/synchronization/memcpyInt.device.cpp
+++ b/hipamd/tests/src/runtimeApi/synchronization/memcpyInt.device.cpp
@@ -1,12 +1,12 @@
 #include <hip/hip_runtime.h>
 
 
-    
+
 extern "C" __global__ void
 memcpyIntKernel(hipLaunchParm lp, int *dst, const int * src, size_t numElements)
 {
-    int gid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    int stride = hipBlockDim_x * hipGridDim_x ;
+    int gid = (blockIdx.x * blockDim.x + threadIdx.x);
+    int stride = blockDim.x * gridDim.x ;
     for (size_t i= gid; i< numElements; i+=stride){
        dst[i] = src[i];
     }
diff --git a/hipamd/tests/src/stress/hipStressAsync.cpp b/hipamd/tests/src/stress/hipStressAsync.cpp
index e06e16809c..a142b41730 100644
--- a/hipamd/tests/src/stress/hipStressAsync.cpp
+++ b/hipamd/tests/src/stress/hipStressAsync.cpp
@@ -30,7 +30,7 @@ THE SOFTWARE.
 #define ITER 1<<10
 
 __global__ void Iter(hipLaunchParm lp, int *Ad, int num){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     if(tx == 0){
         for(int i = 0; i<num;i++){
             Ad[tx] += 1;
@@ -58,7 +58,7 @@ int main(){
     }
 
     std::cout<<"Waitin..."<<std::endl;
-    
+
     hipDeviceSynchronize();
 }
 
diff --git a/hipamd/tests/src/stress/hipStressChain.cpp b/hipamd/tests/src/stress/hipStressChain.cpp
index 97183a97fa..55582d85e5 100644
--- a/hipamd/tests/src/stress/hipStressChain.cpp
+++ b/hipamd/tests/src/stress/hipStressChain.cpp
@@ -26,7 +26,7 @@ THE SOFTWARE.
 static size_t size[NUM_SIZE];
 
 __global__ void Add(hipLaunchParm lp, int *Ad){
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     Ad[tx] = Ad[tx] + tx;
 }
 
diff --git a/hipamd/tests/src/stress/hipStressKernel.cpp b/hipamd/tests/src/stress/hipStressKernel.cpp
index 52d8fa1fe9..90cc10894f 100644
--- a/hipamd/tests/src/stress/hipStressKernel.cpp
+++ b/hipamd/tests/src/stress/hipStressKernel.cpp
@@ -26,7 +26,7 @@ THE SOFTWARE.
 static size_t size[NUM_SIZE];
 
 __global__ void Add(hipLaunchParm lp, int *Ad){
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     Ad[tx] = Ad[tx] + tx;
 }
 
diff --git a/hipamd/tests/src/stress/hipStressSync.cpp b/hipamd/tests/src/stress/hipStressSync.cpp
index 44d004e59b..ef6a340937 100644
--- a/hipamd/tests/src/stress/hipStressSync.cpp
+++ b/hipamd/tests/src/stress/hipStressSync.cpp
@@ -30,7 +30,7 @@ THE SOFTWARE.
 #define ITER 1<<10
 
 __global__ void Iter(hipLaunchParm lp, int *Ad, int num){
-    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    int tx = threadIdx.x + blockIdx.x * blockDim.x;
     if(tx == 0){
         for(int i = 0; i<num;i++){
             Ad[tx] += 1;
@@ -57,7 +57,7 @@ int main(){
     }
 
     std::cout<<"Waitin..."<<std::endl;
-    
+
     hipDeviceSynchronize();
 }
 
diff --git a/hipamd/tests/src/test_common.h b/hipamd/tests/src/test_common.h
index 35f0aedb25..1b9744647b 100644
--- a/hipamd/tests/src/test_common.h
+++ b/hipamd/tests/src/test_common.h
@@ -136,8 +136,8 @@ vectorADD(hipLaunchParm lp,
             T *C_d,
             size_t NELEM)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     for (size_t i=offset; i<NELEM; i+=stride) {
         C_d[i] = A_d[i] + B_d[i];
@@ -153,8 +153,8 @@ vectorADDReverse(hipLaunchParm lp,
             T *C_d,
             size_t NELEM)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     for (int64_t i=NELEM-stride+offset; i>=0; i-=stride) {
         C_d[i] = A_d[i] + B_d[i];
@@ -169,8 +169,8 @@ addCount( const T *A_d,
         size_t NELEM,
         int count)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     // Deliberately do this in an inefficient way to increase kernel runtime
     for (int i=0; i<count; i++) {
@@ -188,8 +188,8 @@ addCountReverse( const T *A_d,
         int64_t NELEM,
         int count)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     // Deliberately do this in an inefficient way to increase kernel runtime
     for (int i=0; i<count; i++) {
@@ -205,8 +205,8 @@ __global__ void
 memsetReverse( T *C_d,  T val,
         int64_t NELEM)
 {
-    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
-    size_t stride = hipBlockDim_x * hipGridDim_x ;
+    size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x ;
 
     for (int64_t i=NELEM-stride+offset; i>=0; i-=stride) {
         C_d[i] = val;
@@ -219,12 +219,12 @@ void setDefaultData(size_t numElements, T *A_h, T* B_h, T *C_h)
 {
     // Initialize the host data:
     for (size_t i=0; i<numElements; i++) {
-        if (A_h) 
+        if (A_h)
             (A_h)[i] = 3.146f + i; // Pi
-        if (B_h) 
+        if (B_h)
             (B_h)[i] = 1.618f + i; // Phi
-        if (C_h) 
-            (C_h)[i] = 0.0f + i;  
+        if (C_h)
+            (C_h)[i] = 0.0f + i;
     }
 }
 
@@ -268,8 +268,8 @@ void initArraysForHost(T **A_h, T **B_h, T **C_h,
 
 template <typename T>
 void initArrays(T **A_d, T **B_d, T **C_d,
-                T **A_h, T **B_h, T **C_h, 
-                size_t N, bool usePinnedHost=false) 
+                T **A_h, T **B_h, T **C_h,
+                size_t N, bool usePinnedHost=false)
 {
     size_t Nbytes = N*sizeof(T);
 
@@ -317,7 +317,7 @@ void freeArraysForHost(T *A_h, T *B_h, T *C_h, bool usePinnedHost)
 
 template <typename T>
 void freeArrays(T *A_d, T *B_d, T *C_d,
-                T *A_h, T *B_h, T *C_h, bool usePinnedHost) 
+                T *A_h, T *B_h, T *C_h, bool usePinnedHost)
 {
     if (A_d) {
         HIPCHECK ( hipFree(A_d) );
@@ -453,9 +453,9 @@ struct Pinned {
 	static const bool isPinned = true;
 	static const char *str() { return "Pinned"; };
 
-    static void *Alloc(size_t sizeBytes) 
+    static void *Alloc(size_t sizeBytes)
 	{
-        void *p; 
+        void *p;
         HIPCHECK(hipHostMalloc((void**)&p, sizeBytes));
         return p;
     };
@@ -463,12 +463,12 @@ struct Pinned {
 
 
 //---
-struct Unpinned 
+struct Unpinned
 {
 	static const bool isPinned = false;
 	static const char *str() { return "Unpinned"; };
 
-    static void *Alloc(size_t sizeBytes) 
+    static void *Alloc(size_t sizeBytes)
 	{
         void *p  = malloc (sizeBytes);
         HIPASSERT(p);
@@ -496,7 +496,7 @@ template<>
 struct MemTraits<Memcpy>
 {
 
-    static void Copy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind,  hipStream_t stream) 
+    static void Copy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind,  hipStream_t stream)
 	{
 		HIPCHECK(hipMemcpy(dest, src, sizeBytes, kind));
 	}
@@ -507,7 +507,7 @@ template<>
 struct MemTraits<MemcpyAsync>
 {
 
-    static void Copy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind,  hipStream_t stream) 
+    static void Copy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind,  hipStream_t stream)
 	{
 		HIPCHECK(hipMemcpyAsync(dest, src, sizeBytes, kind, stream));
 	}
diff --git a/hipamd/tests/src/texture/hipTextureObj2D.cpp b/hipamd/tests/src/texture/hipTextureObj2D.cpp
index 443d708418..9ddafd6b1c 100644
--- a/hipamd/tests/src/texture/hipTextureObj2D.cpp
+++ b/hipamd/tests/src/texture/hipTextureObj2D.cpp
@@ -17,8 +17,8 @@ __global__ void tex2DKernel(float* outputData,
                              int width,
                              int height)
 {
-    int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x*blockDim.x + threadIdx.x;
+    int y = blockIdx.y*blockDim.y + threadIdx.y;
     outputData[y*width + x] = tex2D<float>(textureObject, x, y);
 }
 
diff --git a/hipamd/tests/src/texture/hipTextureRef2D.cpp b/hipamd/tests/src/texture/hipTextureRef2D.cpp
index eb27b23230..4430ca722d 100644
--- a/hipamd/tests/src/texture/hipTextureRef2D.cpp
+++ b/hipamd/tests/src/texture/hipTextureRef2D.cpp
@@ -18,8 +18,8 @@ __global__ void tex2DKernel(float* outputData,
                              int width,
                              int height)
 {
-    int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x*blockDim.x + threadIdx.x;
+    int y = blockIdx.y*blockDim.y + threadIdx.y;
 #ifdef __HIP_PLATFORM_HCC__
     outputData[y*width + x] = tex2D(tex, textureObject, x, y);
 #else

From fc34fd6f03813ff9ce648f98a5c7b4891a71ad52 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sun, 19 Nov 2017 15:21:37 -0600
Subject: [PATCH 7/9] Temporarily disable P2P on nvidia (fails on dual GPU)

---
 hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
index 9d46ccb0d8..b01a0aeb1d 100644
--- a/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
+++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp
@@ -23,7 +23,7 @@ THE SOFTWARE.
  */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp  EXCLUDE_HIP_PLATFORM nvcc
  * RUN: %t
  * HIT_END
  */

From aeadc3f18ff70472360c752e18774c542e4fa1c5 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sun, 19 Nov 2017 15:31:02 -0600
Subject: [PATCH 8/9] Fix test on cuda

---
 hipamd/tests/src/kernel/hipGridLaunch.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/hipamd/tests/src/kernel/hipGridLaunch.cpp b/hipamd/tests/src/kernel/hipGridLaunch.cpp
index f2cea64492..dff6527e2c 100644
--- a/hipamd/tests/src/kernel/hipGridLaunch.cpp
+++ b/hipamd/tests/src/kernel/hipGridLaunch.cpp
@@ -52,7 +52,7 @@ vectorADD2( hipLaunchParm lp,
     size_t stride = blockDim.x * gridDim.x ;
 
     for (size_t i=offset; i<N; i+=stride) {
-        auto foo = __hiloint2double(A_d[i], B_d[i]);
+        double foo = __hiloint2double(A_d[i], B_d[i]);
         C_d[i] = __double2loint(foo) + __double2hiint(foo);//A_d[i] + B_d[i] ;
     }
 }

From 822743845c5e76f7f24aeaeb7f3d1d1935c2368d Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 23 Nov 2017 09:57:24 +0530
Subject: [PATCH 9/9] Fix float2int rounding functions

Change-Id: I67943859a6344c5eec0eaa23418c9b802ef72468
---
 hipamd/src/device_functions.cpp | 13 +++++++++----
 1 file changed, 9 insertions(+), 4 deletions(-)

diff --git a/hipamd/src/device_functions.cpp b/hipamd/src/device_functions.cpp
index 615ae4d0b7..3c9bf334fa 100644
--- a/hipamd/src/device_functions.cpp
+++ b/hipamd/src/device_functions.cpp
@@ -23,6 +23,11 @@ THE SOFTWARE.
 #include <hc_math.hpp>
 #include "device_util.h"
 
+extern "C" float __ocml_floor_f32(float);
+extern "C" float __ocml_rint_f32(float);
+extern "C" float __ocml_ceil_f32(float);
+extern "C" float __ocml_trunc_f32(float);
+
 struct holder64Bit{
   union{
     double d;
@@ -151,19 +156,19 @@ __device__ long long int __double_as_longlong(double x)
 
 __device__ int __float2int_rd(float x)
 {
-  return (int)x;
+  return (int)__ocml_floor_f32(x);
 }
 __device__ int __float2int_rn(float x)
 {
-  return (int)x;
+  return (int)__ocml_rint_f32(x);
 }
 __device__ int __float2int_ru(float x)
 {
-  return (int)x;
+  return (int)__ocml_ceil_f32(x);
 }
 __device__ int __float2int_rz(float x)
 {
-  return (int)x;
+  return (int)__ocml_trunc_f32(x);
 }
 
 __device__ long long int __float2ll_rd(float x)