CUDA Kernel Programming — Complete Tutorial
(with expanded practical code examples + timing)

This is the next important step after understanding GPU → SM → Warp → Thread → Memory.

The goal here is to understand:

textCopy
1CPU Program
2    ↓
3Launch CUDA Kernel
4    ↓
5GPU
6    ↓
7Grid
8    ↓
9Blocks
10    ↓
11Warps
12    ↓
13Threads

1. What is a CUDA Kernel?

A CUDA kernel is a function that is executed on the GPU.

In CUDA C++, you mark a GPU kernel with:

cppCopy
1__global__

Example:

cppCopy
1__global__ void hello_kernel()
2{
3    printf("Hello from GPU\n");
4}

You launch it from CPU code:

cppCopy
1hello_kernel<<<1, 1>>>();

The important difference is:

cppCopy
1void normal_function()
2{
3    // CPU function
4}

versus:

cppCopy
1__global__ void gpu_kernel()
2{
3    // GPU kernel
4}

2. CPU Code vs GPU Code

A CUDA program usually contains two sides:

textCopy
1             CUDA Program
2                  │
3        ┌─────────┴─────────┐
4        ↓                   ↓
5      Host                Device
6       CPU                 GPU
7        │                   │
8        │ launch            │
9        └───────────────────→│
10                           Kernel

CUDA calls the CPU Host and the GPU Device.

cppCopy
1// Host code
2int main()
3{
4    // CPU
5}
6
7// Device code
8__global__ void kernel()
9{
10    // GPU
11}

3. __global__

__global__ is used to define a CUDA kernel.

Properties:

• executes on GPU
• normally launched from host/CPU
• launched with <<< >>> syntax
• launch is asynchronous with respect to the host

Complete example with timing:

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void hello()
5{
6    printf("Hello GPU from thread %d in block %d\n",
7           threadIdx.x, blockIdx.x);
8}
9
10int main()
11{
12    cudaEvent_t start, stop;
13    cudaEventCreate(&start);
14    cudaEventCreate(&stop);
15
16    // Warm-up
17    hello<<<1, 1>>>();
18    cudaDeviceSynchronize();
19
20    // Timed launch
21    cudaEventRecord(start);
22    hello<<<2, 4>>>();               // 2 blocks × 4 threads
23    cudaEventRecord(stop);
24    cudaEventSynchronize(stop);
25
26    float ms = 0.0f;
27    cudaEventElapsedTime(&ms, start, stop);
28    printf("Kernel time: %.3f ms\n", ms);
29
30    cudaEventDestroy(start);
31    cudaEventDestroy(stop);
32    return 0;
33}

4. Why <<< >>>?

This syntax:

cppCopy
1hello<<<1, 1>>>();

is CUDA kernel launch syntax.

textCopy
1<<< blocks, threads >>>

Example:

cppCopy
1hello<<<4, 256>>>();

means:

textCopy
14 blocks × 256 threads/block = 1024 logical threads

5. First CUDA Kernel (Hello World)

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void hello_kernel()
5{
6    printf("Hello from GPU!\n");
7}
8
9int main()
10{
11    hello_kernel<<<1, 1>>>();
12    cudaDeviceSynchronize();
13    return 0;
14}

Compile & run:

bashCopy
1nvcc hello.cu -o hello && ./hello

6. Why cudaDeviceSynchronize()?

A kernel launch such as:

cppCopy
1hello_kernel<<<1, 1>>>();

returns control to the CPU before the GPU has necessarily finished.

textCopy
1CPU
2 │
3 │ launch kernel
4 ↓
5GPU starts kernel
6 │
7 CPU can continue
8 │
9 ↓
10cudaDeviceSynchronize()   ← wait for GPU

Always use it in examples so the program does not exit before GPU work completes, and so errors become visible.

7. __device__

__device__ defines a function that runs on the GPU and is called from GPU code.

