Video summary

Coding on NVIDIA GPUs with CUDA C

Main summary

Key takeaways

Technology

Overview

The video is a beginner walkthrough for writing and running CUDA C code on an NVIDIA GPU (specifically mentioning an RTX 3060) rather than only running programs on the CPU. The goal is to start with a “hello world” that runs on the GPU, then build up to simple parallel computations.

Core Concepts and Workflow (CPU vs GPU)

Where code runs

  • By default, normal C programs execute on the CPU.
  • CUDA enables executing code on the GPU.

Why GPUs are good at parallel math

  • CPU:
    • Fewer cores
    • More cache/control
    • Better for branching and general logic
  • GPU:
    • Thousands of cores
    • Less cache
    • Poor for heavy branching
    • Excellent for massively parallel simple operations (e.g., vector/matrix math)

GPU execution model

  • The host (CPU) launches a kernel.
  • The kernel runs on the GPU as part of a hierarchy:
    • grid → blocks → threads
  • Threads use identifiers such as:
    • threadIdx.x (in a 1D example)
    • to decide which element of work each thread handles.

Product/Technology Setup and Tools

CUDA and the toolchain

  • Uses the NVIDIA CUDA API and the CUDA toolchain.
  • Key facts:
    • CUDA is proprietary.
    • CUDA stands for Compute Unified Device Architecture.
    • It is the de facto standard for GPU programming (with AMD as an alternative mentioned).

Compilation

  • CUDA files are compiled with nvcc (not GCC/Clang).
  • A practical requirement:
    • CUDA syntax is recognized only when the file is named .cu (not regular .c).

Tutorial: “Hello World” Executed on the GPU

Steps in the tutorial

  • Start with a CPU “hello world,” then switch to CUDA.
  • Kernel functions are marked with __global__ (CUDA keyword for kernels launched from the host).
  • Device code can print using CUDA printf (not the standard C printf).

Synchronization issue and fix

  • Initial confusion: the program may compile but print nothing until synchronization is added.
  • Important behavior:
    • Kernel launches are non-blocking
    • the host program can exit before device work completes
  • Fix:
    • use cudaDeviceSynchronize()

Verifying execution with thread info

  • The tutorial prints thread identifiers (via threadIdx.x) to confirm multiple GPU threads are running.

Tutorial: Simple Scalar “GPU Add”

What it demonstrates

  • Implements a CUDA kernel that computes a value (e.g., 5 + 6) and writes it through a pointer.
    • Kernels usually use void return semantics.
    • “Return values” are written to memory passed into the kernel.

Synchronization requirement

  • The result must be synchronized before the host reads/prints it.

Limitation

  • Running the same scalar operation across multiple threads (all writing the same output) is not useful for real computation—but it verifies GPU execution.

Tutorial: Vector Increment and Parallel Indexing

Memory allocation and setup

  • Uses CUDA memory allocation:
    • cudaMallocManaged, described as unified memory (accessible by CPU and GPU through a shared address space).
  • Initializes the input vector on the CPU with random values using rand() / srand(time(NULL)).
  • Uses device memory for input/output (managed/unified memory in this approach).

Parallel “increment each element by a constant”

  • Launch configuration: one block with multiple threads (e.g., 5 threads in the demo).
  • Each thread updates one vector element using its index:
    • result[threadIdx.x] = input[threadIdx.x] + constant

Core takeaway

  • Parallelism works because threads map to independent indices.

Tutorial Extension: Composed GPU Kernels (Increment then Sqrt)

Second kernel/function

  • Adds another kernel/function to compute square root per element.

Highlights

  • Initially avoids in-place modification by using separate input/output buffers, then later experiments with writing back to the same vector.
  • Adds a bounds check:
    • threadIdx.x < n
  • Mentions math.h usage and minor editing mistakes (semicolon issues).

Scaling workload and observing performance

  • Demonstrates heavier workloads (e.g., 50,000 square roots at once).
  • Observes:
    • printf is a bottleneck
    • GPU computation may be fast, but frequent device prints slow the program dramatically
  • Suggests removing prints for speed.

Key Notes / Best Practices

  • Synchronization is required (cudaDeviceSynchronize) before reading kernel results on the host.
  • CUDA kernels typically use:
    • pointers and void return
    • results written to memory passed as arguments
  • File extension matters:
    • use .cu so NVCC recognizes CUDA code
  • Device debug printing is slow:
    • printf in device code is supported, but it harms performance
  • Memory transfers/loading concept:
    • CPU memory and GPU memory are separate
    • work requires copying/loading data (conceptually explained via PCIe and device memory)
    • unified memory simplifies addressing in the tutorial

Main Speakers / Sources

  • Speaker: The video creator/host (single-person walkthrough; no other named co-speakers).
  • Source referenced for code/keywords: NVIDIA CUDA documentation / CUDA Programming Guide.

Original video