GPU Architecture · Code example

Adds two 2D matrices two ways — a flat 1D grid and a 2D grid — producing identical results, showing that grid dimensionality is a programmer convenience over linear memory, not a hardware requirement.

// 03-2d-matrix-add.cu — the SAME 2D matrix add, done two ways:
//   (1) a 1D grid that treats the matrix as one flat array, and
//   (2) a 2D grid that hands each thread a natural (row, col).
// Both produce identical results — proof that the grid's dimensionality is a
// convenience for the programmer, not a requirement of the data or hardware.
//
//   Build:  nvcc -arch=sm_120 03-2d-matrix-add.cu -o mat2d
//   Run:    ./mat2d

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

#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)

// ---- Version 1: 1D grid ----------------------------------------------------
// The matrix is just n = width*height floats in a row. One flat index does it.
__global__ void addMatrices1D(const float* A, const float* B, float* C,
                              int n, int width) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;      // ONE flat index
  if (i < n) {
    // For a pure element-wise add you never need row/col. If the algorithm DID
    // need them, with a 1D grid you must unflatten by hand:
    //     int row = i / width;   int col = i % width;   // extra div + mod
    C[i] = A[i] + B[i];
  }
}

// ---- Version 2: 2D grid ----------------------------------------------------
// One thread per (row, col) cell; the coordinates are handed to you directly.
__global__ void addMatrices2D(const float* A, const float* B, float* C,
                              int width, int height) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;    // x → column
  int row = blockIdx.y * blockDim.y + threadIdx.y;    // y → row
  if (col < width && row < height) {
    int idx = row * width + col;                       // flatten to the SAME 1D layout
    C[idx] = A[idx] + B[idx];
  }
}

int main() {
  const int width  = 1000;              // columns (x)
  const int height = 600;               // rows    (y)  — neither a multiple of 16
  const int n      = width * height;
  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 r = 0; r < height; ++r)
    for (int c = 0; c < width; ++c) {
      hA[r * width + c] = (float)r;      // A holds the row,
      hB[r * width + c] = (float)c;      // B holds the col → C must be r + c
    }

  float *dA, *dB, *dC1, *dC2;
  CHECK(cudaMalloc(&dA, bytes));
  CHECK(cudaMalloc(&dB, bytes));
  CHECK(cudaMalloc(&dC1, bytes));        // output of the 1D kernel
  CHECK(cudaMalloc(&dC2, bytes));        // output of the 2D kernel
  CHECK(cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice));

  // ---- Launch 1: plain 1D grid (int args → dim3(x,1,1) implicitly) ----
  int threads1D = 256;
  int blocks1D  = (n + threads1D - 1) / threads1D;
  addMatrices1D<<<blocks1D, threads1D>>>(dA, dB, dC1, n, width);
  CHECK(cudaGetLastError());

  // ---- Launch 2: 2D grid via dim3 ----
  dim3 block2D(16, 16);                                  // 256 threads, 16x16
  dim3 grid2D((width  + block2D.x - 1) / block2D.x,
              (height + block2D.y - 1) / block2D.y);
  addMatrices2D<<<grid2D, block2D>>>(dA, dB, dC2, width, height);
  CHECK(cudaGetLastError());
  CHECK(cudaDeviceSynchronize());

  printf("Matrix: %d x %d = %d cells\n\n", width, height, n);
  printf("1D grid : %d blocks x %d threads          (flat, n = %d)\n",
         blocks1D, threads1D, n);
  printf("2D grid : %u x %u blocks of %u x %u threads  (row/col direct)\n\n",
         grid2D.x, grid2D.y, block2D.x, block2D.y);

  // Verify both outputs against r+c, AND that they equal each other.
  CHECK(cudaMemcpy(hC, dC1, bytes, cudaMemcpyDeviceToHost));
  int bad1 = 0;
  for (int r = 0; r < height; ++r)
    for (int c = 0; c < width; ++c)
      if (hC[r * width + c] != (float)(r + c)) bad1++;

  CHECK(cudaMemcpy(hC, dC2, bytes, cudaMemcpyDeviceToHost));
  int bad2 = 0;
  for (int r = 0; r < height; ++r)
    for (int c = 0; c < width; ++c)
      if (hC[r * width + c] != (float)(r + c)) bad2++;

  printf("1D kernel: %s (%d mismatches)\n", bad1 == 0 ? "PASS" : "FAIL", bad1);
  printf("2D kernel: %s (%d mismatches)\n", bad2 == 0 ? "PASS" : "FAIL", bad2);
  printf("=> identical results from two different grid shapes.\n");

  free(hA); free(hB); free(hC);
  cudaFree(dA); cudaFree(dB); cudaFree(dC1); cudaFree(dC2);
  return 0;
}