soobook
GPU

CUDA 101

CUDA가 무엇이고, 왜 필요하며, GPU에서 코드가 어떻게 실행되는지 정리

2026.06.25 12 min read

CUDA는 NVIDIA GPU를 범용 계산 장치로 쓰기 위한 programming platform이다. GPU에 일을 맡기는 API, C/C++ 언어 확장, compiler, runtime, driver 연동, 수학 library를 함께 묶어 제공한다.

GPU는 하드웨어이고, CUDA는 그 하드웨어에 계산을 제출하고 실행 결과를 받아오기 위한 소프트웨어 계층이다.

{
  "diagram": "cuda-steps",
  "title": "CUDA execution flow",
  "description": "host code가 GPU kernel을 제출하고 grid, block, warp 단위로 실행되는 흐름이다.",
  "width": 1080,
  "height": 680,
  "interval": 1900,
  "lanes": [
    {
      "id": "software",
      "label": "Software stack",
      "x": 52,
      "y": 44,
      "width": 976,
      "height": 150
    },
    {
      "id": "execution",
      "label": "GPU execution model",
      "x": 122,
      "y": 246,
      "width": 836,
      "height": 356
    }
  ],
  "actors": [
    {
      "id": "app",
      "kind": "host",
      "label": "Application",
      "subLabel": "host code",
      "x": 162,
      "y": 124,
      "width": 184,
      "height": 78
    },
    {
      "id": "runtime",
      "kind": "toolkit",
      "label": "CUDA Runtime",
      "subLabel": "API, library, launch",
      "x": 405,
      "y": 124,
      "width": 200,
      "height": 78
    },
    {
      "id": "driver",
      "kind": "driver",
      "label": "NVIDIA Driver",
      "subLabel": "command submission",
      "x": 662,
      "y": 124,
      "width": 208,
      "height": 78
    },
    {
      "id": "gpu",
      "kind": "gpu",
      "label": "GPU",
      "subLabel": "SMs, HBM, scheduler",
      "x": 913,
      "y": 124,
      "width": 184,
      "height": 78
    },
    {
      "id": "grid",
      "kind": "grid",
      "label": "Grid",
      "subLabel": "one kernel launch",
      "x": 540,
      "y": 312,
      "width": 730,
      "height": 72
    },
    {
      "id": "block0",
      "kind": "block",
      "label": "Block 0",
      "subLabel": "256 threads",
      "x": 315,
      "y": 436,
      "width": 172,
      "height": 82
    },
    {
      "id": "block1",
      "kind": "block",
      "label": "Block 1",
      "subLabel": "256 threads",
      "x": 540,
      "y": 436,
      "width": 172,
      "height": 82
    },
    {
      "id": "blockn",
      "kind": "block",
      "label": "Block N",
      "subLabel": "remaining work",
      "x": 765,
      "y": 436,
      "width": 172,
      "height": 82
    },
    {
      "id": "warp",
      "kind": "warp",
      "label": "Warp",
      "subLabel": "32 threads",
      "x": 430,
      "y": 555,
      "width": 188,
      "height": 64
    },
    {
      "id": "sm",
      "kind": "sm",
      "label": "SM",
      "subLabel": "schedules warps",
      "x": 650,
      "y": 555,
      "width": 188,
      "height": 64
    }
  ],
  "flows": [
    {
      "id": "app-runtime",
      "source": "app",
      "target": "runtime"
    },
    {
      "id": "runtime-driver",
      "source": "runtime",
      "target": "driver"
    },
    {
      "id": "driver-gpu",
      "source": "driver",
      "target": "gpu"
    },
    {
      "id": "gpu-grid",
      "source": "gpu",
      "target": "grid"
    },
    {
      "id": "grid-block0",
      "source": "grid",
      "target": "block0"
    },
    {
      "id": "grid-block1",
      "source": "grid",
      "target": "block1"
    },
    {
      "id": "grid-blockn",
      "source": "grid",
      "target": "blockn"
    },
    {
      "id": "block-warp",
      "source": "block0",
      "target": "warp"
    },
    {
      "id": "warp-sm",
      "source": "warp",
      "target": "sm"
    }
  ],
  "steps": [
    {
      "title": "1. Host prepares work",
      "body": "CPU 쪽 application이 device memory를 준비하고 CUDA Runtime API나 library 호출로 GPU 작업을 표현한다.",
      "activeActors": ["app", "runtime", "gpu"],
      "activeFlows": ["app-runtime"]
    },
    {
      "title": "2. Kernel launch is submitted",
      "body": "Runtime과 driver는 launch configuration, argument, command를 GPU가 이해할 수 있는 제출 단위로 넘긴다.",
      "activeActors": ["app", "runtime", "driver", "gpu", "grid"],
      "activeFlows": ["app-runtime", "runtime-driver", "driver-gpu", "gpu-grid"]
    },
    {
      "title": "3. Grid becomes thread blocks",
      "body": "하나의 kernel launch는 grid를 만들고, grid 안의 block들은 가능한 SM에 순차적으로 배치된다.",
      "activeActors": ["grid", "block0", "block1", "blockn", "gpu"],
      "activeFlows": ["gpu-grid", "grid-block0", "grid-block1", "grid-blockn"]
    },
    {
      "title": "4. SM schedules warps",
      "body": "Block 안의 thread는 warp 단위로 묶이고, SM의 warp scheduler가 실행 가능한 warp를 번갈아 실행한다.",
      "activeActors": ["block0", "warp", "sm"],
      "activeFlows": ["block-warp", "warp-sm"]
    },
    {
      "title": "5. Host observes completion",
      "body": "동기화나 device-to-host copy가 발생하면 host는 GPU 작업 완료와 결과를 관찰할 수 있다.",
      "activeActors": ["gpu", "driver", "runtime", "app"],
      "activeFlows": ["__none"]
    }
  ]
}

