그리드와 블록#

CUDA 코드에서 그리드는 전체 디바이스라 할 수 있고 그리드의 크기만큼 블록이 존재합니다. CUDA에서 블록은 스레드의 집합이라고 할 수 있고 블록의 크기만큼 스레드가 동일한 커널을 공유하며 실행됩니다. 동일한 블록 내에서 스레드들은 shared memory를 공유합니다.

SM#

SM 내부는 여러가지 부품으로 구성되어있습니다.

메모리 관련으로는 레지스터 파일, 공유 메모리, 상수 캐시, 텍스처 캐시, L1 I-Cache(명령어 캐시), L1 데이터 캐시, LD/ST(Load/Store Unit), TMA(Tensor Memory Accelerator)가 존재합니다.

연산 관련으로는 FP32 유닛,INT32 유닛,FP64 유닛,텐서 코어,SFU(Special Function Unit),RT 코어,TMU(Texture Mapping Unit; 텍스처 매핑 유닛)가 있습니다.

스케줄링 및 제어 관련으로는 워프 스케줄러,디스패치 유닛,pc(program counter),Active Mask(활성 마스크),분기 스택(Branch Stack),Register Base Pointer(레지스터 베이스 포인터),Scoreboard(스코어보드),Operand Collector(오퍼랜드 콜렉터),배리어 및 동기화 하드웨어이 있습니다.

SM 내부의 메모리 관련 부품 정리#

• 레지스터 파일 : SM 내의 가장 대역폭이 높은 SRAM 영역으로, 각 활성 스레드의 변수들을 저장합니다. 만약 블록 내 레지스터 사용량이 SM 내부 레지스터 파일에 존재하는 레지스터 개수를 넘을 경우 데이터가 로컬 메모리(Local Memory)로 넘어가는 레지스터 스필링(Register Spilling) 현상이 발생합니다

  • 레지스터 뱅크 : 레지스터 파일은 레지스터 뱅크 구조로 되어있어 워프내 32개의 스레드에 오퍼랜드를 한사이클만에 공급할 수 있습니다.
  • 레지스터 뱅크 충돌 : 만약 워프 내의 스레드가 같은 뱅크에 존재하는 레지스터에 접근한다면 각 뱅크는 한사이클에 1개만 오퍼랜드를 줄 수 있어 지연이 발생합니다. 하지만 일반적으로 컴파일러가 이를 피하여 레지스터를 할당하기에 대부분의 상황에서는 신경쓸 필요가 없습니다.

• 공유 메모리 : 블록 내 공유 메모리 또한 레지스터 파일과 같이 뱅크로 이루어져있으며 만약 한 워프에서 같은 공유 메모리 뱅크에 접근한다면 요청이 직렬화되며 지연이 발생합니다.

• 상수 캐시 : 읽기 전용 글로벌 메모리 영역의 데이터를 캐싱하며, 동일 주소 참조 시 브로드캐스트를 통해 모든 스레드에 동시에 공급합니다.

• 텍스처 캐시 : 공간적 2D/3D 지역성이 높은 데이터 접근에 최적화된 읽기 전용 캐시입니다.

• L1 I-Cache(명령어 캐시) : SM 레벨에서 커널의 전체 바이너리 명령어를 캐싱하며, 서브 코어의 L0 I-Cache로 명령어를 공급합니다.

• L1 데이터 캐시 : 글로벌 메모리 데이터의 지역성을 활용해 캐싱합니다.

  • Volta 이후 아키텍처에서는 L1 데이터 캐시, 공유 메모리, 텍스처 캐시가 128KB SRAM으로 융합되었습니다. cudaFuncSetCacheConfig(const void* func, cudaFuncCache cacheConfig)와cudaFuncCachePreferNone, cudaFuncCachePreferShared, cudaFuncCachePreferL1과 같은 열거형 변수를 활용해 L1 데이터 캐시와 공유 메모리 중 어디에 더 메모리를 많이 할당할지 우선순위를 지정할 수 있습니다.

• LD/ST(Load/Store) Unit : 워프 단위의 메모리 요청 주소를 계산하고, 연속된 접근을 묶어 단일 트랜잭션으로 처리하는 병합된 메모리 접근(Coalesced Memory Access)를 하드웨어적으로 수행합니다.

• L0 I-Cache(L0 명령어 캐시) : 통합 L1 명령어 캐시와 별개로, SM 내부를 구성하는 4개의 독립된 서브 코어 단위마다 물리적으로 내장된 초고속 명령어 버퍼입니다. 거대한 L1 캐시에 동시에 접근할 때 발생하는 대역폭 병목을 해소하고 명령어 인출 지연을 최소화합니다.

SM 내부의 연산 관련 유닛 정리#

• FP32 유닛 : 단정밀도 부동소수점 연산을 전담합니다.

