From bebf1a2bc93b76e58fb43b4988785dc9e6f062be Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 23:28:10 +0530 Subject: [PATCH] SWDEV-339657 - Update catch codes with correct coordinates (#2735) Change-Id: I8ca78d770a742bf6c2abede494759caa923fcc19 [ROCm/hip-tests commit: d31aca651d338ba0d8e435a329a4e1385f4527c5] --- .../stress/memory/hipMemPrftchAsyncStressTst.cc | 4 ++-- .../stress/printf/Stress_printf_ComplexKernels.cc | 6 +++--- .../stress/printf/Stress_printf_SimpleKernels.cc | 12 ++++++------ .../graph/hipGraphExecEventRecordNodeSetEvent.cc | 2 +- .../catch/unit/memory/hipMemPrefetchAsync.cc | 4 ++-- .../catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc | 4 ++-- .../catch/unit/texture/hipBindTex2DPitch.cc | 4 ++-- .../catch/unit/texture/hipNormalizedFloatValueTex.cc | 2 +- 8 files changed, 19 insertions(+), 19 deletions(-) diff --git a/projects/hip-tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc b/projects/hip-tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc index bc2cb027da..0a159b8dbb 100644 --- a/projects/hip-tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc +++ b/projects/hip-tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc @@ -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]; } diff --git a/projects/hip-tests/catch/stress/printf/Stress_printf_ComplexKernels.cc b/projects/hip-tests/catch/stress/printf/Stress_printf_ComplexKernels.cc index 16796d7972..eec8e162f7 100644 --- a/projects/hip-tests/catch/stress/printf/Stress_printf_ComplexKernels.cc +++ b/projects/hip-tests/catch/stress/printf/Stress_printf_ComplexKernels.cc @@ -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__ diff --git a/projects/hip-tests/catch/stress/printf/Stress_printf_SimpleKernels.cc b/projects/hip-tests/catch/stress/printf/Stress_printf_SimpleKernels.cc index e116593a4a..1ec20ec873 100644 --- a/projects/hip-tests/catch/stress/printf/Stress_printf_SimpleKernels.cc +++ b/projects/hip-tests/catch/stress/printf/Stress_printf_SimpleKernels.cc @@ -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); } diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc index 69c60635bf..5a1cbe9997 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc @@ -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]; } diff --git a/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsync.cc b/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsync.cc index 6f8fe65c74..17ef618b77 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsync.cc @@ -20,8 +20,8 @@ THE SOFTWARE. #include // 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]; } diff --git a/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc b/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc index c5a51d9488..482a73cd02 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc @@ -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]; } diff --git a/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc b/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc index eff6a9f939..c1d82ebfbe 100644 --- a/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc +++ b/projects/hip-tests/catch/unit/texture/hipBindTex2DPitch.cc @@ -28,8 +28,8 @@ texture 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); } diff --git a/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc index e9cf002eae..8113117a64 100644 --- a/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc +++ b/projects/hip-tests/catch/unit/texture/hipNormalizedFloatValueTex.cc @@ -43,7 +43,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 = elementID/static_cast(numElements);