A kernel is a function that runs on the GPU. It is marked __global__, called from host code with the <<<>>> launch syntax, and executed once per thread by every thread in the grid you asked for.

__global__ void scale(float *a, float k) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    a[i] *= k;
}
 
scale<<<64, 256>>>(d_a, 2.0f);   // 64 blocks x 256 threads = 16384 threads

There is no loop in the kernel body. The loop is the launch itself, and each thread runs one iteration of it. This inversion is the main mental shift coming from OpenMP, where you keep the loop and annotate it.

Kernels have some hard restrictions:

• The return type must be void. Results come back through pointers to device memory.
• Arguments are passed by value into a small parameter buffer, so pass pointers, not host structs holding pointers.
• A launch is asynchronous. The host continues immediately, and errors show up later (Error handling (CUDA)).
• Recursion and function pointers work only in limited forms, and dynamic allocation inside a kernel is possible but slow.

A __device__ function is callable only from device code and is normally inlined. A __host__ __device__ function is compiled twice, once for each side, which is how small utility functions get shared between CPU and GPU code.