• FP64 유닛 : 배정밀도 부동소수점 연산을 전담합니다.

• INT32 유닛 : 정수 계산, 메모리 주소 계산, 제어 흐름 평가를 담당하며 FP32 유닛과 독립적으로 작동하여 동시 병행 실행이 가능합니다.

• SFU(Special Function Unit; 초월함수 유닛) : 삼각함수, 지수함수, 제곱근 등의 복잡한 초월함수를 하드웨어 레벨에서 고속으로 근사 연산합니다.

• 텐서 코어(Tensor Core) : 딥러닝과 행렬 연산에 특화되어, 1사이클 내의 혼합 정밀도 기반의 행렬 곱셈-누산(MMA)을 타일 단위로 일괄 처리합니다.

• RT 코어(Ray Tracing Core) : BVH(Bounding Volume Hierarchy) 트리 순회와 광선-삼각형 교차 테스트를 전용 하드웨어 로직으로 처리하여 SM의 기존 범용 연산 유닛에 가해지는 부하를 제거합니다. 워프가 RT 코어에 명령을 하달하면, RT 코어가 독립적으로 메모리에서 트리를 읽어 연산하고 최종 결과만 반환하므로 그동안 워프는 다른 연산을 병행할 수 있습니다.

• TMU(Texture Mapping Unit; 텍스처 매핑 유닛) : 컴퓨터 그래픽스 및 GPGPU 환경에서 정규화된 텍스처 좌표 입력 시, 소프트웨어적 분기 없이 메모리 경계 처리와 다중 샘플링을 통한 선형 보간 연산을 하드웨어적으로 수행합니다.

SM 내부의 스케줄링 및 제어 관련 부품 정리#

• 워프 스케줄러(Warp Scheduler) : 매 클럭 사이클마다 SM에 할당된 활성 워프들의 메모리 대기 및 데이터 종속성 상태를 평가하여, 즉시 실행 가능한 워프를 파이프라인에 투입합니다. 이를 통해 메모리 접근으로 인한 유휴 시간을 숨기는 지연 시간 은닉을 달성합니다.

• 디스패치 유닛(Dispatch Unit) : 워프 스케줄러가 선택한 워프의 명령어를 디코딩하고, 실제 연산을 수행할 하부 실행 유닛(FP32, INT32 등)으로 명령어를 발행합니다.

• PC(Program Counter) : 실행할 명령어의 메모리 주소를 가리키는 하드웨어 레지스터입니다. Volta 아키텍처 이후 도입된 독립 스레드 스케줄링(Independent Thread Scheduling) 체제에서는 워프 내 32개의 스레드가 각각 독립적인 PC와 호출 스택을 물리적으로 보유하여, 스레드 간의 세밀한 동기화와 병행 실행을 지원합니다.

• Active Mask(활성 마스크) : 32비트 상태 레지스터로, 워프 내 32개 스레드 중 현재 사이클의 명령어 실행에 참여하여 연산 결과를 레지스터나 메모리에 기록할 대상을 결정합니다. 분기문으로 인해 Warp Divergence가 발생하거나 조기 종료된 스레드의 비트는 0으로 비활성화됩니다.

• 분기 스택(Branch Stack / SIMT Stack) : 조건 분기로 인한 워프 다이버전스 발생 시, 지연된 분기 경로의 PC와 활성 마스크 상태를 push하고 병합 지점에서 pop하여 실행 흐름을 통제하는 LIFO 구조의 하드웨어 내부 스택입니다.

• Register Base Pointer(레지스터 베이스 포인터) : 워프가 스케줄링될 때 할당받은 거대한 물리적 레지스터 파일 내의 연속된 주소 블록 시작점을 가리킵니다. 논리적 레지스터 주소를 물리적 주소로 실시간 매핑하며, 메모리 대기 시 이 포인터의 참조 대상만 교체하는 방식으로 무비용 컨텍스트 스위칭을 가능하게 합니다.

• Scoreboard(스코어보드) : 발행 예정인 명령어가 사용할 레지스터의 데이터 종속을 비트 형태로 추적합니다. 데이터 해저드를 방지하고 명령어 수준 병렬성을 하드웨어 레벨에서 관리합니다.

• Operand Collector(오퍼랜드 콜렉터) : 레지스터 파일과 연산 유닛 사이에 위치하여, 레지스터 뱅크 충돌이 발생했을 때 여러 사이클에 걸쳐 피연산자들을 읽어들여 내부 버퍼에 모은 뒤 연산 유닛으로 한 번에 디스패치하여 파이프라인 버블을 방지하는 중재 유닛입니다.

