diff --git a/tests/src/deviceLib/hipTestAtomicAdd.cpp b/tests/src/deviceLib/hipTestAtomicAdd.cpp index 9d03249e10..25f818cf8b 100644 --- a/tests/src/deviceLib/hipTestAtomicAdd.cpp +++ b/tests/src/deviceLib/hipTestAtomicAdd.cpp @@ -75,7 +75,7 @@ bool p_atomicNoRet = false; template __global__ void atomicnoret_manywaves(T* C_d) { - size_t tid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + size_t tid = (blockIdx.x * blockDim.x + threadIdx.x); switch (tid % 9) { case 0: atomicAddNoRet(C_d, INCREMENT_VALUE); diff --git a/tests/src/dynamicLoading/bit_extract_kernel.cpp b/tests/src/dynamicLoading/bit_extract_kernel.cpp index 3a5141d817..2647a6c441 100644 --- a/tests/src/dynamicLoading/bit_extract_kernel.cpp +++ b/tests/src/dynamicLoading/bit_extract_kernel.cpp @@ -21,8 +21,8 @@ THE SOFTWARE. extern "C" __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* 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) { #ifdef __HIP_PLATFORM_AMD__ diff --git a/tests/src/dynamicLoading/complex_loading_behavior.cpp b/tests/src/dynamicLoading/complex_loading_behavior.cpp index f13f4b5c18..d83e98a882 100644 --- a/tests/src/dynamicLoading/complex_loading_behavior.cpp +++ b/tests/src/dynamicLoading/complex_loading_behavior.cpp @@ -37,8 +37,8 @@ THE SOFTWARE. #include "test_common.h" __global__ void vector_add(float* C, float* A, float* B, 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[i] = A[i] + B[i]; } diff --git a/tests/src/hostcall/hipHostcallFuncCall.cpp b/tests/src/hostcall/hipHostcallFuncCall.cpp index fd9923e372..d7ffc04bbc 100644 --- a/tests/src/hostcall/hipHostcallFuncCall.cpp +++ b/tests/src/hostcall/hipHostcallFuncCall.cpp @@ -36,7 +36,7 @@ static void callee(uint64_t* output, uint64_t* input) { } __global__ void kernel(uint64_t fptr, uint64_t* retval0, uint64_t* retval1) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; uint64_t arg0 = (uint64_t)fptr; uint64_t arg1 = tid; uint64_t arg2 = 42; diff --git a/tests/src/hostcall/hipHostcallPrintThings.cpp b/tests/src/hostcall/hipHostcallPrintThings.cpp index 9e53aeb9eb..02119079c5 100644 --- a/tests/src/hostcall/hipHostcallPrintThings.cpp +++ b/tests/src/hostcall/hipHostcallPrintThings.cpp @@ -43,7 +43,7 @@ void print_things_0(ulong* output, ulong* input) { } __global__ void kernel0(ulong fptr, ulong* retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; ulong arg0 = fptr; const char* str = "(%lu -> %lu)\n"; @@ -77,7 +77,7 @@ void print_things_1(ulong* output, const ulong* input) { } __global__ void kernel1(ulong fptr, ulong name, ulong* retval) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; ulong arg0 = fptr; ulong arg1 = name; ulong arg2 = tid; diff --git a/tests/src/kernel/hipShflTests.cpp b/tests/src/kernel/hipShflTests.cpp index 92205c4054..e1467aaa17 100644 --- a/tests/src/kernel/hipShflTests.cpp +++ b/tests/src/kernel/hipShflTests.cpp @@ -40,7 +40,7 @@ THE SOFTWARE. // Device (Kernel) function, it must be void template __global__ void matrixTranspose(T* out, T* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int x = blockDim.x * blockIdx.x + threadIdx.x; T val = in[x]; for (int i = 0; i < width; i++) { for (int j = 0; j < width; j++) out[i * width + j] = __shfl(val, j * width + i); diff --git a/tests/src/printf/hipPrintfBasic.cpp b/tests/src/printf/hipPrintfBasic.cpp index 8ad578cecc..6169dd01e4 100644 --- a/tests/src/printf/hipPrintfBasic.cpp +++ b/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) { diff --git a/tests/src/printf/hipPrintfManyDevices.cpp b/tests/src/printf/hipPrintfManyDevices.cpp index 2e7d99461f..004757cac3 100644 --- a/tests/src/printf/hipPrintfManyDevices.cpp +++ b/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]); diff --git a/tests/src/printf/hipPrintfManyWaves.cpp b/tests/src/printf/hipPrintfManyWaves.cpp index 0cb2b16c24..3bcb7fb4c1 100644 --- a/tests/src/printf/hipPrintfManyWaves.cpp +++ b/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); diff --git a/tests/src/texture/hipBindTex2DPitch.cpp b/tests/src/texture/hipBindTex2DPitch.cpp index 41c99a8a61..3fb8d21f8b 100644 --- a/tests/src/texture/hipBindTex2DPitch.cpp +++ b/tests/src/texture/hipBindTex2DPitch.cpp @@ -33,8 +33,8 @@ texture tex; // texture object is a kernel argument __global__ void texture2dCopyKernel( TYPE_t* dst) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT - 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); } diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index 2430f09906..0447c6aae2 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -60,7 +60,7 @@ template __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 =(float) elementID/numElements;