SWDEV-339113 - Update sampl codes with correct kernel coordinate (#2704)

Change-Id: Ibfc0fd285441cd3d79b312d2b739729a039a6f84

samples/0_Intro/bit_extract/bit_extract.cpp

@@ -35,8 +35,8 @@ THE SOFTWARE.
     }
 
 __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__

samples/0_Intro/module_api/vcpy_kernel.cpp

@@ -23,6 +23,6 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 
 extern "C" __global__ void hello_world(float* a, float* b) {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     b[tx] = a[tx];
 }

samples/0_Intro/module_api_global/vcpy_kernel.cpp

@@ -28,11 +28,11 @@ __device__ float myDeviceGlobal;
 __device__ float myDeviceGlobalArray[16];
 
 extern "C" __global__ void hello_world(const float* a, float* b) {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     b[tx] = a[tx];
 }
 
 extern "C" __global__ void test_globals(const float* a, float* b) {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
     b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE];
 }

samples/0_Intro/square/square.hipref.cpp

@@ -38,8 +38,8 @@ THE SOFTWARE.
  */
 template <typename T>
 __global__ void vector_square(T* C_d, const 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) {
         C_d[i] = A_d[i] * A_d[i];

samples/2_Cookbook/0_MatrixTranspose/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];
 }

samples/2_Cookbook/0_MatrixTranspose/Readme.md

@@ -27,8 +27,8 @@ __global__ void matrixTranspose(float *out,
                                 const int width,
                                 const int height)
 {
-    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+    int x = blockDim.x * blockIdx.x + bhreadIdx.x;
+    int y = blockDim.y * blockIdx.y + bhreadIdx.y;
 
     out[y * width + x] = in[x * height + y];
 }
@@ -39,7 +39,7 @@ other function-type qualifiers are:
 `__device__` functions are Executed on the device and Called from the device only
 `__host__` functions are Executed on the host and Called from the host
 
-`__host__` can combine with `__device__`, in which case the function compiles for both the host and device. These functions cannot use the HIP grid coordinate functions (for example, "hipThreadIdx_x", will talk about it latter). A possible workaround is to pass the necessary coordinate info as an argument to the function.
+`__host__` can combine with `__device__`, in which case the function compiles for both the host and device. These functions cannot use the HIP grid coordinate functions (for example, "threadIdx.x", will talk about it latter). A possible workaround is to pass the necessary coordinate info as an argument to the function.
 `__host__` cannot combine with `__global__`.
 
 `__global__` functions are often referred to as *kernels*, and calling one is termed *launching the kernel*.
@@ -47,9 +47,9 @@ other function-type qualifiers are:
 Next keyword is `void`. HIP `__global__` functions must have a `void` return type. Global functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute.
 
 The kernel function begins with
-`    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;`
-`    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;`
-here the keyword hipBlockIdx_x, hipBlockIdx_y and hipBlockIdx_z(not used here) are the built-in functions to identify the threads in a block. The keyword hipBlockDim_x, hipBlockDim_y and hipBlockDim_z(not used here) are to identify the dimensions of the block.
+`    int x = blockDim.x * blockIdx.x + threadIdx.x;`
+`    int y = blockDim.y * blockIdx.y + threadIdx.y;`
+here the keyword blockIdx.x, blockIdx.y and blockIdx.z(not used here) are the built-in functions to identify the threads in a block. The keyword blockDim.x, blockDim.y and blockDim.z(not used here) are to identify the dimensions of the block.
 
 We are familiar with rest of the code on device-side.
 

samples/2_Cookbook/10_inline_asm/inline_asm.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;
 
     asm volatile("v_mov_b32_e32 %0, %1" : "=v"(out[x * width + y]) : "v"(in[y * width + x]));
 }

samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp

@@ -32,49 +32,49 @@ texture<int4, hipTextureType2D, hipReadModeElementType> texInt4;
 texture<float4, hipTextureType2D, hipReadModeElementType> texFloat4;
 
 extern "C" __global__ void tex2dKernelChar(char* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texChar, x, y);
 }
 
 extern "C" __global__ void tex2dKernelShort(short* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texShort, x, y);
 }
 
 extern "C" __global__ void tex2dKernelInt(int* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texInt, x, y);
 }
 
 extern "C" __global__ void tex2dKernelFloat(float* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texFloat, x, y);
 }
 
 extern "C" __global__ void tex2dKernelChar4(char4* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texChar4, x, y);
 }
 
 extern "C" __global__ void tex2dKernelShort4(short4* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texShort4, x, y);
 }
 
 extern "C" __global__ void tex2dKernelInt4(int4* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texInt4, x, y);
 }
 
 extern "C" __global__ void tex2dKernelFloat4(float4* outputData, int width, int height) {
-    int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-    int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
     outputData[y * width + x] = tex2D(texFloat4, x, y);
 }

samples/2_Cookbook/12_cmake_hip_add_executable/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];
 }

samples/2_Cookbook/14_gpu_arch/gpuarch.cpp

@@ -36,7 +36,7 @@ THE SOFTWARE.
 // 'out'
 //              but it will update with "NOT_SUPPORTED" for any other gfx archs.
 __global__ void incrementKernel(int32_t* in, int32_t* out, int32_t value, size_t buffSize) {
-  int index = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+  int index = blockDim.x * blockIdx.x + threadIdx.x;
   if (index < buffSize) {
 #if defined(__gfx908__)
     out[index] = in[index] + value;

samples/2_Cookbook/6_dynamic_shared/dynamic_shared.cpp

@@ -37,8 +37,8 @@ THE SOFTWARE.
 __global__ void matrixTranspose(float* out, float* in, const int width) {
     extern __shared__ float sharedMem[];
 
-    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];
 

samples/2_Cookbook/7_streams/stream.cpp

@@ -37,8 +37,8 @@ __global__ void matrixTranspose_static_shared(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];
 
@@ -51,8 +51,8 @@ __global__ void matrixTranspose_dynamic_shared(float* out, float* in,
                                                const int width) {
     extern __shared__ float sharedMem[];
 
-    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];
 

samples/2_Cookbook/8_peer2peer/peer2peer.cpp

@@ -110,8 +110,8 @@ __global__ void matrixTranspose_static_shared(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];
 
@@ -124,8 +124,8 @@ __global__ void matrixTranspose_dynamic_shared(float* out, float* in,
                                                const int width) {
     extern __shared__ float sharedMem[];
 
-    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];
 

samples/2_Cookbook/9_unroll/unroll.cpp

@@ -43,7 +43,7 @@ void matrixRowSum(int* input, int* output, int width) {
 
 // Device (kernel) function
 __global__ void gpuMatrixRowSum(int* input, int* output, int width) {
-    int index = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int index = blockDim.x * blockIdx.x + threadIdx.x;
 #pragma unroll
     for (int i = 0; i < width; i++) {
         output[index] += input[index * width + i];