• 배리어 및 동기화 하드웨어 : 비동기 연산의 일관성을 제어하기 위한 하드웨어 장치입니다.

  • MBarrier : 공유 메모리에 상주하며, 비동기 연산의 완료 상태를 ‘메모리 트랜잭션 바이트 수’의 도달을 기준으로 추적하고 동기화합니다.
  • cp.async : 레지스터 파일을 우회하여 글로벌 메모리에서 공유 메모리로 데이터를 직접 라우팅하는 비동기 메모리 복사 파이프라인입니다.
  • DSMEM(Distributed Shared Memory) : 클러스터 내의 여러 SM이 고속 네트워크를 통해 서로의 물리적 공유 메모리 공간에 직접 접근하고 동기화(클러스터 배리어)할 수 있도록 지원합니다.

병합된 메모리 접근 (Coalesced Memory Access)#

워프 내 32개의 스레드가 글로벌 메모리의 데이터를 읽거나 쓸 때, 하드웨어의 LD/ST 유닛은 스레드들의 메모리 요청 주소를 분석합니다.

• 트랜잭션 병합 : 만약 32개의 스레드가 요청하는 메모리 주소가 연속적이고 올바르게 정렬되어 있다면, 메모리 컨트롤러는 이를 개별적으로 처리하지 않고 32바이트, 64바이트, 또는 128바이트 단위의 최소화된 거대한 메모리 트랜잭션으로 병합하여 한 번에 데이터를 가져옵니다.
• 비병합 접근(Uncoalesced Access)의 페널티 : 만약 스레드들이 흩어진 주소를 요청하거나, 분기문으로 인해 접근이 불규칙해지면 하드웨어는 데이터를 가져오기 위해 수많은 작은 메모리 트랜잭션을 발행해야 합니다. 이는 요구 데이터 대비 버스 전송량을 급증시켜 막대한 대역폭 낭비를 초래하며, 커널 성능을 수십 배까지 저하시키는 주요 원인이 됩니다.

병합된 메모리 접근을 위한 데이터 레이아웃 (AoS vs SoA)#

하드웨어 LD/ST 유닛이 32바이트/128바이트 단위의 단일 메모리 트랜잭션으로 데이터를 가져오기(Coalescing) 위해서는 메모리상에 데이터가 배치된 구조가 매우 중요합니다.

• AoS (Array of Structures) : struct { float x, y, z; } point[1024];와 같이 객체 지향 프로그래밍에서 흔히 사용하는 방식입니다. 연속된 스레드(Thread 0, 1, 2…)가 각자의 point[i].x에 접근할 때, 메모리상에서 x 값들은 y, z 크기만큼 떨어져 존재하게 됩니다. 이 불연속적인 스트라이드 접근 패턴은 비병합 접근을 유발하여 메모리 대역폭을 크게 낭비합니다.
• SoA (Structure of Arrays) : struct { float x[1024], y[1024], z[1024]; } points;와 같이 동일한 속성들을 연속된 배열로 모아두는 방식입니다. 워프 내 32개의 스레드가 동시에 points.x[threadIdx.x]에 접근할 때, 요구하는 메모리 주소가 완벽하게 연속적이므로 LD/ST 유닛이 단 한 번의 128바이트 트랜잭션으로 데이터를 병합하여 가져올 수 있습니다.

SM 외부의 디바이스 메모리 계층: L2 캐시#

SM 내부의 레지스터, 공유 메모리, L1 캐시를 거쳐, VRAM(글로벌 메모리)에 도달하기 직전의 마지막 온칩 하드웨어 캐시 계층입니다.

• 통합 공유 구조 : 독립적인 자원을 갖는 각 SM들과 달리, L2 캐시는 GPU 칩 내의 모든 SM이 공유하는 단일 메모리 공간입니다.
• 역할 및 동작 원리 : 모든 글로벌 메모리 읽기/쓰기 요청은 반드시 L2 캐시를 통과합니다. 여러 SM이 글로벌 메모리의 동일한 주소 영역을 참조할 때 메모리 트랜잭션의 중복을 방지하며, 원자적 연산의 직렬화와 가시성 제어가 L2 캐시 레벨에서 이루어집니다.
• 캐시 라인 : L2 캐시는 기본적으로 32바이트 단위의 캐시 섹터로 관리됩니다. 따라서 커널이 단 1바이트의 데이터를 요청하더라도 하드웨어는 VRAM에서 32바이트를 읽어옵니다. L1 캐시가 활성화된 상태라면 128바이트 트랜잭션 단위로 인출됩니다.

명령어 수준의 실행 및 연산 최적화#

물리적 연산 유닛(FP32, SFU 등)을 유휴 시간 없이 최대로 구동하기 위해서는 컴파일러에 의존하는 제어 흐름 최적화가 필요합니다.

워프 셔플 명령어#

블록 내 스레드 간 통신을 위해 일반적으로 공유 메모리를 사용하지만, 공유 메모리 역시 메모리 할당, 뱅크 충돌, 그리고 동기화(__syncthreads())에 따른 지연 시간이 발생합니다. 최상의 성능을 내기 위해 이를 하드웨어 레벨에서 우회하는 기법이 존재합니다.

