Hello CUDA: Profile Your First Kernel

Start with a small vector-add kernel: the CUDA equivalent of a practical Hello World. Edit the kernel body, choose profiling passes, and use the launch configuration below the kernel to experiment with grid and block dimensions. The host harness supplies A and B, launches your kernel for several input sizes, and verifies that every C[i] equals A[i] + B[i].

Languages: CUDA C++, Rust

Lesson

GPUFlight compiles this editable kernel and runs the generated CUDA program under gpufl trace. Trace records the activity timeline and launch metadata. RangeProfilerKernelReplay adds per-kernel hardware counters such as occupancy, SM throughput, and cache behavior. SassMetrics attributes supported execution and memory metrics to SASS instructions. PcSampling statistically samples program counters and warp stall reasons. Trace is required as the anchor; the other selected passes are linked into the same analysis.

PC sampling reports samples only after a few thousand kernel launches, and this harness launches each input size once. So when PcSampling is selected, its pass relaunches the 16,777,216-element case up to 4,096 times while the other passes still launch each case once. That pass therefore shows more launches and GPU time than Trace, and the analysis marks the passes as non-deterministic for that reason.

Run the starter with Trace first. Add the other passes to answer specific profiling questions, noting their different overheads. Then edit the launch configuration below the kernel: try a block x dimension of 64, 128, or 256 threads while keeping the grid's x dimension derived from N and the block size. Compare launch metadata, GPU time, and occupancy. You can also rewrite the body as long as it still produces the expected vector sum. Compare with a Trace-only run and observe which deeper profiler tabs no longer have data.

Start with Trace to locate work and transfers. Add a metric pass only for a specific question; replay/capture work can add runtime and credit cost.

What you will learn

Run editable CUDA code through GPUFlight and open the resulting analysis.
Explain what Trace, RangeProfilerKernelReplay, SassMetrics, and PcSampling collect.
Edit the launch configuration while preserving complete input coverage.
Compare two runs after changing one code or launch variable.

Starter code

    // One thread handles one vector element. Try another correct body, or keep
    // this code and edit the launch-configuration function below to see how
    // launch shape affects metrics.
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        float sum = A[i] + B[i];
        C[i] = sum;
    }