Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Update Kernel Launch Example #3631

Open
wants to merge 2 commits into
base: docs/develop
Choose a base branch
from
Open
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 10 additions & 9 deletions docs/reference/cpp_language_extensions.rst
Original file line number Diff line number Diff line change
Expand Up @@ -104,7 +104,7 @@ You can include your kernel arguments after these parameters.
.. code-block:: cpp

// Example hipLaunchKernelGGL pseudocode:
__global__ MyKernel(float *A, float *B, float *C, size_t N)
__global__ void MyKernel(float *A, float *B, float *C, size_t N)
{
...
}
Expand All @@ -131,30 +131,31 @@ Kernel launch example

// Example showing device function, __device__ __host__
// <- compile for both device and host
float PlusOne(float x)
#include <hip/hip_runtime.h>
// Example showing device function, __device__ __host__
__host__ __device__ float PlusOne(float x) // <- compile for both device and host
{
return x + 1.0;
}

__global__
void
MyKernel (hipLaunchParm lp, /*lp parm for execution configuration */
const float *a, const float *b, float *c, unsigned N)
__global__ void MyKernel (const float *a, const float *b, float *c, unsigned N)
{
unsigned gid = threadIdx.x; // <- coordinate index function
const int gid = threadIdx.x + blockIdx.x * blockDim.x; // <- coordinate index function
if (gid < N) {
c[gid] = a[gid] + PlusOne(b[gid]);
}
}

void callMyKernel()
{
float *a, *b, *c; // initialization not shown...
unsigned N = 1000000;
const unsigned blockSize = 256;
const int gridSize = (N + blockSize - 1)/blockSize;

MyKernel<<<dim3(gridDim), dim3(groupDim), 0, 0>>> (a,b,c,n);
MyKernel<<<dim3(gridSize), dim3(blockSize), 0, 0>>> (a,b,c,N);
// Alternatively, kernel can be launched by
// hipLaunchKernelGGL(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0, a,b,c,N);
// hipLaunchKernelGGL(MyKernel, dim3(gridSize), dim3(blockSize), 0, 0, a,b,c,N);
}

Variable type qualifiers
Expand Down
Loading