SWDEV-341956 - update HIP Dtest with correct coordinates (#2755)

Change-Id: Iaccc4251a984e690d680ce884c6bfcc1093703af

tests/src/printf/hipPrintfBasic.cpp

@@ -36,7 +36,7 @@ THE SOFTWARE.
 DECLARE_DATA();
 
 __global__ void kernel_uniform0(int *retval) {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   retval[tid] = printf("Hello World\n"); // In Hip-Rocclr, printf returns number of characters printed.
                                          // In Cuda, printf returns the number of arguments parsed.
 }
@@ -78,7 +78,7 @@ static void test_uniform0(int *retval, uint num_blocks,
 }
 
 __global__ void kernel_uniform1(int *retval) {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   retval[tid] = printf("Six times Eight is %d\n", 42);
 }
 
@@ -119,7 +119,7 @@ static void test_uniform1(int *retval, uint num_blocks,
 }
 
 __global__ void kernel_divergent0(int *retval) {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   retval[tid] = printf("Thread ID: %d\n", tid);
 }
 
@@ -171,7 +171,7 @@ static void test_divergent0(int *retval, uint num_blocks,
 }
 
 __global__ void kernel_divergent1(int *retval) {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   if (tid % 2) {
     retval[tid] = printf("Hello World\n");
   } else {
@@ -222,7 +222,7 @@ static void test_divergent1(int *retval, uint num_blocks,
 __global__ void kernel_series(int *retval) {
   DECLARE_DATA();
 
-  const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  const uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   int result = 0;
   result += printf("%s\n", msg_long1);
   result += printf("%s\n", msg_short);
@@ -269,7 +269,7 @@ static void test_series(int *retval, uint num_blocks, uint threads_per_block) {
 }
 
 __global__ void kernel_divergent_loop() {
-  const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  const uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   int result = 0;
 
   for (int i = 0; i <= tid; ++i) {

tests/src/printf/hipPrintfManyDevices.cpp

@@ -34,7 +34,7 @@ DECLARE_DATA();
 __global__ void print_things() {
   DECLARE_DATA();
 
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   const char *msg[] = {msg_short, msg_long1, msg_long2};
 
   printf("%s\n", msg[tid % 3]);

tests/src/printf/hipPrintfManyWaves.cpp

@@ -38,7 +38,7 @@ DECLARE_DATA();
 __global__ void kernel_mixed0(int *retval) {
   DECLARE_DATA();
 
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
 
   // Three strings passed as divergent values to the same hostcall.
   const char *msg;
@@ -107,7 +107,7 @@ static void test_mixed0(int *retval, uint num_blocks, uint threads_per_block) {
 __global__ void kernel_mixed1(int *retval) {
   DECLARE_DATA();
 
-  const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  const uint tid = threadIdx.x + blockIdx.x * blockDim.x;
 
   // Three strings passed to divergent hostcalls.
   switch (tid % 3) {
@@ -173,7 +173,7 @@ static void test_mixed1(int *retval, uint num_blocks, uint threads_per_block) {
 __global__ void kernel_mixed2(int *retval) {
   DECLARE_DATA();
 
-  const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  const uint tid = threadIdx.x + blockIdx.x * blockDim.x;
 
   // Three different strings. All workitems print all three, but
   // in different orders.
@@ -230,7 +230,7 @@ static void test_mixed2(int *retval, uint num_blocks, uint threads_per_block) {
 __global__ void kernel_mixed3(int *retval) {
   DECLARE_DATA();
 
-  const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  const uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   int result = 0;
 
   result += printf("%s\n", msg_long1);
@@ -300,7 +300,7 @@ static void test_mixed3(int *retval, uint num_blocks, uint threads_per_block) {
 }
 
 __global__ void kernel_numbers() {
-  uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+  uint tid = threadIdx.x + blockIdx.x * blockDim.x;
   for (uint i = 0; i != 7; ++i) {
     uint base = tid * 21 + i * 3;
     printf("%d %d %d\n", base, base + 1, base + 2);