Your First GPU Kernel: CUDA and HIP Side by Side

A GPU program has two halves. The kernel is a function that runs on the device, once per thread, and is written as if you were a single thread. The host code runs on the CPU and does the logistics: allocate device memory, copy data across, launch the kernel over a grid of threads, copy results back, free everything.

The surprise for most people is how little of this differs between Nvidia and AMD. The kernel body is essentially identical. Almost all of the portability work lives in the host half.

Vector addition is the canonical first kernel because it is the simplest thing with real parallelism: every output element is independent of every other one, which is exactly the shape the hardware wants.

__global__ void vecAdd(const float *a, const float *b, float *c, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}
            

Here the dialects diverge slightly, and only in syntax.

// CUDA
vecAdd<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, n);

// HIP, portable macro form
hipLaunchKernelGGL(vecAdd, numBlocks, threadsPerBlock, 0, 0, d_a, d_b, d_c, n);

// HIP also accepts the CUDA-style form when built with hipcc
vecAdd<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, n);
            

Every GPU program of this shape does the same five things, and this is where the API names actually differ.

The pattern is a straight cuda to hip rename across the board, which is deliberate and is what makes automated porting viable.

#include <hip/hip_runtime.h>   // or <cuda_runtime.h>
#include <vector>
#include <cstdio>

__global__ void vecAdd(const float *a, const float *b, float *c, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

int main()
{
    const int n = 1 << 20;
    const size_t bytes = n * sizeof(float);

    std::vector<float> h_a(n, 1.0f), h_b(n, 2.0f), h_c(n);

    float *d_a, *d_b, *d_c;
    hipMalloc(&d_a, bytes);
    hipMalloc(&d_b, bytes);
    hipMalloc(&d_c, bytes);

    hipMemcpy(d_a, h_a.data(), bytes, hipMemcpyHostToDevice);
    hipMemcpy(d_b, h_b.data(), bytes, hipMemcpyHostToDevice);

    const int threadsPerBlock = 256;
    const int numBlocks = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, n);

    hipMemcpy(h_c.data(), d_c, bytes, hipMemcpyDeviceToHost);

    printf("c[0] = %f, c[n-1] = %f\n", h_c[0], h_c[n-1]);

    hipFree(d_a); hipFree(d_b); hipFree(d_c);
    return 0;
}
            

Three things are worth carrying forward from this small program.

The copy is not free. This kernel does one addition per two loaded floats, which is almost no arithmetic per byte moved. Vector addition is bandwidth-bound and always will be, and on a problem this small the host transfers cost far more than the arithmetic. It is a teaching example, not a demonstration of speedup.

No error was checked. Every call above returns a status code and this program ignores all of them, which is why a broken GPU program so often fails silently and far from the cause.

The grid was sized to the data. One thread per element is the obvious mapping and not always the best one; the grid-stride alternative is on the next page.