GPU Architecture · Code example

Queries the GPU and prints SMs, cores, register file, caches, and memory bandwidth — confirming the lesson numbers on your own hardware — then runs a first vector-add kernel.

// 01-device-query.cu — prove the toolchain works and see Lessons 1–5 on your real GPU.
//
//   Build:  nvcc -arch=sm_120 01-device-query.cu -o devq
//   Run:    ./devq
//
// (sm_120 = Blackwell / compute capability 12.0. nvcc 13.3 also defaults to it.)

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

// Minimal error check: every CUDA call can fail; never ignore the return code.
#define CHECK(call) do {                                              \
    cudaError_t _e = (call);                                          \
    if (_e != cudaSuccess) {                                          \
      printf("CUDA error at line %d: %s\n", __LINE__,                 \
             cudaGetErrorString(_e));                                 \
      return 1;                                                       \
    }                                                                 \
  } while (0)

// FP32 lanes per SM by compute capability (Lesson 2). Blackwell/Hopper = 128.
static int coresPerSM(int major, int minor) {
  if (major == 12) return 128;                 // Blackwell
  if (major == 9)  return 128;                 // Hopper
  if (major == 8)  return (minor == 0) ? 64 : 128;  // Ampere GA100 vs GA10x/Ada
  return 64;                                   // conservative fallback
}

// A trivial kernel. Each thread computes its unique global index (Lesson 3) and
// does one add. The `if (i < n)` guard matters because the grid rounds up past n.
__global__ void vadd(const float* a, const float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;   // the single most-used line in CUDA
  if (i < n) c[i] = a[i] + b[i];
}

int main() {
  cudaDeviceProp p;
  CHECK(cudaGetDeviceProperties(&p, 0));

  int cores = coresPerSM(p.major, p.minor);

  // CUDA 13 removed memoryClockRate / memoryBusWidth from cudaDeviceProp;
  // query them via cudaDeviceGetAttribute instead.
  int memClockKHz = 0, busWidth = 0;
  CHECK(cudaDeviceGetAttribute(&memClockKHz, cudaDevAttrMemoryClockRate, 0));
  CHECK(cudaDeviceGetAttribute(&busWidth,   cudaDevAttrGlobalMemoryBusWidth, 0));
  // Standard deviceQuery bandwidth estimate: 2 (DDR) * clock * (bus bytes).
  double bwGBs = 2.0 * memClockKHz * (busWidth / 8) / 1.0e6;

  printf("== %s  (compute capability %d.%d) ==\n\n", p.name, p.major, p.minor);
  printf("  SMs                     : %d\n", p.multiProcessorCount);
  printf("  CUDA cores (estimated)  : %d   (%d SMs x %d lanes)\n",
         p.multiProcessorCount * cores, p.multiProcessorCount, cores);
  printf("  Warp size               : %d threads\n", p.warpSize);
  printf("  Max threads / block     : %d\n", p.maxThreadsPerBlock);
  printf("  Registers / SM          : %d (32-bit)\n", p.regsPerMultiprocessor);
  printf("  Registers / block (max) : %d\n", p.regsPerBlock);
  printf("  Shared mem / SM         : %zu KB\n", p.sharedMemPerMultiprocessor / 1024);
  printf("  Shared mem / block      : %zu KB\n", p.sharedMemPerBlock / 1024);
  printf("  L2 cache                : %d MB\n", p.l2CacheSize / (1024 * 1024));
  printf("  Global memory (DRAM)    : %.1f GB\n", p.totalGlobalMem / 1.0e9);
  printf("  Memory bus width        : %d-bit\n", busWidth);
  printf("  Memory bandwidth (est.) : %.0f GB/s\n", bwGBs);

  // --- Prove kernel execution end to end (host<->global movement + a launch). ---
  const int n = 1 << 20;                       // 1,048,576 elements
  size_t bytes = (size_t)n * sizeof(float);

  float *ha = (float*)malloc(bytes);
  float *hb = (float*)malloc(bytes);
  float *hc = (float*)malloc(bytes);
  for (int i = 0; i < n; ++i) { ha[i] = 1.0f; hb[i] = 2.0f; }

  float *da, *db, *dc;
  CHECK(cudaMalloc(&da, bytes));
  CHECK(cudaMalloc(&db, bytes));
  CHECK(cudaMalloc(&dc, bytes));
  CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));   // host -> global
  CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));

  int threads = 256;                           // block shape (Lesson 3)
  int blocks  = (n + threads - 1) / threads;   // grid shape, rounded up to cover n
  vadd<<<blocks, threads>>>(da, db, dc, n);
  CHECK(cudaGetLastError());                    // catch a bad launch configuration
  CHECK(cudaDeviceSynchronize());               // wait, and surface any runtime error

  CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));   // global -> host
  printf("\n  Kernel check: c[0]=%.1f  c[%d]=%.1f   (expected 3.0)\n",
         hc[0], n - 1, hc[n - 1]);

  free(ha); free(hb); free(hc);
  cudaFree(da); cudaFree(db); cudaFree(dc);
  return 0;
}