From a94de8f202dedeedb28eae5d5bc6969b42113ec3 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 11:22:27 +0530 Subject: [PATCH] SWDEV-339113 - update coordinate in HIP sample kernels (#2736) Change-Id: I8ea179b4ba8f1c0ebec830a5aa5947e843f06e42 --- samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp | 4 ++-- samples/2_Cookbook/1_hipEvent/hipEvent.cpp | 4 ++-- samples/2_Cookbook/3_shared_memory/sharedMemory.cpp | 4 ++-- samples/2_Cookbook/4_shfl/shfl.cpp | 2 +- samples/2_Cookbook/5_2dshfl/2dshfl.cpp | 4 ++-- samples/2_Cookbook/5_2dshfl/Readme.md | 2 +- 6 files changed, 10 insertions(+), 10 deletions(-) diff --git a/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp b/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp index 012debc148..fe71a03c5f 100644 --- a/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp +++ b/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp @@ -37,8 +37,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - 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; out[y * width + x] = in[x * width + y]; } diff --git a/samples/2_Cookbook/1_hipEvent/hipEvent.cpp b/samples/2_Cookbook/1_hipEvent/hipEvent.cpp index a83b99beb6..81013c1dd1 100644 --- a/samples/2_Cookbook/1_hipEvent/hipEvent.cpp +++ b/samples/2_Cookbook/1_hipEvent/hipEvent.cpp @@ -35,8 +35,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - 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; out[y * width + x] = in[x * width + y]; } diff --git a/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp b/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp index cff03f0a72..8bd489dbf3 100644 --- a/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp +++ b/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp @@ -38,8 +38,8 @@ THE SOFTWARE. __global__ void matrixTranspose(float* out, float* in, const int width) { __shared__ float sharedMem[WIDTH * WIDTH]; - 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; sharedMem[y * width + x] = in[x * width + y]; diff --git a/samples/2_Cookbook/4_shfl/shfl.cpp b/samples/2_Cookbook/4_shfl/shfl.cpp index 7902fd153a..de1ff7a950 100644 --- a/samples/2_Cookbook/4_shfl/shfl.cpp +++ b/samples/2_Cookbook/4_shfl/shfl.cpp @@ -36,7 +36,7 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int x = blockDim.x * blockIdx.x + threadIdx.x; float val = in[x]; diff --git a/samples/2_Cookbook/5_2dshfl/2dshfl.cpp b/samples/2_Cookbook/5_2dshfl/2dshfl.cpp index d2c26e7933..269ad58383 100644 --- a/samples/2_Cookbook/5_2dshfl/2dshfl.cpp +++ b/samples/2_Cookbook/5_2dshfl/2dshfl.cpp @@ -36,8 +36,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - 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; float val = in[y * width + x]; out[x * width + y] = __shfl(val, y * width + x); diff --git a/samples/2_Cookbook/5_2dshfl/Readme.md b/samples/2_Cookbook/5_2dshfl/Readme.md index fa10c71d6c..0419ef736c 100644 --- a/samples/2_Cookbook/5_2dshfl/Readme.md +++ b/samples/2_Cookbook/5_2dshfl/Readme.md @@ -31,7 +31,7 @@ We will be using the Simple Matrix Transpose application from the previous tutor In the same sourcecode, we used for MatrixTranspose. We'll add the following: ``` - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int y = blockDim.y * blockIdx.y + threadIdx.y; out[x*width + y] = __shfl(val,y*width + x); ```