Saturday, July 18, 2026

CUDA Hello World - Gain Experience of 10 Years

If you already know the fundamentals from your earlier work with Nsight Eclipse. A lot has changed since then, and your current setup is far better going by current standards.

The IDE situation in 2026

About 10 years back, you needed Nsight Eclipse because it was the only CUDA-aware IDE. That's no longer true. VS Code is now the go-to for CUDA development — and you already have it installed 

Nsight Eclipse is Linux-only and essentially deprecated for Windows. On Windows, NVIDIA moved to Nsight Visual Studio Edition (integrates into Visual Studio) or the standalone Nsight Systems/Compute profilers — which you already have Nsight Compute installed.

Open a terminal and run:

cmd
nvcc --version

Already confirmed ✓. Now try compiling a tiny .cu file:

cuda
// hello.cu
#include <stdio.h>

__global__ void hello() {
    printf("Hello from thread %d, block %d\n", threadIdx.x, blockIdx.x);
}

int main() {
    hello<<<2, 4>>>();
    cudaDeviceSynchronize();
    return 0;
}

Save it anywhere and compile with:

cmd
    nvcc hello.cu -o hello.exe
    hello.exe

If you see 8 lines printing (2 blocks × 4 threads), your entire chain — driver, CUDA toolkit, nvcc, and GPU — is working end to end. That's your "Hello World" moment, much more meaningful than the bit-reverse template because you'll understand exactly why 8 lines print. BIT REVERSE TEMPLATE DEMO


What's different from 2016

2016 (your 410M)2026 (your GTX 1650)
CUDA 5.5, CC 2.1        CUDA 12.2, CC 7.5
Nsight Eclipse on Ubuntu        VS Code + nvcc on Windows
No Tensor Cores        Tensor Cores available
~192 CUDA cores        896 CUDA cores
No printf in kernels easily     printf from device works natively

The printf inside the kernel above would not have worked cleanly on your old 410M — it's a nice demonstration that your new hardware is a proper leap forward.

C:\Users\usrname\OneDrive\Documents\GitHub\LearningC>nvcc helloWorld.cu -o helloWorld.exe

helloWorld.cu

tmpxft_00001d74_00000000-10_helloWorld.cudafe1.cpp

   Creating library helloWorld.lib and object helloWorld.exp


C:\Users\usrname\OneDrive\Documents\GitHub\LearningC>helloWorld.exe

Hello from thread 0, block 0

Hello from thread 1, block 0

Hello from thread 2, block 0

Hello from thread 3, block 0

Hello from thread 0, block 1

Hello from thread 1, block 1

Hello from thread 2, block 1

Hello from thread 3, block 1


Let's Go!!

got 8 hello messages

That's your GPU running code you wrote. 8 lines = 2 blocks × 4 threads, all executing in parallel on the GTX 1650.

Notice what just happened in that launch syntax:

cuda
hello<<<2, 4>>>();

Those angle brackets are CUDA's extension to C++ — not standard C, not anything your MinGW or MSVC understands alone. Only nvcc parses them. The two numbers are:

  • 2 → number of blocks in the grid
  • 4 → number of threads per block

So total threads = 2 × 4 = 8. Each thread ran the same hello() function, but with a different threadIdx.x and blockIdx.x — that's how every thread knows who it is and what data to work on.

Now let's make it real — Lesson 2: Vector Addition

This is the "Hello World" of parallel computing that actually means something. You'll add two arrays of 1 million numbers simultaneously on the GPU.

This exercise teaches you the four fundamental steps every CUDA program ever written follows:

1. Allocate memory on GPU        (cudaMalloc)
2. Copy data CPU → GPU           (cudaMemcpy)
3. Launch kernel                 (kernel<<<blocks, threads>>>)
4. Copy result GPU → CPU         (cudaMemcpy back)

Create a new file vector_add.cu and type this out (don't copy-paste — typing it builds muscle memory):

#include <stdio.h>

#include <stdlib.h>


// DEVICE function — runs on GPU

// Each thread adds ONE pair of elements

__global__ void vectorAdd(float *a, float *b, float *c, int n) {

    // Calculate this thread's unique index

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


    // Guard: don't go out of bounds

    if (i < n) {

        c[i] = a[i] + b[i];

    }

}


int main() {

    int n = 1 << 20;          // 1,048,576 elements (1M)

    size_t bytes = n * sizeof(float);


    // --- HOST memory (CPU RAM) ---

    float *h_a = (float*)malloc(bytes);

    float *h_b = (float*)malloc(bytes);

    float *h_c = (float*)malloc(bytes);


    // Fill input arrays

    for (int i = 0; i < n; i++) {

        h_a[i] = (float)i;

        h_b[i] = (float)(i * 2);

    }


    // --- DEVICE memory (GPU VRAM) ---

    float *d_a, *d_b, *d_c;

    cudaMalloc(&d_a, bytes);

    cudaMalloc(&d_b, bytes);

    cudaMalloc(&d_c, bytes);


    // Step 2: Copy CPU → GPU

    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);

    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);


    // Step 3: Launch kernel

    int threads = 256;

    int blocks  = (n + threads - 1) / threads;   // ceiling division

    vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);


    // Step 4: Copy GPU → CPU

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);


    // Verify first 5 and last 5 results

    printf("First 5 results:\n");

    for (int i = 0; i < 5; i++)

        printf("  c[%d] = %.0f  (expected %.0f)\n", i, h_c[i], h_a[i] + h_b[i]);


    printf("Last 5 results:\n");

    for (int i = n-5; i < n; i++)

        printf("  c[%d] = %.0f  (expected %.0f)\n", i, h_c[i], h_a[i] + h_b[i]);


    // Cleanup

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

    free(h_a); free(h_b); free(h_c);


    printf("\nDone. %d elements added in parallel.\n", n);

    return 0;

}

Compile and run:

cmd
nvcc vector_add.cu -o vector_add.exe
vector_add.exe

The one line to understand deeply

cuda
int i = blockIdx.x * blockDim.x + threadIdx.x;

This is the most important line in all of CUDA. Every thread computes its own unique global index from three built-in variables:

VariableMeaningIn your launch
blockIdx.xWhich block am I in?0 to 4095
blockDim.xHow many threads per block?256
threadIdx.xWhich thread within my block?0 to 255

So thread 0 of block 3 gets index 3 * 256 + 0 = 768 — it adds a[768] + b[768]. Thread 1 of block 3 gets index 769. And so on. Every thread handles exactly one element, all simultaneously.

Run it and share the output — then we'll talk about why blocks = (n + threads - 1) / threads and what happens if you skip the if (i < n) guard.

Output:

C:\Users\usrname\OneDrive\Documents\GitHub\LearningC>nvcc vector_add.cu -o vector_add.exe

vector_add.cu

tmpxft_00003078_00000000-10_vector_add.cudafe1.cpp

   Creating library vector_add.lib and object vector_add.exp

C:\Users\usrname\OneDrive\Documents\GitHub\LearningC>vector_add.exe

First 5 results:

  c[0] = 0  (expected 0)

  c[1] = 3  (expected 3)

  c[2] = 6  (expected 6)

  c[3] = 9  (expected 9)

  c[4] = 12  (expected 12)

Last 5 results:

  c[1048571] = 3145713  (expected 3145713)

  c[1048572] = 3145716  (expected 3145716)

  c[1048573] = 3145719  (expected 3145719)

  c[1048574] = 3145722  (expected 3145722)

  c[1048575] = 3145725  (expected 3145725)

Done. 1048576 elements added in parallel.