What CUDA Is

CUDA는 Compute Unified Device Architecture의 약자다. NVIDIA가 만든 parallel computing platform이자 programming model이다.

일반적인 CPU 프로그램은 CPU core에서 실행된다. CUDA 프로그램은 CPU에서 시작하지만, 계산 일부를 GPU로 보낸다. CPU 쪽 코드를 host code, GPU 쪽에서 실행되는 코드를 device code라고 부른다.

CUDA가 제공하는 것은 크게 네 가지다.

  • GPU에서 실행할 함수를 쓰는 방법
  • CPU code가 GPU 작업을 제출하는 API
  • GPU memory를 할당하고 복사하는 API
  • cuBLAS, cuDNN, NCCL 같은 고성능 library

딥러닝 프레임워크를 쓸 때도 이 계층은 숨어 있다. PyTorch에서 tensor.to("cuda")를 호출하거나 GPU tensor 연산을 실행하면 내부에서는 CUDA runtime, CUDA library, NVIDIA driver가 함께 움직인다.

Why CUDA Exists

GPU는 원래 그래픽 처리를 위해 발전했다. 같은 종류의 계산을 많은 데이터에 반복해서 적용하는 데 강하다. 픽셀, vertex, matrix, tensor처럼 모양이 비슷한 데이터가 많을수록 GPU가 잘 맞는다.

딥러닝 학습과 추론도 이런 형태다. 큰 matrix multiplication, convolution, attention 연산은 많은 원소에 비슷한 계산을 반복한다. CPU가 순서가 복잡한 작업에 강하다면, GPU는 같은 모양의 계산을 대량으로 밀어 넣는 데 강하다.

CUDA는 이 GPU를 그래픽 API 없이 일반 계산에 쓰기 위해 만들어졌다.

CUDA가 없으면 개발자는 GPU를 직접 다루기 어렵다. GPU memory를 어떻게 잡을지, CPU memory와 GPU memory 사이를 어떻게 복사할지, 어떤 함수를 몇 개의 thread로 실행할지, 실행이 끝났는지 어떻게 알지 모두 직접 맞춰야 한다.

CUDA는 이 작업들을 programming model로 정리한다. 개발자는 GPU에서 실행할 kernel을 쓰고, CUDA runtime이나 library를 통해 GPU에 작업을 제출한다.

CUDA Is Not The GPU

CUDA를 이해할 때 가장 먼저 분리할 것은 hardware와 software다.

GPU는 물리 장치다. SM, CUDA Core, Tensor Core, HBM, L2 cache 같은 부품이 GPU 안에 있다.

CUDA는 그 장치를 쓰는 방법이다. CUDA C++, CUDA Runtime API, CUDA Driver API, CUDA compiler, CUDA library가 여기에 들어간다.

NVIDIA driver는 operating system과 GPU 사이에 있다. CUDA runtime은 application이 쓰기 쉬운 API를 제공하고, driver API는 더 낮은 수준의 제어를 제공한다. NVIDIA 문서에서는 runtime API가 driver API 위에 만들어져 있다고 설명한다.

대략 이런 층으로 볼 수 있다.

Application
  PyTorch, TensorFlow, custom CUDA C++ program

CUDA Toolkit
  compiler, runtime API, libraries

NVIDIA Driver
  GPU command submission, device management

GPU
  SM, warp scheduler, memory, execution units

대부분의 개발자는 CUDA runtime이나 PyTorch 같은 framework를 통해 CUDA를 만난다. 하지만 문제가 생기면 어느 층에서 막혔는지 나눠 봐야 한다.

Toolkit

CUDA Toolkit은 CUDA 개발에 필요한 도구 묶음이다.

nvcc는 CUDA C++ code를 compile하는 compiler driver다. GPU에서 실행될 device code와 CPU에서 실행될 host code를 함께 다룬다.