Complete example:

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__device__ int square(int x)
5{
6    return x * x;
7}
8
9__global__ void kernel(int* output)
10{
11    int tid = threadIdx.x;
12    output[tid] = square(tid + 1);   // 1², 2², 3² ...
13}
14
15int main()
16{
17    const int N = 8;
18    int *d_out, h_out[N];
19
20    cudaMalloc(&d_out, N * sizeof(int));
21
22    kernel<<<1, N>>>(d_out);
23    cudaDeviceSynchronize();
24
25    cudaMemcpy(h_out, d_out, N * sizeof(int), cudaMemcpyDeviceToHost);
26
27    for (int i = 0; i < N; i++)
28        printf("square(%d) = %d\n", i + 1, h_out[i]);
29
30    cudaFree(d_out);
31    return 0;
32}

8. __global__ vs __device__

textCopy
1__global__   :  CPU ─────────→ GPU
2__device__    :  GPU ─────────→ GPU

9. __host__

__host__ specifies a function that executes on the CPU.
Ordinary functions are host functions by default, so this qualifier is usually unnecessary.

cppCopy
1__host__ void cpu_function()
2{
3    printf("Running on CPU\n");
4}

10. The Three CUDA Function Qualifiers

Combinations such as __host__ __device__ are possible but not needed for beginners.

11. Understanding Kernel Launch

cppCopy
1kernel<<<4, 256>>>();

creates:

textCopy
1Grid
2│
3├── Block 0 → 256 threads
4├── Block 1 → 256 threads
5├── Block 2 → 256 threads
6└── Block 3 → 256 threads

Total: 4 × 256 = 1024 threads.

12. What is the Grid?

The complete collection of blocks launched for one kernel is called the grid.

cppCopy
1kernel<<<4, 256>>>();

→ Grid containing 4 blocks.

13. What is a Block?

A block is a group of threads that can cooperate via shared memory and __syncthreads().

textCopy
1Block 0
2│
3├── Thread 0
4├── Thread 1
5├── ...
6└── Thread 255

14. What is a Thread?

Every thread executes the kernel code.

Practice – print IDs:

cppCopy
1__global__ void print_ids()
2{
3    printf("Block %d, Thread %d → global %d\n",
4           blockIdx.x, threadIdx.x,
5           blockIdx.x * blockDim.x + threadIdx.x);
6}
7
8int main()
9{
10    print_ids<<<2, 4>>>();   // 8 threads total
11    cudaDeviceSynchronize();
12    return 0;
13}

15. threadIdx

cppCopy
1threadIdx.x   // for 1-D blocks

Example:

cppCopy
1__global__ void kernel()
2{
3    printf("Thread: %d\n", threadIdx.x);
4}
5
6int main()
7{
8    kernel<<<1, 4>>>();
9    cudaDeviceSynchronize();
10    return 0;
11}

16. blockIdx

cppCopy
1blockIdx.x

Example:

cppCopy
1__global__ void kernel()
2{
3    printf("Block: %d\n", blockIdx.x);
4}
5
6int main()
7{
8    kernel<<<4, 1>>>();
9    cudaDeviceSynchronize();
10    return 0;
11}

17. blockDim

blockDim.x gives the number of threads per block.

If you launch kernel<<<4, 256>>(), then blockDim.x == 256 inside every block.

18. The Most Important CUDA Formula

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;

This is the fundamental 1-D global thread index.

19. Why Do We Need Global Indexing?

With kernel<<<4, 4>>>() you have 16 threads.
The formula gives each thread a unique index from 0 to 15.

20–23. First Useful Kernel — Vector Addition (Complete + Timed)

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void vector_add(const float* A, const float* B, float* C, int N)
5{
6    int i = blockIdx.x * blockDim.x + threadIdx.x;
7    if (i < N)
8        C[i] = A[i] + B[i];
9}
10
11int main()
12{
13    const int N = 1 << 20;          // 1M elements
14    size_t bytes = N * sizeof(float);
15
16    float *h_A = new float[N];
17    float *h_B = new float[N];
18    float *h_C = new float[N];
19
20    for (int i = 0; i < N; i++) {
21        h_A[i] = 1.0f;
22        h_B[i] = 2.0f;
23    }
24
25    float *d_A, *d_B, *d_C;
26    cudaMalloc(&d_A, bytes);
27    cudaMalloc(&d_B, bytes);
28    cudaMalloc(&d_C, bytes);
29
30    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
31    cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
32
33    int threads = 256;
34    int blocks  = (N + threads - 1) / threads;
35
36    // ---- Timing ----
37    cudaEvent_t start, stop;
38    cudaEventCreate(&start);
39    cudaEventCreate(&stop);
40
41    // Warm-up
42    vector_add<<<blocks, threads>>>(d_A, d_B, d_C, N);
43    cudaDeviceSynchronize();
44
45    // Timed run
46    cudaEventRecord(start);
47    vector_add<<<blocks, threads>>>(d_A, d_B, d_C, N);
48    cudaEventRecord(stop);
49    cudaEventSynchronize(stop);
50
51    float ms = 0.0f;
52    cudaEventElapsedTime(&ms, start, stop);
53    printf("Vector add time: %.3f ms\n", ms);
54
55    cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
56
57    printf("C[0] = %.1f (should be 3.0)\n", h_C[0]);
58    printf("C[N-1] = %.1f\n", h_C[N-1]);
59
60    // Cleanup
61    cudaEventDestroy(start);
62    cudaEventDestroy(stop);
63    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
64    delete[] h_A; delete[] h_B; delete[] h_C;
65    return 0;
66}

