CUDA 프로그래밍에서 CPU는 Host, GPU는 Device로 칭합니다. 초보자가 가장 먼저 접하는 개념은 병렬 아키텍처이며, Grid와 Block의 차이는 혼란을 야기할 수 있습니다. 본 글에서는 이들 개념을 명확히 정리합니다.

Grid, Block, Thread 관계

• Thread: 병렬 연산의 최소 단위 (경량 쓰레드)
• Block: 협력하는 Thread들의 그룹. 한 Block 내 Thread는 동기화 및 빠른 데이터 교환 가능. 최대 512개 Thread 지원
• Grid: Block의 집합. 전역 메모리(global memory)를 공유
• Kernel: GPU에서 실행되는 프로그램. 하나의 Kernel은 하나의 Grid에 대응

Block과 Thread는 각각 blockIdx(1D, 2D)와 threadIdx(1D, 2D, 3D)라는 고유 ID를 가집니다. 또한 blockDim과 threadDim으로 차원 정보를 표현하며, x, y, z 세 가지 성분이 있습니다. Block 내 모든 Thread를 동기화하려면 __syncthreads()를 사용합니다.

각 Thread는 자체 레지스터(register)와 로컬 메모리(local memory) 공간을 보유합니다. 여러 Thread가 하나의 Block을 구성하며, 이들은 공유 메모리(shared memory)를 함께 사용합니다. 모든 Thread(다른 Block의 Thread 포함)는 전역 메모리(global memory), 상수 메모리(constant memory), 텍스처 메모리(texture memory)를 공유합니다. 서로 다른 Grid는 각자의 전역 메모리, 상수 메모리, 텍스처 메모리를 가집니다.

메모리 계층

메모리 할당은 cudaMalloc과 cudaFree로 전역 메모리를 대상으로 하며, Host-Device 간 데이터 교환은 cudaMemcpy로 수행합니다.

변수 한정자

• __device__: GPU 전역 메모리 공간, Grid 내 모든 Thread 접근 가능
• __constant__: GPU 상수 메모리 공간, Grid 내 모든 Thread 접근 가능
• __shared__: GPU Thread Block 공간, Block 내 모든 Thread 접근 가능
• local: SM 내에 위치, 해당 Thread만 접근 가능

데이터 타입

내장 벡터 타입: int1, int2, int3, int4, float1, float2, float3, float4 등
텍스처 타입: texture<Type, Dim, ReadMode> texRef;
내장 dim3 타입: Grid와 Block의 구성을 정의. 예시:

dim3 dimGrid(2, 2);
dim3 dimBlock(4, 2, 2);
kernelFoo<<<dimGrid, dimBlock>>>(argument);

함수 정의

• __device__: Device에서 실행, Device에서만 호출 가능. 주소(&) 사용 불가, 재귀 미지원, 정적 변수 및 가변 인자 미지원
• __global__: void 반환, Device에서 실행, Host에서만 호출 가능. 반드시 void 반환
• __host__: Host에서 실행, Host에서만 호출 가능 (기본값)

Kernel 함수 실행 시 실행 구성(execution configuration)인 <<<....>>>를 반드시 제공해야 합니다.

__global__ void KernelFunc(...);
dim3 DimGrid(100, 50);  // 5000 thread blocks
dim3 DimBlock(4, 8, 8); // 256 threads per block
size_t SharedMemBytes = 64; // 64 bytes of shared memory
KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);

수학 함수

CUDA는 sin, pow 등 수학 함수를 제공하며, 각 함수는 일반 버전과 빠르지만 덜 정확한 __sin 같은 버전이 따로 존재합니다.

내장 변수

gridDim, blockIdx, blockDim, threadIdx, warpSize는 읽기 전용이며 할당이 불가능합니다.

프로그램 작성

CUDA는 C 언어를 잘 지원합니다. CUDA 코드를 포함하는 프로그램은 cuda_runtime_api.h를 포함하고, 파일 확장자는 .cu이며 nvcc 컴파일러로 컴파일합니다.

GPU 하드웨어 구조

Streaming Processor(SP)

GPU의 최소 연산 단위로, 완전 파이프라인 단일 이벤트 비순차 마이크로프로세서입니다. 두 개의 ALU와 하나의 FPU, 다중 레지스터 파일을 포함하며 캐시는 없습니다. 현대 GPU는 SP의 배열(SPA)로 구성됩니다. 각 SP는 하나의 Thread를 실행합니다.

Streaming Multiprocessor(SM)

