Kernel is a function that runs on the GPU, marked __global__ and called from host code with the <<<>>> launch syntax. It is executed once per thread by every thread in the grid you ask 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 (CUDA Error handling).
• 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.