From 5ba644381c9fc9d94e13eb1d58c68e7ccff6c33c Mon Sep 17 00:00:00 2001 From: Harkirat Gill Date: Thu, 3 Oct 2024 14:37:25 -0400 Subject: [PATCH 1/2] Update Kernel Launch Example --- docs/reference/cpp_language_extensions.rst | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/docs/reference/cpp_language_extensions.rst b/docs/reference/cpp_language_extensions.rst index 48a817c31e..c8591d1dd9 100644 --- a/docs/reference/cpp_language_extensions.rst +++ b/docs/reference/cpp_language_extensions.rst @@ -131,30 +131,31 @@ Kernel launch example // Example showing device function, __device__ __host__ // <- compile for both device and host - float PlusOne(float x) + #include + // 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<<>> (a,b,c,n); + MyKernel<<>> (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 From a06e67300a35f04793739a415afa2299e40d1f6c Mon Sep 17 00:00:00 2001 From: Harkirat Gill Date: Thu, 3 Oct 2024 15:31:24 -0400 Subject: [PATCH 2/2] Update Kernel Launch Example --- docs/reference/cpp_language_extensions.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/reference/cpp_language_extensions.rst b/docs/reference/cpp_language_extensions.rst index c8591d1dd9..44f10d9120 100644 --- a/docs/reference/cpp_language_extensions.rst +++ b/docs/reference/cpp_language_extensions.rst @@ -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) { ... }