CUDA Runtime API는 memory allocation, memory copy, kernel launch, stream, event 같은 기능을 제공한다.

CUDA library는 자주 쓰는 고성능 연산을 미리 구현해둔 것이다. 예를 들어 matrix multiplication은 직접 kernel을 짤 수도 있지만, 실제 딥러닝에서는 cuBLAS나 cuDNN 같은 library가 핵심 경로에 들어가는 경우가 많다.

그래서 CUDA를 배울 때는 직접 kernel을 쓰는 execution model과 framework가 CUDA library를 호출하는 경로를 함께 봐야 한다.

Kernel

CUDA에서 kernel은 GPU에서 실행되는 함수다.

CPU에서 함수를 호출하면 현재 CPU thread가 그 함수를 실행한다. CUDA kernel은 다르다. CPU 쪽 프로그램이 kernel launch를 요청하면 GPU가 많은 CUDA thread를 만들어 같은 kernel code를 실행한다.

간단한 vector addition을 보자.

__global__ void add(float* x, float* y, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = x[i] + y[i];
    }
}

이 함수는 원소 하나를 더하는 코드처럼 보인다. 실제 실행에서는 많은 thread가 동시에 같은 함수를 실행하고, 각 thread가 자기 index에 해당하는 원소를 처리한다.

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

add<<<blocks, threads>>>(x, y, out, n);

<<<blocks, threads>>>는 이 kernel을 몇 개의 block과 몇 개의 thread로 실행할지 정하는 launch configuration이다.

Grid and Block

CUDA kernel을 launch하면 하나의 grid가 생긴다. grid 안에는 여러 thread block이 있고, block 안에는 여러 thread가 있다.

kernel launch
└── grid
    ├── block 0
    │   ├── thread 0
    │   ├── thread 1
    │   └── ...
    ├── block 1
    └── ...

같은 block 안의 thread들은 서로 협력할 수 있다. shared memory를 함께 쓰고, __syncthreads()로 같은 block 안의 thread들을 기다리게 만들 수 있다.

서로 다른 block은 기본적으로 독립적이어야 한다. CUDA는 block들이 어떤 순서로 SM에 배치될지 보장하지 않는다. 같은 kernel 안에서 block A가 block B의 결과를 기다리는 식으로 짜면 기본 모델에 맞지 않는다.

각 CUDA thread는 자기가 처리할 데이터 위치를 직접 계산한다.

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

threadIdx.x는 block 안에서의 thread 번호다. blockIdx.x는 grid 안에서의 block 번호다. blockDim.x는 block 하나에 들어 있는 thread 수다.

Warp and SM

thread는 block 안에 들어 있지만, 실제 실행은 warp 단위로 움직인다.

NVIDIA GPU에서 warp 하나는 32개 thread다. CUDA의 실행 방식은 SIMT라고 부른다. Single Instruction, Multiple Threads의 약자다. 같은 kernel code를 여러 thread가 실행하지만, 각 thread는 자기 data와 자기 index를 가진다.

정리 한 번 하고 넘어가자.

  • Thread: CUDA Kernel 을 실행하는 가장 작은 논리단위.
  • Warp: NVIDIA GPU 가 Thread 를 실행할 때 묶는 단위. 보통 32개 Thread.
  • Thread Block: 여러 Thread 의 묶음, 여러 Warp 로 나뉨.
  • SM: Block 과 Warp 를 실제로 실행하는 GPU 안의 Hardware unit.

예를 들어 block 하나가 256 Thread 라면:

1 block = 256 threads = 8 warps\textbf{\text{1 block = 256 threads = 8 warps}}

이 block 은 하나의 SM 에 배치된다. block 하나가 여러 SM 에 쪼개져 들어가지는 않는다.

그리고 SM 은 그 block 안의 warp 들을 실행한다. 한 SM 에는 resource 가 허용하는 만큼 여러 block 이 동시에 올라갈 수 있고, 그러면 SM 안에는 여러 block 에서 나온 여러 warp 가 대기한다.

SM 의 warp scheduler 가 이 warp 들을 번갈아 실행한다.

Thread 32개 = Warp 1개
Warp 여러 개 = Thread Block 1개
Thread Block은 SM 하나에 배치됨
SM은 여러 block/warp를 올려두고 warp 단위로 실행함

조건문이 있으면 같은 warp 안에서도 어떤 thread는 if 안으로 들어가고, 어떤 thread는 들어가지 않을 수 있다. 이런 상황을 warp divergence라고 부른다.

if (i % 2 == 0) {
    out[i] = x[i] * 2;
}

같은 warp 안에서 일부 thread만 활성화되면 나머지 lane은 잠시 mask된다. 그래서 같은 warp 안의 thread들이 비슷한 control flow를 타는 편이 효율적이다.

