SWDEV-339657 - Update catch codes with correct coordinates (#2735)

Change-Id: I8ca78d770a742bf6c2abede494759caa923fcc19 [ROCm/hip-tests commit: d31aca651d]
2022-06-15 23:28:10 +05:30
@@ -26,8 +26,8 @@ THE SOFTWARE.

 // Kernel function
 __global__ void MemPrftchAsyncKernel1(int* Hmm, 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) {
    Hmm[i] = Hmm[i] * Hmm[i];
  }
@@ -230,19 +230,19 @@ __device__ __host__ struct printInfo startPrint(uint32_t tid,
 // This kernel is launched only in X dimension
 __global__ void kernel_complex_opX(uint32_t *a, uint32_t *b,
                                  uint32_t iterCount) {
-  uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x;
  startPrint(tid, iterCount, a, b);
 }
 // This kernel is launched only in Y dimension
 __global__ void kernel_complex_opY(uint32_t *a, uint32_t *b,
                                  uint32_t iterCount) {
-  uint32_t tid = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y;
+  uint32_t tid = threadIdx.y + blockIdx.y * blockDim.y;
  startPrint(tid, iterCount, a, b);
 }
 // This kernel is launched only in Z dimension
 __global__ void kernel_complex_opZ(uint32_t *a, uint32_t *b,
                                  uint32_t iterCount) {
-  uint32_t tid = hipThreadIdx_z + hipBlockIdx_z * hipBlockDim_z;
+  uint32_t tid = threadIdx.z + blockIdx.z * blockDim.z;
  startPrint(tid, iterCount, a, b);
 }
 #ifdef __linux__
@@ -85,7 +85,7 @@ __global__ void kernel_printf_conststr(uint iterCount) {
 // 'g' grid size such that (total bytes per iteration)*n*b*g ≈ N GB,
 // where N is user input.
 __global__ void kernel_printf_two_conditionalstr(uint iterCount) {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
  uint mod_tid = (tid % 2);
  if (0 == mod_tid) {
    for (uint count = 0; count < iterCount; count++) {
@@ -101,7 +101,7 @@ __global__ void kernel_printf_two_conditionalstr(uint iterCount) {
 // iterations per thread using 'b' block size and 'g' grid size such that
 // (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
 __global__ void kernel_printf_single_conditionalstr(uint iterCount) {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
  uint mod_tid = (tid % 2);
  if (0 == mod_tid) {
    for (uint count = 0; count < iterCount; count++) {
@@ -115,7 +115,7 @@ __global__ void kernel_printf_single_conditionalstr(uint iterCount) {
 // iterations per thread using 'b' block size and 'g' grid size such
 // that (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
 __global__ void kernel_printf_variablestr(uint iterCount, int *ret) {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
  int retlocal = 0;
  const char *const_str =
  "Hello World from Device.Iam printing (threadID,number)=";
@@ -134,7 +134,7 @@ __global__ void kernel_printf_variablestr(uint iterCount, int *ret) {
 // size and 'g' grid size such that
 // (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
 __global__ void kernel_dependent_calc(uint32_t iterCount, int *ret) {
-  uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x;
  int retlocal = 0;
  const char *const_str =
  "Hello World from Device.Iam printing number=";
@@ -158,7 +158,7 @@ __global__ void kernel_dependent_calc(uint32_t iterCount, int *ret) {
 // (total bytes per iteration)*n*b*g ≈ N GB, where N is user input.
 __global__ void kernel_dependent_calc_atomic(uint32_t iterCount,
                                             int *ret) {
-  uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x;
  int retlocal = 0;
  const char *const_str =
  "Hello World from Device.Iam printing number=";
@@ -209,7 +209,7 @@ __global__ void kernel_shared_mem() {
  __shared__ uint32_t sharedMem;
  sharedMem = 0;
  __syncthreads();
-  atomicAdd(&sharedMem, hipThreadIdx_x);
+  atomicAdd(&sharedMem, threadIdx.x);
  __syncthreads();
  printf("%s%u\n", CONST_STR3, sharedMem);
 }
@@ -57,7 +57,7 @@ Testcase Scenarios :
 * Kernel Functions to copy.
 */
 static __global__ void copy_ker_func(int* a, int* b) {
-  int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
+  int tx = blockIdx.x*blockDim.x + threadIdx.x;
  if (tx < LEN) b[tx] = a[tx];
 }

@@ -20,8 +20,8 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 // Kernel function
 __global__ void MemPrftchAsyncKernel(int* C_d, const int* A_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] * A_d[i];
  }
@@ -37,8 +37,8 @@ THE SOFTWARE.
 // Kernel function

 __global__ void MemPrftchAsyncKernel1(int* Hmm, 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) {
    Hmm[i] = Hmm[i] * Hmm[i];
  }
@@ -28,8 +28,8 @@ texture<TYPE_t, 2, hipReadModeElementType> tex;

 // texture object is a kernel argument
 static __global__ void texture2dCopyKernel(TYPE_t* dst) {
-    int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
-    int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y;
+    int x = threadIdx.x + blockIdx.x * blockDim.x;
+    int y = threadIdx.y + blockIdx.y * blockDim.y;
    if ( (x < SIZE_W) && (y < SIZE_H) ) {
        dst[SIZE_W*y+x] = tex2D(tex, x, y);
    }
@@ -43,7 +43,7 @@ template<typename T>
 __global__ void normalizedValTextureTest(unsigned int numElements,
                                         float* pDst) {
 #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
-  unsigned int elementID = hipThreadIdx_x;
+  unsigned int elementID = threadIdx.x;
  if (elementID >= numElements)
    return;
  float coord = elementID/static_cast<float>(numElements);