TL;DR

  • CUDA는 그래픽에서 비롯된 처리량 중심의 병렬 컴퓨팅을 범용 문제에 적용하며, 배열 원소마다 하나의 작업을 할당하는 방식으로 GPU의 강점을 활용함.
  • CPU는 소수의 명령어 흐름을 빠르게 처리하도록 캐시, 분기 예측, 비순차 실행 등을 갖추며, GPU는 단일 스레드의 속도보다 수많은 경량 스레드의 동시 실행에 트랜지스터를 배분함.
  • SIMT(단일 명령 다중 스레드) 모델에서는 커널이 논리적 스레드 하나의 작업을 표현하고, 각 스레드는 고유 인덱스를 이용해 서로 다른 데이터에 같은 연산을 수행함.
  • CUDA 실행은 그리드, 블록, 스레드로 구성되며, 블록당 최대 스레드는 1,024개이고 올림 나눗셈과 인덱스 경계 검사로 임의 길이 배열을 처리함.
  • 완전한 CUDA 프로그램은 호스트와 디바이스 사이에서 메모리를 할당·복사하고 커널을 실행한 뒤 결과를 복사하며, 하드웨어의 스레드 스케줄링과 성능 최적화는 다음 편의 주제임.

두 가지 트랜지스터 활용 방식

  • 그래픽 처리 장치(GPU)는 그래픽 연산에서 출발했으며, 그래픽은 일반 목적 연산과 다른 계산 문제였음. 중앙 처리 장치(CPU)가 문서 열기, 파일 저장, 메시지 전송처럼 낮은 지연 시간이 중요한 작업에 강한 반면, 그래픽은 많은 독립 요소를 한꺼번에 처리하는 데 초점을 둠.
  • 4K(울트라 HD) 영상의 흑백 필터 적용에는 프레임당 3,840 × 2,160, 약 830만 픽셀을 처리해야 함. 각 픽셀에 거의 같은 행렬 변환을 적용하므로 단일 입력의 처리 속도보다 다음 프레임 전에 전체 픽셀 배열을 처리하는 처리량이 중요함.
  • 처리량 최적화는 그래픽 외에도 자연어 처리, 자동 음성 인식, 컴퓨터 비전, 의료 영상, 물리 시뮬레이션, 시장 모델링, 신약 개발, 이미지·영상 편집 등 대규모 병렬 작업에 적용됨.
  • 하드웨어 설계에서는 지연 시간 최소화와 처리량 극대화 사이의 선택이 불가피함. 실리콘 면적이 한정돼 있어 한 가지 지표를 최적화하면 다른 측면의 자원이 줄어듦.
  • CPU 코어는 상대적으로 적고 큼. 단일 명령어 흐름에서 명령어 간 의존성이나 메모리 지연으로 생기는 정지를 줄이기 위해 캐시, 분기 예측, 비순차 실행, 레지스터 이름 변경, 정교한 명령어 스케줄링 등에 실리콘을 사용함.
  • GPU 코어는 더 작고 수가 많음. 하나의 스레드 내부에서 병렬성을 찾기보다 수많은 스레드를 실행하는 방식으로, 각 스레드가 배열의 서로 다른 원소를 처리하도록 함.
  • 원소별 연산이 서로 독립적이면 어떤 순서로든, 심지어 동시에 실행할 수 있음. 이러한 속성인 데이터 병렬성을 찾아내는 일이 GPU 가속 작업의 주요 개념적 과정임.
  • 이후 설명은 NVIDIA GPU와 CUDA를 바탕으로 하며, CUDA는 C/C++ 함수처럼 보이는 코드를 수백만 개의 경량 GPU 스레드에서 실행하는 NVIDIA의 범용 GPU 컴퓨팅 플랫폼임.