24. What's Happening in This Program?

textCopy
1CPU
2 │
3 ├── Allocate host arrays
4 ├── Fill A and B
5 │
6 ↓
7cudaMalloc → GPU global memory
8 │
9 ↓
10cudaMemcpy H→D
11 │
12 ↓
13Launch kernel  (blocks × threads)
14 │
15 ↓
16GPU executes: each thread computes one C[i]
17 │
18 ↓
19cudaMemcpy D→H
20 │
21 ↓
22CPU prints results

25. __device__ Function Example (with timing)

cppCopy
1__device__ float square(float x)
2{
3    return x * x;
4}
5
6__global__ void square_kernel(const float* input, float* output, int N)
7{
8    int i = blockIdx.x * blockDim.x + threadIdx.x;
9    if (i < N)
10        output[i] = square(input[i]);
11}
12
13// Host code is identical to the vector-add example;
14// just replace the kernel name and launch the same way.

26. __host__ + __device__

cppCopy
1__host__ __device__ float square(float x)
2{
3    return x * x;
4}

Can be called from both host and device code (subject to restrictions).

27. Important Difference: Kernel vs Normal Function

cppCopy
1// Normal function
2float square(float x) { return x * x; }
3float y = square(x);               // direct call
4
5// CUDA kernel
6__global__ void square_kernel(...);
7square_kernel<<<blocks, threads>>>(...);   // launch

28–29. 1D / 2D / 3D Kernel Launches + 2D Indexing

2D matrix scale example (complete + timed):

cppCopy
1__global__ void matrix_scale(float* matrix, int width, int height, float factor)
2{
3    int col = blockIdx.x * blockDim.x + threadIdx.x;
4    int row = blockIdx.y * blockDim.y + threadIdx.y;
5
6    if (row < height && col < width) {
7        int idx = row * width + col;
8        matrix[idx] *= factor;
9    }
10}
11
12int main()
13{
14    const int WIDTH  = 1024;
15    const int HEIGHT = 1024;
16    size_t bytes = WIDTH * HEIGHT * sizeof(float);
17
18    float *h_mat = new float[WIDTH * HEIGHT];
19    // ... initialize h_mat ...
20
21    float *d_mat;
22    cudaMalloc(&d_mat, bytes);
23    cudaMemcpy(d_mat, h_mat, bytes, cudaMemcpyHostToDevice);
24
25    dim3 threads(16, 16);
26    dim3 blocks( (WIDTH  + 15) / 16,
27                 (HEIGHT + 15) / 16 );
28
29    cudaEvent_t start, stop;
30    cudaEventCreate(&start);
31    cudaEventCreate(&stop);
32
33    // Warm-up
34    matrix_scale<<<blocks, threads>>>(d_mat, WIDTH, HEIGHT, 2.0f);
35    cudaDeviceSynchronize();
36
37    // Timed
38    cudaEventRecord(start);
39    matrix_scale<<<blocks, threads>>>(d_mat, WIDTH, HEIGHT, 2.0f);
40    cudaEventRecord(stop);
41    cudaEventSynchronize(stop);
42
43    float ms = 0.0f;
44    cudaEventElapsedTime(&ms, start, stop);
45    printf("2-D matrix scale time: %.3f ms\n", ms);
46
47    // ... copy back, free, etc. ...
48    return 0;
49}