GPU에는 여러 SM이 있다. kernel launch로 만들어진 block들은 SM에 배치된다. 일반적인 CUDA kernel에서 block 하나는 하나의 SM에서 실행된다. 한 SM은 register, shared memory, 최대 thread 수 같은 resource가 허용하는 만큼 여러 block을 동시에 올려둘 수 있다.

CUDA는 block 실행 순서를 프로그래머에게 보장하지 않는다.

grid has 1000 blocks
GPU has 80 SMs

처음에는 일부 block만 SM에 올라간다.
block이 끝나면 남은 block이 빈 SM에 올라간다.
이 과정을 모든 block이 끝날 때까지 반복한다.

Memory

CUDA 프로그램은 CPU memory와 GPU memory를 구분해서 봐야 한다.

CPU process가 가진 일반 RAM은 host memory다. GPU가 직접 계산에 쓰는 memory는 device memory다. PyTorch에서는 CPU tensor와 CUDA tensor가 나뉘어 보이는 이유도 여기에 있다.

CUDA C++에서는 보통 이런 흐름이 나온다.

cudaMalloc(&x_gpu, size);
cudaMemcpy(x_gpu, x_cpu, size, cudaMemcpyHostToDevice);

add<<<blocks, threads>>>(x_gpu, y_gpu, out_gpu, n);

cudaMemcpy(out_cpu, out_gpu, size, cudaMemcpyDeviceToHost);

계산이 GPU에서 빨라도, CPU memory와 GPU memory 사이를 자주 오가면 전체 시간은 느려질 수 있다. 그래서 GPU code를 볼 때는 “계산을 GPU에서 하는가”뿐 아니라 “데이터가 이미 GPU memory에 있는가”를 같이 봐야 한다.

PyTorch에서 x.to("cuda")는 주로 데이터 이동이고, model(x)는 내부적으로 여러 CUDA kernel이나 CUDA library call을 실행하는 계산 흐름이다.

Synchronization

CUDA 작업은 host 입장에서 비동기로 진행되는 경우가 많다. CPU가 kernel launch를 요청했다고 해서 CPU가 항상 kernel 완료까지 기다리는 것은 아니다.

add<<<blocks, threads>>>(x, y, out, n);

// host code may continue here

CPU는 GPU에 일을 제출하고 다음 일을 준비할 수 있다. CUDA stream을 쓰면 data copy와 kernel 실행을 겹치게 만들 수도 있다.

결과를 CPU에서 바로 읽어야 한다면 어느 지점에서는 기다려야 한다. cudaDeviceSynchronize(), stream synchronization, device-to-host copy 같은 동작이 그 경계가 된다.

block 안의 synchronization과 host-GPU synchronization도 구분해야 한다.

__syncthreads()는 같은 block 안의 thread들을 맞추는 장치다. GPU 전체나 CPU와 GPU 사이를 맞추는 함수가 아니다.

__shared__ float tile[256];

tile[threadIdx.x] = x[i];
__syncthreads();

int next = (threadIdx.x + 1) % blockDim.x;
out[i] = tile[next] * 2;

이 코드에서는 각 thread가 shared memory에 값을 쓴 뒤, 옆 thread가 쓴 값을 읽는다. 읽기 전에 같은 block의 thread들이 쓰기를 끝냈는지 맞춰야 한다.

Frameworks

대부분의 ML engineer는 CUDA kernel을 직접 쓰지 않는다. PyTorch, TensorFlow, JAX 같은 framework가 CUDA를 호출한다.

예를 들어 PyTorch에서 다음 코드를 실행한다고 하자.

import torch

x = torch.randn(1024, 1024, device="cuda")
y = x @ x

이 코드는 Python 한 줄처럼 보이지만 내부에서는 CUDA memory allocation, CUDA library call, kernel launch, stream scheduling이 일어난다. matrix multiplication은 직접 작성한 kernel이 아니라 cuBLAS 같은 library 경로를 탈 수 있다.

그래서 CUDA를 이해하면 framework를 버리고 저수준 code를 쓰자는 뜻이 아니다. GPU memory usage, asynchronous execution, synchronization, kernel launch overhead 같은 현상을 읽는 데 도움이 된다.

CUDA의 네 층

CUDA를 읽을 때는 네 층으로 나누면 된다.

Application은 PyTorch나 CUDA C++ 프로그램이다. 여기서 GPU tensor를 만들거나 kernel을 launch한다.

CUDA Toolkit은 runtime API, compiler, library를 제공한다. 개발자는 이 계층을 통해 GPU 작업을 표현한다.

NVIDIA driver는 OS와 GPU 사이에서 command submission과 device 관리를 맡는다.

GPU는 제출된 kernel을 grid, block, thread, warp 단위로 실행하고, device memory를 읽고 쓴다.

References