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

Change-Id: Ibfc0fd285441cd3d79b312d2b739729a039a6f84
This commit is contained in:
ROCm CI Service Account
2022-06-02 17:07:52 +05:30
committed by GitHub
parent 7b0ed9e8d3
commit 5da168300f
14 changed files with 47 additions and 47 deletions
@@ -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];
}
@@ -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.