여러 SP가 SM을 구성합니다. 하나의 SM은 하나의 Block을 실행하며, 8개의 SP, 2개의 Special Function Unit(SFU, 초월 함수 및 보간 계산용 FPU 4개 포함), MultiThreading Issue Unit(쓰레드 명령 분배), 명령 및 상수 캐시, 공유 메모리를 포함합니다.

Texture Processor Cluster(TPC)

특정 다른 유닛을 포함하는 SM 그룹입니다.

SPMD(Single-Program Multiple-Data) 모델

CPU는 순차적으로 코드를 실행하지만, GPU는 Thread Block 단위로 동시 실행되는 코드를 구성합니다. Kernel 프로그램은 Thread Block들의 Grid 내에서 실행됩니다. Thread Block은 협력하는 Thread들의 집합으로, __syncthreads를 통한 동기화와 공유 메모리를 통한 변수 공유가 가능하지만, 다른 Block과는 동기화할 수 없습니다. Thread Block은 1~512개의 동시 Thread를 포함하며, 고유 Block ID(1D, 2D, 3D 가능)를 가집니다. 같은 Block 내 Thread는 동일한 프로그램을 실행하지만 다른 피연산자를 사용하며 동기화가 가능하고, 각 Thread는 고유 ID를 가집니다.

Thread 하드웨어 동작 원리

GPU는 Global Block Scheduler로 Block을 SM에 할당합니다. 각 SM은 최대 8개 Block, 최대 768개 Thread를 처리할 수 있습니다(예: 1 Block x 512 Thread, 또는 3 Block x 256 Thread). 같은 SM 위 Block의 크기는 동일해야 합니다. Thread 스케줄링과 ID는 SM이 관리합니다.

SM의 부하를 최대로 하려면 적절한 Block 크기를 선택해야 합니다. 예: 8x8(64 Thread) Block은 SM에서 8 Block(512 Thread)만 처리 가능하여 비효율적이며, 16x16(256 Thread) Block은 3 Block(768 Thread)으로 최적입니다. 32x32(1024 Thread) Block은 SM이 처리 불가능합니다.

Block은 독립적으로 실행되며, 각 Block 내 Thread는 협력 가능합니다. 각 Thread는 SM 내 SP에서 실행되지만, SM에 8개의 SP만 있으므로 768개 Thread는 Warp 단위로 실행됩니다. Warp는 32개 Thread로 구성된 SM의 기본 스케줄링 단위이며, 사실상 32-way SIMD 명령입니다. 기본 단위는 half-warp입니다. SM이 768개 Thread로 만부하 시 24개 Warp가 존재하며, 매 순간 하나의 Warp 그룹만 실행됩니다. Warp의 모든 Thread는 동일한 명령을 실행하며, 각 명령은 4 clock cycle이 소요됩니다.

Thread의 생애: Grid 시작 → Block이 SM에 할당 → SM이 Thread를 Warp로 구성 → SM이 Warp 스케줄링 및 실행 → 실행 종료 후 자원 해제 → 반복.

Thread 저장 모델

Register 및 Local Memory

Thread 전용이며 프로그래머에게 투명합니다. 각 SM에는 8192개의 레지스터가 있으며 특정 Block에 할당되고, Block 내 Thread는 할당된 레지스터만 사용합니다. Thread 수가 많을수록 각 Thread가 사용할 레지스터 수는 줄어듭니다.

Shared Memory

Block 내에서 공유되며 동적 할당 가능합니다. 예: __shared__ float region[N];. Shared Memory는 16개의 Bank(각 Bank는 half-warp 길이)로 나뉘며, 연속된 32-bit 워드는 연속된 Bank에 매핑됩니다. 동일 Bank 동시 접근은 Bank Conflict이며 최소화해야 합니다.

Global Memory

캐시되지 않아 병목이 발생하기 쉬우며 최적화가 중요합니다. Half-warp 내 16개 Thread의 Global Memory 접근은 다음 조건에서 한 번의 큰 메모리 접근(Coalesce)으로 병합 가능합니다: 데이터 길이 4, 8, 16 bytes; 연속된 주소; 시작 주소 정렬; N번째 Thread가 N번째 데이터 접근. Coalesce는 성능을 크게 향상시킵니다. 비정형 접근(Uncoalesced) 시 모든 Thread가 동일 주소를 읽으면 Constant Memory를, 불규칙 읽기는 Texture Memory를, 구조체 크기가 4/8/16의 배수가 아니면 __align__(X)(X=4,8,16)로 강제 정렬을 사용할 수 있습니다.