[Documentation]: examples in "Reference" contain multiple syntax errors

[samuelpmish]

Description of errors

I'm a long-time CUDA developer trying to learn more about HIP and ROCm, but the documentation is making that difficult.

The most fundamental construct in CUDA is the kernel, so I started by reading https://rocm.docs.amd.com/projects/HIP/en/latest/reference/kernel_language.html

which contains a few example snippets illustrating how to write kernels for HIP. The first is:

// Example pseudo code introducing hipLaunchKernelGGL:
__global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N)
{
...
}

MyKernel<<<dim3(gridDim), dim3(groupDim), 0, 0>>> (a,b,c,n);
// Alternatively, kernel can be launched by
// hipLaunchKernelGGL(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n);

Maybe I'm missing something, but this code doesn't seem to work at all.

It has a handful of problems:

• it's not syntax highlighted (please just add cpp to the markdown code fences)
• the kernel definition is missing a void return type
• the triple-chevron launch example complains about the wrong number of parameters (maybe related to hipLaunchParam?)
• the hipLaunchKernelGGL call also complains about the wrong number of parameters

---

The next snippet on the page:

// Example showing device function, __device__ __host__
// <- compile for both device and host
float PlusOne(float x)
{
    return x + 1.0;
}

__global__
void
MyKernel (hipLaunchParm lp, /*lp parm for execution configuration */
          const float *a, const float *b, float *c, unsigned N)
{
    unsigned gid = threadIdx.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;

    MyKernel<<<dim3(gridDim), dim3(groupDim), 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);
}

also has a lot of errors

• it's also not syntax highlighted
• the PlusOne function is missing the __host__ __device__ annotations
• gridDim, groupDim not defined
• wrong variable names (N vs n)
• too few arguments in the kernel launch examples

Attach any links, screenshots, or additional evidence you think will be helpful.

No response

[ppanchad-amd]

@samuelpmish Internal ticket has been created to fix the documentation. Thanks!

[harkgill-amd]

Hi @samuelpmish, the examples have been updated in our docs staging branch which will be released in the near future. The changes can be found at https://github.com/ROCm/HIP/blob/docs/develop/docs/reference/cpp_language_extensions.rst#calling-global-functions. If you feel all the issues have not been addressed, please leave a comment and I will re-open this ticket. Thanks!