CUDA와 SIMT 패러다임

  • CUDA는 Compute Unified Device Architecture의 약자이며, 프로그래밍 언어 하나가 아니라 프로그래밍 모델, C/C++ 및 다른 언어의 확장, 런타임, 컴파일러, 드라이버 API, 라이브러리와 프레임워크를 포함하는 구성임.
  • CPU의 벡터 덧셈은 반복문 하나로 배열을 순회하며 각 인덱스에서 C[i] = A[i] + B[i]를 수행하는 형태임. CUDA에서는 전체 배열을 순회하는 프로그램 대신 논리적 스레드 하나가 수행할 작업을 커널로 작성함.
  • 커널은 반복문 없이 원소 하나의 덧셈을 표현하며, 같은 커널을 각자 다른 상태와 데이터를 가진 여러 스레드에서 실행하는 방식이 SIMT(Single Instruction Multiple Threads, 단일 명령 다중 스레드)임.
  • 벡터 덧셈에서 C[i]는 i와 다른 인덱스의 결과에 의존하지 않으므로, 스레드마다 같은 연산을 서로 다른 데이터에 적용할 수 있음.
  • CUDA는 각 스레드에 식별자를 제공함. threadIdx.x는 블록 안의 스레드 인덱스이며, 각 스레드는 이 값에 따라 서로 다른 배열 원소를 처리함.
  • 커널 호출의 <<<1, 32>>> 표기는 블록 하나에 스레드 32개를 실행한다는 뜻임. 블록 하나의 스레드 수는 하드웨어 제한으로 최대 1,024개이며, 이 제한의 이유는 다음 편에서 다룰 내용임.

블록과 그리드

  • CUDA 커널 실행은 블록들의 그리드로 구성되며, 각 블록은 스레드들의 집합임. threadIdx.x는 블록마다 0부터 다시 시작하므로, 블록 인덱스와 블록 크기를 함께 사용해야 전체 실행에서 고유한 스레드 인덱스를 얻을 수 있음.
  • blockIdx.x는 현재 블록의 인덱스이며 blockDim.x는 블록당 스레드 수임. 전체 인덱스는 blockIdx.x * blockDim.x + threadIdx.x로 계산함.
  • 임의 길이 배열을 처리하려면 블록당 스레드 수를 정하고, 필요한 블록 수를 (N + threadsPerBlock - 1) / threadsPerBlock로 계산함. 이는 나머지가 있으면 블록 수를 올리는 올림 나눗셈임.
  • 블록당 스레드가 256개일 때 원소가 1,050개면 블록 5개와 초과 스레드 230개, 원소가 1,024개면 블록 4개와 초과 스레드 없음, 원소가 1,000개면 블록 4개와 초과 스레드 24개가 필요함.
  • 배열 길이가 블록 크기의 배수가 아니면 존재하지 않는 원소를 가리키는 스레드가 생길 수 있으므로, 커널에서 i < N을 확인해 유효한 인덱스만 계산함.

나머지 코드

  • 커널은 단독으로 실행되지 않음. CPU인 호스트가 GPU인 디바이스 메모리를 할당하고, 데이터를 복사한 뒤 커널을 실행하고, 결과를 호스트로 다시 복사해야 함.
  • 명시적 메모리 모델에서는 호스트 배열과 GPU 배열이 서로 다른 메모리에 존재함. cudaMalloc으로 디바이스 메모리를 할당하고, cudaMemcpy의 호스트-디바이스 방향으로 입력을 복사한 뒤 커널을 실행하고, 디바이스-호스트 방향으로 결과를 복사하며, 마지막에 디바이스 메모리를 해제함.
  • 변수 이름의 d_ 접두사는 디바이스에 있는 메모리를 나타내는 관례임. CUDA에는 메모리 공간의 분리를 완화하는 다른 모델도 있지만 여기서는 다루지 않음.
  • 일반 CPU 포인터를 커널에 전달해도 그대로 작동하지 않으며, 이 점이 첫 CUDA 커널을 작성할 때 흔히 혼란을 일으키는 부분임.

전체 과정 통합

  • 완전한 예제는 호스트에서 원소가 각각 1과 2인 배열 A와 B를 각각 1,000만 개 생성하고, GPU에 A·B·결과 배열을 할당한 뒤 입력을 복사함.
  • 블록당 256개 스레드와 올림 나눗셈으로 블록 수를 정해 벡터 덧셈 커널을 실행하고, 결과를 호스트로 복사함. 결과 배열의 첫 원소와 마지막 원소가 모두 3인지 확인한 뒤 디바이스 메모리를 해제함.
  • 이 프로그램은 배열 덧셈 커널 하나로 1,000만 번의 덧셈을 수행하는 완전한 CUDA 프로그램임.
  • 여기까지는 프로그래머가 보는 GPU의 실행 모델을 다룸. 수백만 개의 스레드가 하드웨어에 어떻게 배치되고 스케줄링되는지, 모두 동시에 실행되는지, 성능을 어떻게 추론하고 커널을 최적화하는지는 다음 편에서 다룰 내용임.