• 레지스터 직접 통신 : Kepler 아키텍처부터 도입된 워프 셔플 명령어(__shfl_sync,__shfl_down_sync 등)를 사용하면, 동일한 워프 내에 속한 32개의 스레드가 공유 메모리를 거치지 않고 서로의 레지스터 값을 직접 읽고 쓸 수 있습니다.
• 동작 원리 및 이점 : 하드웨어의 데이터 경로를 통해 레지스터에서 레지스터로 데이터가 1~2 사이클 만에 즉시 전달됩니다. 공유 메모리를 할당할 필요가 없어 SM 점유율 하락을 막을 수 있고, 동기화 오버헤드가 없으며, 리덕션(Reduction, 배열의 총합이나 최댓값을 구하는 연산)과 같은 알고리즘을 구현할 때 대역폭과 속도를 극단적으로 끌어올릴 수 있습니다.

원자적 연산과 하드웨어 직렬화#

여러 스레드가 동일한 메모리 주소에 동시에 값을 읽고 쓰는 작업을 수행할 때, 데이터의 무결성을 보장하기 위해 원자적 연산(atomicAdd(), atomicMax() 등)을 사용합니다.

• 하드웨어 구현 위치 : 원자적 연산은 SM 내부의 연산 유닛이 처리하는 것이 아니라, 타겟 메모리가 위치한 메모리 컨트롤러(L2 캐시 또는 공유 메모리 뱅크) 레벨에서 하드웨어적으로 처리됩니다.
• 직렬화 병목 : 수백 개의 스레드가 글로벌 메모리의 단일 주소에 동시에 atomicAdd를 요청하면, L2 캐시 컨트롤러는 이 요청들을 병렬로 처리하지 못하고 큐에 넣어 한 번에 하나씩 직렬로 처리해야 합니다. 이를 원자적 충돌이라고 하며, 병렬 처리의 이점을 완전히 상실하게 만들어 심각한 성능 저하를 유발합니다.

따라서 원자적 연산의 사용을 최소화하거나, 글로벌 메모리 대신 지연 시간이 짧은 공유 메모리 레벨에서 원자적 연산을 1차적으로 수행한 후 그 결과만 글로벌 메모리에 한 번 원자적으로 반영하는 방식으로 병목을 분산시켜야 합니다.

커널 실행 오버헤드#

일반적으로 호스트에서 <<<...>>> 문법이나 cudaLaunchKernel()을 통해 GPU 커널을 실행할 때마다 약 5~10 마이크로초의 CPU 오버헤드가 발생합니다. 커널의 연산 시간이 매우 짧고 반복 횟수가 많은 워크로드(예: 딥러닝 추론, 반복적인 물리 시뮬레이션)에서는 GPU가 연산을 끝내고 다음 명령을 기다리는 유휴 시간이 발생하여 전체 성능이 CPU에 의해 병목되는 현상이 나타납니다.

• 그래프 기반 실행 모델 : CUDA Graphs는 수천 개의 커널 실행, 메모리 복사(cudaMemcpy), 동기화 노드들을 하나의 거대한 ‘작업 그래프(DAG)‘로 사전에 정의합니다.
• 하드웨어 직접 디스패치 : 런타임에 이 그래프를 통째로 1번만 실행하면, 이후의 작업 스케줄링과 노드 간의 종속성 해결은 CPU의 개입 없이 GPU 하드웨어 내부의 스케줄러(Grid Management Unit)가 직접 처리합니다. 이를 통해 커널 실행 오버헤드를 극단적으로 줄이고 GPU 활용률을 극대화할 수 있습니다.

글로벌 메모리#

• 논리적 접근 범위 : GPU 디바이스의 메인 VRAM입니다. 실행 중인 모든 그리드, 블록, 스레드에서 자유롭게 읽고 쓸 수 있으며, 호스트와 데이터를 주고받는 메인 통로 역할을 합니다.
• 물리적 실체 : SM 외부에 존재하는 오프칩(Off-chip) 메모리이므로, 용량이 가장 큰 대신 대역폭이 가장 낮고 접근 지연 시간이 수백 사이클에 달할 정도로 매우 느립니다.

로컬 메모리#

• 논리적 접근 범위 : 각 스레드에게 개별적이고 독립적으로 할당되는 논리적 메모리 공간입니다. 스레드 내부에서 선언된 변수나 배열이 너무 커서 레지스터 스필링 발생시 임시 저장소로 사용됩니다.
• 물리적 실체 : 로컬 메모리는 SM 내부에 존재하는 독립적인 물리적 공간이 아닌 가장 느린 글로벌 메모리(VRAM)의 특정 영역을 스레드별로 분할하여 사용하는 것에 불과합니다.