30. Kernel Launch Configuration

cppCopy
1// 1-D
2kernel<<<100, 256>>>(args);
3
4// Explicit dim3
5dim3 blocks(100);
6dim3 threads(256);
7kernel<<<blocks, threads>>>(args);
8
9// 2-D
10dim3 blocks(10, 10);
11dim3 threads(16, 16);
12kernel<<<blocks, threads>>>(args);

31. What Does NOT Happen?

cppCopy
1kernel<<<4, 256>>>();

does not mean “use 4 SMs”.
It means “launch 4 logical blocks”. The hardware scheduler maps blocks to SMs.

32. Kernel Launch Is Not the Same as CPU Function Call

textCopy
1CPU
2 │
3 │ enqueue kernel
4 ↓
5CUDA runtime / driver
6 │
7 ↓
8GPU scheduler
9 │
10 ↓
11SMs → Warps → Threads

33. Kernel Errors (Recommended Pattern)

cppCopy
1kernel<<<blocks, threads>>>(...);
2
3cudaError_t err = cudaGetLastError();
4if (err != cudaSuccess) {
5    printf("Launch Error: %s\n", cudaGetErrorString(err));
6}
7
8err = cudaDeviceSynchronize();
9if (err != cudaSuccess) {
10    printf("Runtime Error: %s\n", cudaGetErrorString(err));
11}

Reusable macro (recommended):

cppCopy
1#define CUDA_CHECK(call)                                                     \
2    do {                                                                     \
3        cudaError_t err = call;                                              \
4        if (err != cudaSuccess) {                                            \
5            printf("CUDA error at %s:%d – %s\n",                             \
6                   __FILE__, __LINE__, cudaGetErrorString(err));             \
7            exit(EXIT_FAILURE);                                              \
8        }                                                                    \
9    } while (0)

Usage:

cppCopy
1CUDA_CHECK(cudaMalloc(&d_ptr, bytes));
2kernel<<<blocks, threads>>>(...);
3CUDA_CHECK(cudaGetLastError());
4CUDA_CHECK(cudaDeviceSynchronize());

34. The Most Important Concepts From This Tutorial

textCopy
1CUDA Program
2     │
3     ├── Host → CPU
4     │
5     └── Device → GPU
6              │
7              └── Kernel
8                    │
9                    ↓
10                   Grid
11                    │
12             ┌──────┴──────┐
13             ↓             ↓
14          Block 0       Block 1
15             │             │
16          Threads       Threads
17             │
18           Warps
19             │
20         GPU execution

Key items to memorize:

• Function qualifiers: __host__, __device__, __global__
• Launch syntax: kernel<<<blocks, threads>>>(args);
• Global index: int i = blockIdx.x * blockDim.x + threadIdx.x;
• Always check errors + synchronize in development

35. What You Should Practice Now

Write these yourself (start with the timed versions shown above):

Beginner

1. Hello GPU
2. Print thread ID / block ID / global ID
3. Vector addition
4. Vector subtraction
5. Vector multiplication / scaling

Intermediate (element-wise)

6. ReLU
7. SiLU / Swish
8. GELU (approximate)
9. Square / Absolute value / Clamp

2-D

10. Matrix element-wise addition
11. Matrix scaling
12. Matrix transpose (naive)

Important kernel concepts (next phase)

13. Shared memory + __syncthreads()
14. Reduction (sum)
15. Tiled matrix multiplication

Suggested Practice Skeleton (copy-paste ready)

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void your_kernel(/* args */)
5{
6    int i = blockIdx.x * blockDim.x + threadIdx.x;
7    // your computation
8}
9
10int main()
11{
12    // 1. Allocate & initialize host data
13    // 2. cudaMalloc + cudaMemcpy H→D
14    // 3. Calculate blocks / threads
15    // 4. Warm-up launch
16    // 5. Timed launch with cudaEvent
17    // 6. cudaMemcpy D→H + verify
18    // 7. Cleanup
19    return 0;
20}

Once you can write and time the beginner + intermediate kernels without looking at notes, you are ready for shared memory, reductions, and the first real optimization techniques that appear in Transformer kernels (RMSNorm, Softmax, etc.).