결국 Wave Intrinsics는 하드웨어의 물리적 구조를 소프트웨어 알고리즘의 영역으로 끌어들여, GPU 성능을 한계까지 끌어올린 SM6의 정점입니다.
지난 시간, DX11(SM5)의 등장과 함께 Compute Shader와 Tessellation이 어떻게 데이터와 워크로드 모델의 패러다임을 바꿨는지 살펴보았습니다. 이는 렌더링 엔진의 책임이 커졌음을 의미하며, 동시에 GPU 프로그래밍의 주도권이 하드웨어(고정 기능)에서 소프트웨어(개발자의 설계) 로 크게 이동하는 서막이기도 했습니다. DX12는 통제권을 개발자에게 돌려준 Low-level API의 전환이었고, SM6는 Wave Intrinsics로 하드웨어의 웨이브 실행 단위를 셰이더 코드에 노출시켰습니다.
Low-Level API, Thin API
DX11에서의 리소스 상태 관리(추적 / 검증)는 드라이버의 몫이었습니다. 드라이버가 내부에서 암묵적으로 상태를 관리하고 검증과 최적화를 대신 수행해주었습니다. 이는 필연적으로 CPU 오버헤드를 동반했습니다. DX12는 이 API 레이어를 아주 얇게(Thin) 걷어냈습니다. 이제 리소스 배리어(Resource Barrier) 설정부터 메모리 배치(Heap 기반)까지 과거 드라이버가 수행하던 저수준 제어를 개발자가 명시적으로 다룰 수 있게 된 것입니다.
Command Queue와 병렬화
단일 컨텍스트(Immediate Context) 기반이었던 DX11은 아무리 코어가 많아도 렌더링 명령은 제한적인 경로로 GPU에 전달되어야만 했습니다. 특히 중간에 상태 변경이나 검증 비용이 누적되면 CPU 병목이 커지기 쉬웠습니다. 반면 DX12는 멀티스레드 환경에서 여러 개의 Command List를 동시에 기록하고 이를 GPU로 제출할 수 있는 구조를 택했습니다. 이는 현대 멀티코어 프로세서의 성능을 최대로 끌어내며 CPU 병목을 완화한 핵심 동력이 되었습니다.
Bindless Resource
과거 Draw Call 오버헤드의 큰 부분은 리소스 바인딩과 그에 따른 드라이버 검증 및 상태 변경 비용이었습니다. 즉, 드로우를 호출할 때마다 텍스처·버퍼 같은 리소스를 슬롯에 바인딩하고, 드라이버가 그 상태를 확인·정리하는 작업이 반복되면서 CPU 비용이 누적되었습니다. DX12에서는 이를 Descriptor Heap와 Descriptor Table 이라는 거대한 메모리 풀에 리소스를 몰아넣고 인덱스만 전달하는 Bindless 방식으로 개선했습니다. 덕분에 리소스 전환 비용을 최소화하고 GPU가 연산에만 집중할 수 있는 환경을 구축했습니다.
LDS vs Wave Intrinsics
Wave Intrinsics
SM6의 등장과 함께 추가된 기능 중 가장 혁신적인 것을 꼽으라면 단연 Wave Intrinsics입니다. 이는 GPU의 서브그룹 실행 단위(Warp 혹은 Wavefront)를 셰이더 코드에서 직접 다룰 수 있게 해줍니다. 같은 Wave 내 Lane들 사이에서 값을 교환하거나 집계(Reduction)할 수 있어서 그룹 공유 메모리(groupshared / LDS)와 배리어에 의존하던 패턴을 같은 Wave 내부에서는 더 싸게 연산할 수 있는 선택지를 제공하게 됐습니다.
groupshared / LDS + 배리어 의존도 SM5까지의 모델은 스레드 그룹 단위로 데이터를 공유할 때 보통 groupshared(LDS)에 쓰고, GroupMemoryBarrierWithGroupSync()로 동기화한 뒤 읽는 패턴이 일반적이었습니다. 이때 메모리 접근 지연과 배리어 대기 시간이 최적화의 비용으로 작동하기 쉬웠습니다.
Wave 내부에서 레지스터 수준 교환 및 연산을 제공하다. SM6의 Wave Intrinsics는 일부 패턴에서는 이 메모리 단계를 대체하거나 줄일 수 있습니다. Wave 내부에 한해 Lane 간 값을 레지스터 수준으로 교환 및 연산할 수 있게 합니다. (단, wave 밖에서는 여전히 LDS/배리어가 필요합니다.)
WaveActiveSum() : Wave 내 합 (Reduction)
WaveReadLaneAt() : 특정 Lane 값 읽기 (Shuffle)
WavePrefixSum() : Prefix 연산
통제가 만든 최적화의 선택지 결과적으로 개발자는 하드웨어가 실제로 실행하는 서브그룹 단위를 전제로 알고리즘을 설계할 수 있게 되었고, 특정 워크로드에서는 groupshared(LDS)나 배리어 비용을 줄이는 방식으로 성능을 개선할 여지가 생겼습니다. (Occlusion / Visibility 계열 처리, 타일링 / 스캔 / 리덕션 기반 작업 등)
(참고: wave 크기는 벤더와 아키텍처, 설정에 따라 달라질 수 있습니다.)
맺음말
우리는 지금까지 DX11에서 DX12로 넘어오며, GPU의 실행 모델과 리소스와 동기화 책임이 어떻게 더 명시화되어 왔는지 살펴보았습니다. 누군가에게 DX12는 드라이버가 해주던 일을 개발자가 더 많이 떠맡게 된 불친절한 변화였을지 모릅니다. 하지만 역설적으로 이 통제권은, UE5의 거대한 두 기둥인 나나이트(Nanite)와 루멘(Lumen) 같은 현대적 렌더링 기술을 뒷받침하는 중요한 기반이 되었습니다.
Tessellation은 기하를 '런타임 생성 데이터'로 바꿨고, Compute Shader는 연산을 '정식 워크로드'로 분리했다.
지난 시간에는 DX10/SM4로 넘어가며 Unified Shader라는 공용 셰이더 코어 풀에 대해서 다뤘습니다. 이번 시간에는 DX11(SM5)에 새롭게 추가된 Tessellation과 Compute Shader에 대해 알아보겠습니다.
Tessellation의 등장
그래픽이 발전함에 따라 거리에 따른 최적화인 Level Of Detail(LOD)가 많이 발전하게 되었습니다. 거리에 따라 폴리곤 수를 줄임으로써 깜빡임도 완화시키며 기하 aliasing 문제도 어느정도 해결하고 렌더링을 진행할 폴리곤 수도 줄어드니 1석 2조의 해결책이었죠. GPU Level에서도 기하 디테일을 정적인 폴리곤(Asset Level)로 들고가는 방식의 비용을 줄이고, 런타임에서 연속적으로 LOD와 실루엣을 개선하려는 니즈가 촉발했습니다.
DX11(SM5) 탄생 직전에 와서는 Programmable Pipeline 자체도 성숙해졌고, 단계 추가를 감당할 수 있는 설계 및 검증 역량도 증가하였습니다. 그리고 GPU 또한 메모리와 대역폭의 압박이 완화되어 초고폴리 저장 대신 필요시 생성이 더 유리해졌죠. 그리하여 고정기능으로 Tessellator + Hull Shader / Domain Shader를 결합하여 표준 파이프라인에 추가할 수 있었습니다.
Compute Shader의 등장
DX11/SM5에서 또 하나의 큰 축은 Compute Shader(DirectCompute)입니다. 이 시점부터 GPU의 병렬 연산을 그래픽 파이프라인(픽셀 셰이더 등)에 억지로 끼워 넣는 방식이 아니라, 별도의 “계산 워크로드”로 분리해 다룰 수 있는 표준 경로가 강화됩니다.
Compute Shader가 보급되기 전에는 VS / PS로 GPGPU(General-Purpose computing on Graphics Processing Units)를 진행하여 알고리즘을 돌린 사례가 많았습니다. 데이터를 텍스처나 렌더 타겟으로 포장하고 VS/PS를 커널처럼 사용해 여러 패스로 핑퐁하는 방식이 흔했습니다. 게임에서는 블룸, 톤매핑, 가우시안 블러, SSAO 같은 화면 공간 처리들은 ‘그래픽 파이프라인 기반의 데이터 병렬 계산’ 패턴으로 구현되어 왔고, 산업 영역에서도 영상/신호 처리, 과학 계산, 금융 등에서 GPU 병렬 연산 수요가 점점 커지고 있었습니다.
이런 니즈가 커지는 상황에서, GPU의 대량 병렬 실행을 안정적으로 수행하는 기반이 성숙하고(스케줄링/실행 모델), 리소스·버퍼 접근 모델(쓰기/구조화 데이터 등)도 강화되면서, 그래픽 셰이더에 억지로 넣던 작업을 Compute 워크로드로 분리할 수 있게 됩니다. 결과적으로 이후의 GPGPU/머신러닝 같은 데이터 병렬 분야가 GPU를 더 직접적으로 활용할 수 있는 토대가 마련됩니다.
DX11 Graphics Pipeline(ref. Microsoft)
데이터 모델의 변화
DX11/SM5 이전에는 Asset(Mesh, Texture, ...)처럼 "미리 만든 데이터"를 가져와 렌더링하는 비중이 컸습니다. 하지만 DX11/SM5 세대에 들어서며 테셀레이션, 시뮬레이션, 후처리 등으로 인해 GPU가 프레임 안에서 새 데이터를 생성 및 갱신하는 비중이 커졌습니다. 즉, 데이터는 더 이상 입력만이 아니라 "중간 산출물"이 되었습니다.
데이터 구조가 단순했던 텍스처나 상수에서 "구조화 데이터"로 확장됐습니다. 조명 리스트, 파티클 상태, 타일 / 클러스터 정보, 누적 결과 등. "그림"이 아니라 "테이블 혹은 레코드"에 가까워졌습니다. SM5 세대는 이런 데이터를 GPU에서 다루기 위한 구조가 더 자연스러워지고, 결과적으로 GPU 워크로드는 점점 그래픽 + 데이터 처리의 혼합이 되게 됩니다.
워크로드 모델의 변화
그렇게 프레임은 [기하 처리] => [래스터라이즈] => [픽셀 셰이딩] => [출력]이라는 단일 파이프라인이 아닌 여러 작업(기하 세밀화, 조명 누적, 후처리, 시뮬레이션, 누적/정리 연산 등)으로 쪼개지는 "작업들의 그래프"가 되었습니다. 그 작업들이 서로의 결과를 다시 입력받습니다. 그렇게 프레임은 점점 데이터 의존성을 가진 패스들의 연결 구조가 됩니다. Compute Shader 또한 계산 워크로드로 분리되면서 렌더링 단계의 제약(출력 형태, 파이프라인 규칙)에 덜 묶인 방식으로 설계할 수 있게 됩니다.
거기에 읽기 전용 리소스 중심에서 "쓰기 / 갱신 가능한 리소스" 비중 또한 커지면서 GPU는 단순 출력 장치가 아니라 중간 결과를 저장하고 재사용하는 계산 공간이 됩니다. 그것은 단순히 그리기(Draw) 단위가 아니라 Dispatch와 Pass 단위로 사고해야 한다는 것입니다. 화면 결과는 더 이상 "몇 번 그렸나"만으로 설명되지 않습니다. 어떤 작업을 그래픽으로 수행했고, 어떤 작업을 컴퓨트로 수행했으며, 그 사이에 어떤 데이터가 오갔는지가 핵심이 됩니다. 엔진 관점에서는 프레임을 Pass의 집합으로 보고, 각 Pass의 입력과 출력 리소스를 명시적으로 관리하게 됩니다.
병렬성의 핵심은 "연산량"에서 "데이터의 이동, 합치기"로 이동합니다. 단순히 많이 돌리는 것에서 끝나는 것이 아닙니다. 여러 스레드가 만든 결과를 합치거나 동일 자원에 동시에 접근할 때 규칙이 필요한데 SM5 세대 이후로는 "무엇을 계산하냐"만큼 "결과를 어떻게 축적하고 정리하는지"가 워크로드 설계의 중심으로 올라옵니다.
화면 품질의 무게중심 또한 "픽셀 셰이딩" 단일 단계에서 "프레임 전체 구성"으로 이동합니다. DX10(SM4)에서 이미 픽셀 작업과 화면 공간 패스 비중이 커졌다면, DX11(SM5)에서는 그 흐름이 더 강화되어 품질은 특정 셰이더 한 방이 아니라 여러 패스의 조합과 데이터 재사용에서 만들어집니다. 조명, 후처리, 시뮬레이션이 서로 데이터를 주고받는 구조가 자연스러워졌습니다.
렌더링 엔진의 책임
결국 렌더링 엔진의 책임이 단순히 "셰이더 작성"에서 그치지 않고 "파이프라인과 리소스 오케스트레이션"으로 확장되게 됩니다. 이젠 셰이더 한 개를 잘 짜는 것만으로는 부족해집니다. 어떤 리소스가 언제 생성되고, 어떤 단계에서 읽히고, 어떤 단계에서 갱신되는지, 그리고 그 순서가 맞는지(의존성)가 엔진 구조의 핵심이 됩니다. 즉, 엔진은 점점 워크로드 스케줄러 + 리소스 매니저의 성격을 띄게 됩니다.
끝으로
테셀레이션은 기하를 런타임 데이터 생성 대상으로 만들었고, 컴퓨트는 비그래픽 연산을 정식 워크로드로 분리했습니다. 이 둘이 합쳐지면서, 프레임은 데이터 흐름을 가진 작업 그래프로 재정의되고, 이후의 렌더링 / 시뮬레이션 / 머신러닝 같은 데이터 병렬 워크로드가 GPU로 더 자연스럽게 이동할 수 있는 기반이 형성될 수 있었습니다.
Unified Shader는 단계 통합이 아니다. VS / GS / PS가 같은 공용 코어 풀을 공유하는 구조 전환이다.
지난 시간, SIMT모델과 그것의 가장 좋은 사례인 UE3에 대해서 알아보았습니다. 이번 시간에는 DX9(SM3) 이후 DX10(SM4)이 등장하면서 Unified Shader가 왜, 어떻게 전환됐는지 알아보도록 하겠습니다.
Unified Shader의 등장(DX10, SM4)
VS / PS 같은 단계별로 유닛을 분리하여 Programmable Shader를 실행해 연산을 처리하던 DX9에도 문제점은 있었습니다. 바로 각 단계에서 워크로드가 한쪽 단계로 쏠리면 다른 단계의 유닛이 유휴(idle) 상태가 되는 로드 불균형이 발생하는 점이었습니다. 기존 DX9시절 GPU는 기본적으로 Vertex 쪽과 Pixel 쪽 연산 코어는 분리 설계가 더 저렴하고 빠르고 구현이 단순했기 때문에 분리되어 있었습니다. FFP의 유산으로 단계별 파이프를 만들기가 더 쉬웠죠.
이 부분을 더 자세히 들여다보면 SM2 / SM3 시절엔 실제로 VS / PS가 지원하는 기능이 미묘하게 달랐습니다. Pixel 쪽은 화면상의 2x2 quad 기반 파생값(ddx/ddy), 보간, 텍스처 접근 패턴 등 "래스터 기반 규칙"이 강했습니다. 하지만 Vertex 쪽은 그런 제약이 상대적으로 적고, 대신 기하 데이터 처리에 최적화 돼 있었습니다. 즉, "둘 다 셰이더니까 코어를 공유하면 되지 않나?"가 오늘 관점에서는 쉬워보여도 당시에는 실행 조건이 달라서 코어를 공용화하려면 추가 하드웨어, 규칙 정리가 필요했습니다.
(참고로 SM3의 VS에서도 텍스처 룩업 자체는 가능했지만, 픽셀 단계의 2x2 quad 기반 파생값(ddx/ddy) 같은 래스터 규칙은 본질적으로 픽셀 단계에 묶여 있었습니다.)
게다가 부하가 한쪽으로 쏠릴 때 발생하는 손해가 그때는 덜 치명적이기도 했습니다. DX9 시절의 게임은 평균 해상도, 셰이딩 복잡도, 포스트프로세스가 지금보다 낮았고, 파이프라인도 지금보다 단순했습니다. VS / PS 고정 분리 비효율이 확실한 손해로 체감되는 구간이 지금보다 덜했습니다. 그러다보니 "일을 동적으로 나눠주는 로직(스케줄러, 디스패처, 레지스터 파일, 컨텍스트 스위칭 등)" 을 넣는 비용은 당시 공정, 전력, 검증 대비 이득이 불확실했죠.
하지만 DX10 / SM4에서는 셰이더 길이와 픽셀 작업 비중이 커지면서 로드 불균형이 더 자주 체감되기 시작했습니다. 동시에 기술의 발전으로 비용도 낮아지면서 스테이지별 제약을 정리(거의 동일한 명령어, 레지스터, 리소스 접근)하고 통합된 셰이더 모델을 전제로 설계를 재정의하면서 Unified Shader 구조가 자리잡게 됩니다.
DX10(SM4) Pipeline Stages
Direct3D 10 그래픽 파이프라인의 공식 단계 순서를 그대로 배치하고 그 위에 Unified Shader(= Common Shader Cores)이 무엇을 뜻하는지 함께 표시한 그림입니다. 핵심은 두 층으로 나뉩니다.
위쪽 : DX10의 파이프라인 단계(stage)
아래쪽 : VS/GS/PS가 실행되는 공용 연산 자원(core pool)
즉, 단계(stage)가 통합된 것이 아니라 Programmable 단계가 같은 연산 자원을 공유하게 된 것이 Unified shader의 포인트 입니다.
단계별 의미
1) Input Assembler (IA)
버텍스 버퍼/인덱스 버퍼에서 데이터를 읽어와서
“정점 스트림”을 조립해 다음 단계로 넘긴다.
2) Vertex Shader (VS) (Programmable)
정점 단위 연산을 수행한다. (좌표 변환, 스키닝, 정점 속성 계산 등)
3) Geometry Shader (GS) (Programmable)
프리미티브(삼각형/라인/포인트) 단위로 연산한다.
필요하면 프리미티브를 증감/변형할 수 있다.
4) Stream Output (SO) [Optional]
GS의 결과(또는 파이프라인 중간 결과)를 버퍼로 기록하는 옵션 단계다.
그림의 **점선 화살표(Feedback)**는 SO에 저장한 데이터를 다시 IA로 재투입할 수 있음을 나타낸다. (즉 “GPU 안에서 결과를 저장→다시 입력으로 사용”하는 루프가 가능)
5) Rasterizer (RS)
정점/프리미티브를 **픽셀 후보(프래그먼트)**로 변환한다.
화면 공간으로 투영하고, 클리핑/컬링 같은 고정 기능 처리도 포함된다.
6) Pixel Shader (PS) (Programmable)
픽셀(프래그먼트) 단위 연산을 수행한다.
텍스처 샘플링, 조명 계산, 머티리얼 평가 같은 대부분의 “화면 품질” 연산이 여기서 발생한다.
7) Output Merger (OM)
PS의 출력 색/깊이 결과를 최종 렌더 타깃에 합성한다.
블렌딩, 깊이/스텐실 테스트 같은 “최종 합성 규칙”이 적용된다.
Common Shader Cores가 의미하는 것
DX10 / SM4에서 VS, GS, PS는 같은 종류의 Programmable Core에서 실행될 수 있습니다. 따라서 어떤 프레임에서 픽셀 작업이 많으면 PS쪽에 코어가 더 배정되고, 정점 / 지오메트리 작업이 많으면 VS / GS 쪽에 코어가 더 배정되는 식입니다. 따라서 로드 불균형을 줄이는 방향으로 하드웨어가 설계된 것입니다.
중요한 점
파이프라인 "단계"는 여전히 존재합니다.(IA > VS > GS > SO (Optional) > RS > PS > OM 순서 유지)
Unified Shader는 단계 통합이 아니라 연산 자원 통합(공용 코어 풀)입니다.
맺음말
정리하면 DX10/SM4의 Unified Shader는 파이프라인 단계가 사라진 것이 아니라, VS / GS / PS가 같은 실행 자원(공용 코어 풀)을 공유하도록 바뀐 구조 변화입니다. 그 결과 프레임마다 달라지는 워크로드에 맞춰 연산 자원을 더 유연하게 배분할 수 있게 되었고, DX9 시절의 로드 불균형 문제도 완화되는 방향으로 설계가 이동했습니다. 다음 글에서는 이 통합 구조 위에서 SM5 시대로 넘어가며 무엇이 확장됐는지 이어서 보겠습니다.
CPU는 SIMD를 ‘더 강한 벡터’로 키웠고, GPU는 SIMD를 ‘스레드 추상화’로 바꿨다.
지난 시간 FFP에서 고전 SIMD의 등장까지 알아보았습니다. 이제 SIMD에서 SIMT로 넘어가던 DirectX9 ~ DirectX10 시절을 알아보도록 하겠습니다.
개요
하나의 명령어로 여러 데이터를 동시에 처리하던 고전 SIMD는 구조적인 한계가 명확했습니다.
데이터 길이가 workload에 맞지 않으면 cost 낭비가 발생
스레드 개념이 없음
분기 시 lane이 갈라지면 실행이 직렬화되어 효율(활성 lane 비율)이 크게 하락
gather/scatter 미지원 => 다양한 메모리 주소 접근이 불가능
이 구조는 단순한 벡터 연산에는 강력하지만, 픽셀·조명·텍스처·조건문이 많은 실제 그래픽스 workload에는 치명적이었다.
CPU는 SIMD 폭을 확장하고 ISA를 강화하는 방식으로 발전 방향을 잡았습니다.
벡터 폭 증가
gather/scatter 지원 (일반 load/store 관점)
mask 연산 강화
분기 최소화 등
즉, CPU는 SIMD를 더 강한 벡터 처리 장치로 발전시키는 방향을 택했다.
SIMT의 등장
하지만 GPU는 완전히 다른길을 선택했습니다. 그것은 바로 SIMD를 Thread로 추상화하여 재해석하는 방법입니다. GPU는 수천 ~ 수만 개의 Pixel Fragment를 병렬 처리해야 했기 때문에 PC(ProgramCounter)를 Thread마다 하나씩 둘 수 없었습니다. PC를 Thread마다 하나씩 두면 면적과 전력소모가 폭발적으로 증가하기 때문이었습니다.
그래서 GPU는 SIMD를 다음과 같이 재해석했습니다.
"SIMD lane을 Thread처럼 보이게 만들고, 프로그래머에게는 각 lane이 독립적인 스레드처럼 보이게 하자."
이 방식이 바로 SIMT 모델입니다.
about SIMT
Warp란 무엇인가?
CPU는 Thread마다 PC가 하나씩 있습니다. 그렇기 때문에 완전한 MIMD 실행 모델로 각자 다른 위치에서 다른 코드를 실행할 수 있습니다. 하지만 GPU의 Warp의 구조는 정반대입니다.
예를 들어 Warp = 32 thread가 있다면,
(벤더/세대에 따라 warp/wavefront 크기가 다를 수 있음)
32개의 스레드는 각자 독립적으로 보이지만, Warp 전체가 PC 1개, 그렇기 때문에 32개의 스레드는 한번에 동일한 명령어만 실행 가능.
각 Thread는 고유한 Register 값을 갖지만 명령어 흐름은 공유.
이렇게 동일한 명령어를 동시에 실행하는 한 단계가 Lockstep입니다.
Branch Divergence
이 조건을 보면 Lockstep의 의미가 바로 드러납니다.
아래와 같은 조건이 나왔다고 합시다.
if (x < 16)
A 실행
else if(x < 32)
B 실행
thread 0 ~ 15 => A 실행
thread 16 ~ 31 => B 실행
이렇게 CPU 에서는 스레드 별로 명령어를 다르게 실행할 수 있기 때문에 각 스레드에서 A와 B를 나눠서 실행하게 됩니다. 하지만 GPU는 Warp는 하나의 PC만 있으므로 두 경로를 동시에 실행할 수 없습니다.
그렇기에 SIMT는 다음과 같은 방식으로 처리합니다.
Thread 32개를 컨트롤 하는 Warp의 PC를
A 경로로 이동
true thread만 활성 (mask = on)
false thread는 비활성 (mask = off)
Warp 실행(A만 실행 됨)
B 경로로 이동
false thread만 활성 (mask = on)
true thread는 비활성 (mask = off)
Warp 실행 (B만 실행 됨)
이렇게 한 warp 안에서 분기가 갈리면, warp는 A와 B를 직렬화해 순차적으로 실행합니다. 이것이 lockstep 실행의 제약인 Branch Divergence입니다. 고전 SIMD에서도 분기 직렬화는 실행 효율을 떨어뜨렸습니다. 다만 SIMT에서는 warp 단위로 작업이 쪼개져 있고, 실행 유닛이 다른 warp를 교대로 실행하며 파이프라인을 유지하기 때문에, “전체가 멈춘다”기보다는 활성 lane 감소로 효율이 하락하는 형태로 나타납니다.
맺음말
SIMT의 등장은 고전 SIMD의 한계를 '극복'했다기보다는, 그래픽스라는 특수한 workload를 처리하기 위한 현실적인 타협이었습니다.
SIMD lane을 Thread처럼 추상화하여 프로그래머에게는 MIMD처럼 보이게 만들고
내부에서는 여전히 lockstep 기반 SIMD의 효율을 유지하며
Massive Parallel Pixel Workload를 감당할 수 있도록 한 구조
즉, SIMT는 "GPU스럽게 동작하는 MIMD의 환상"을 만들어낸 모델이라고 할 수 있습니다.
Branch divergence는 여전히 존재하며 성능을 떨어뜨립니다. 하지만 FFP 시대에는 상상할 수 없었던 복잡한 조명, 그림자, 포스트 프로세싱, 물리 기반 셰이딩을 GPU가 처리할 수 있게 된 것도 결국 이 SIMT 모델 덕분입니다.
다음 글에서는 이 SIMT 모델이 어떻게 구체적으로 구현되었는지, 그리고 DirectX10~11, UE3~UE4 시대를 지나면서 하드웨어와 셰이더 모델이 어떤 방향으로 진화했는지를 이어서 다뤄보겠습니다.
Programmable Pipeline의 토대, 스칼라 중심 GPU → SIMD Lane 재구성
그래픽카드의 한계
2000년대 초반의 GPU는 FFP(Fixed Function Pipeline)중심이었고, 지금처럼 복잡한 셰이더 기반 렌더링이 불가능했습니다. VRAM 용량도 32~64MB 수준이라 SSAA나 고급 조명 기법을 적용하기 어려웠죠. 이 글에서는 당시 GPU의 한계를 살펴보고, 어떻게 SIMD 기반의 Programmable Pipeline 으로 발전했는지 정리합니다.
MSAA (Multi-Sampling Anti-Aliasing)
MSAA vs EQAA 샘플 패턴 비교 이미지
Pixel Boundary
한 픽셀을 크게 확대한 사각형 틀. (픽셀은 점이 아니라 “면적”을 가진 사각형입니.)
Color Sample Location
픽셀 내부의 특정 지점(Point)에서 실제 색상(Color)을 샘플링하는 위치. 여러 Color Sample을 블렌딩하여 최종 픽셀 색을 만든다.
Coverage Sample Location
폴리곤이 픽셀을 얼마나 덮고 있는지(Coverage) 판단하는 지점. 각 샘플이 폴리곤 내부면 1, 외부면 0에 해당하며, 이를 평균내어 블렌딩 가중치(강도)로 사용한다. Color Sample보다 계산 비용이 훨씬 낮다.
초기 GPU의 VRAM 용량이 32 ~ 64MB 하던 시절에는 픽셀 단위로 복잡한 처리를 수행할 수 없었습니다. 그래서 이미지 전체를 더 높은 해상도로 여러 번 렌더링하는 SSAA(Super-Sampling Anti-Aliasing)을 다운 샘플링을 하는 방식이 유일한 선택지였습니다. 해당 방식은 해상도 배율 x 샘플 수 만큼의 비용증가가 있었고, 이는 픽셀 셰이딩, 텍스쳐 샘플링, ROP(Blend / Depth / Color Write) 등 '픽셀 이후 단계'의 비용을 전부 O(S) 수준으로 폭증시켰습니다. 즉, 샘플 수만큼만 비용이 증가하는 것이 아니라, 렌더타겟 자체도 커지기 때문에 파이프라인의 대부분이 정비례로 비싸지는 구조였습니다.
하지만 VRAM 용량이 어느정도 여유가 생기자, 픽셀 내부의 Color Sample과 Coverage Sample을 저장할 수 있는 구조가 가능해졌고, MSAA는 Pixel Shader는 한 번만 수행하고, Coverage Test와 Color Resolve만 샘플 수에 따라 반복하는 방식으로 SSAA 대비 큰 폭의 최적화를 달성했습니다.
Fixed Function Pipeline > Programmable Pipeline 의 과도기
FFP(Fixed Function Pipeline)
Transform & Lighting
Texture Stage State 기반 2~3단 텍스처 처리
멀티패스 조명
DOT3 Bump Mapping(FFP LOD 바탕)
Lightmap 기반의 Static Lighting
UI, HUD, 간단한 머터리얼 > FFP로 렌더링
하지만 UE2 후반 버전은...
Programmable Pipeline 일부 적용
Shader-driven Material Effects
Normal Mapping
Specular Mask
Detail Normal Layer
Color Modulation
Hardware Skinning
NVIDIA/ATI 전용 Shader Path
이것이 가능해진 이유는 GPU의 주요한 변화 덕분.
Scalar ALU vs SIMD Lane
기존
Scalar ALU에 단순한 microcode 실행기를 사용하여 float4연산은 4개의 스칼라 연산을 순차적으로 연산하도록 실행, 단지 GPU에 Scalar ALU를 많이 담아서 단순 병렬처리를 수행했을 뿐.
SIMD(Single Instruction Multiple Data stream)의 등장
하나의 명령어(Single Instruction)를 Scalar ALU의 묶음인 SIMD Lane에 전달하여 연산에 사용될 Multi-Data를 레지스터/버퍼(L1/L2, GPR)에 담아 연산하도록 실행
요즘 레이트레이싱에 관련된 작업을 시작하게 되어 정리차원에서 블로깅하게 되었습니다. 현재는 TinyRaytracer를 마개조 중이었고, 이번 글은 깃헙의 TinyRaytracer를 기반으로 레이트레이싱 기법에 대해 간략하게 설명해보도록 하겠습니다. 최적화 기법은 추후에 다룰 예정입니다. 다른 포스팅에서 확인 부탁드립니다. (fresnel equation이 추가되었습니다.)
ray_intersect함수의 경우, 레이의 original position, direction을 받아서 해당 오브젝트의 충돌검출(구와 직선의 충돌 알고리즘)과 함께 해당 오브젝트까지의 거리를 t0로 넘겨줍니다.
2. 기능 설명
- render
void render(const std::vector<const sdf_model*> &models, const std::vector<Light> &lights) {
const int width = 1024;
const int height = 768;
const float fov = M_PI/3.;
std::vector<Vec3f> framebuffer(width*height);
#pragma omp parallel for
for (size_t j = 0; j<height; j++) { // actual rendering loop
for (size_t i = 0; i<width; i++) {
float dir_x = (i + 0.5) - width/2.;
float dir_y = -(j + 0.5) + height/2.; // this flips the image at the same time
float dir_z = -height/(2.*tan(fov/2.));
framebuffer[i+j*width] = cast_ray(Vec3f(0,0,0), Vec3f(dir_x, dir_y, dir_z).normalize(), models, lights);
}
}
std::ofstream ofs; // save the framebuffer to file
ofs.open("./out.ppm",std::ios::binary);
ofs << "P6\n" << width << " " << height << "\n255\n";
for (size_t i = 0; i < height*width; ++i) {
Vec3f &c = framebuffer[i];
float max = std::max(c[0], std::max(c[1], c[2]));
if (max > 1) c = c * (1. / max);
for (size_t j = 0; j<3; j++) {
ofs << (char)(255 * std::max(0.f, std::min(1.f, framebuffer[i][j])));
}
}
ofs.close();
}
카메라에서 레이를 발사하여 frame buffer에 기록하고, 그 데이터를 ppm 파일에 저장합니다.
- cast_ray
Vec3f cast_ray(const Vec3f &orig, const Vec3f &dir, const std::vector<Sphere> &spheres, const std::vector<Light> &lights, size_t depth=0) {
Vec3f point, N;
Material material;
if (depth>4 || !scene_intersect(orig, dir, spheres, point, N, material)) {
return Vec3f(0.2, 0.7, 0.8); // background color
}
Vec3f refract_color(0.f, 0.f, 0.f);
// compute fresnel
float kr;
fresnel(dir, N, material.refractive_index, kr);
// compute refraction if it is not a case of total internal reflection
if (kr < 1) {
Vec3f refract_dir = refract(dir, N, material.refractive_index).normalize();
Vec3f refract_orig = refract_dir * N < 0 ? point - N * EPSILON : point + N * EPSILON;
refract_color = cast_ray(refract_orig, refract_dir, models, lights, depth + 1);
}
Vec3f reflect_dir = reflect(dir, N).normalize();
Vec3f reflect_orig = reflect_dir * N < 0 ? point - N * EPSILON : point + N * EPSILON; // offset the original point to avoid occlusion by the object itself
Vec3f reflect_color = cast_ray(reflect_orig, reflect_dir, models, lights, depth + 1);
float diffuse_light_intensity = 0, specular_light_intensity = 0;
for (size_t i=0; i<lights.size(); i++) {
Vec3f light_dir = (lights[i].position - point).normalize();
float light_distance = (lights[i].position - point).norm();
Vec3f shadow_orig = light_dir*N < 0 ? point - N*EPSILON : point + N*EPSILON; // checking if the point lies in the shadow of the lights[i]
Vec3f shadow_pt, shadow_N;
Material tmpmaterial;
if (scene_intersect(shadow_orig, light_dir, spheres, shadow_pt, shadow_N, tmpmaterial) && (shadow_pt-shadow_orig).norm() < light_distance)
continue;
diffuse_light_intensity += lights[i].intensity * std::max(0.f, light_dir*N);
specular_light_intensity += powf(std::max(0.f, -reflect(-light_dir, N)*dir), material.specular_exponent)*lights[i].intensity;
}
return material.diffuse_color * diffuse_light_intensity * material.albedo[0]
+ Vec3f(1., 1., 1.)*specular_light_intensity * material.albedo[1]
+ reflect_color * material.albedo[2] * kr
+ refract_color * material.albedo[3] * (1 - kr);
}
레이를 발사하는 함수입니다. 먼저 장면에 충돌시켜 충돌된 위치(point)와 해당 point의 normal값을 구합니다. 그 후, 위 함수는 세 연산으로 나뉩니다.
첫번째는 레이 연산입니다. 레이 연산은 다시 세가지로 나뉘는데 fresnel equation, 반사(reflection), 투과(refraction입니다. 이 부분을 설명하기엔 길어지므로 아래의 각 파트에서 다루도록 하겠습니다.
두번째는 빛 연산입니다. 먼저 변수에 대한 설명을 진행하도록 하겠습니다.
diffuse_light_intensity는 분산광의 정도를 저장하는 변수입니다.
specular_light_intensity는 반사광의 정도를 저장하는 변수입니다.
light_dir는 빛이 diffuse light이므로 빛의 방향이 됩니다.
light_distance는 빛에서부터 point까지의 거리로 shadow casting을 할 때, 앞에 있는 물체를 검출하기 위해 사용할 값입니다. 앞에 물체가 있다면 light값을 적용하지 않습니다.
shadow_origin은 point의 위치를 EPSILON으로 보정한 값입니다. point의 앞에서 발사한 빛인지 뒤에서 발사한 빛인지를 내적으로 검출하여 EPSILON으로 보정하여 light_distance 비교연산을 할 때 오차를 줄여줍니다. light_distance 비교연산은 변수 설명이 끝난 후 설명하겠습니다.
shadow_pt는 빛에서 point로 레이를 발사하는 과정에서 부딪힌 point입니다. light_distance보다 짧은 곳에서 충돌했다면 point 앞에 물체가 있는 것으로 음영에 diffuse_light_intensity와 이 필요합니다.
tmpmaterial은 scene_intersect에 넣어주기 위한 용도로 TinyRaytracer에서는 사용하지 않습니다.
마지막으로 light_distance 비교연산을 진행하고 그 결과에 따라 light_intensity를 적용하게 됩니다. 빛에서부터 해당 point까지 scene_intersect를 진행하여 light_distance보다 짧은 거리에 위치한 물체가 있다면 light_intensity 연산을 하지않게 되는 것이지요.
세번째는 위 연산에 따라 색상을 혼합해주는 연산입니다. 여기서 주목해야할 부분은 reflect_color와 refract_color에 fresnel equation의 결과값인 kr을 적용해주는 부분입니다. 아래 fresnel편에서 자세히 다루도록 하겠습니다.
- scene_intersect
bool scene_intersect(const vec3 &orig, const vec3 &dir, const std::vector<Sphere> &spheres, vec3 &hit, vec3 &N, Material &material) {
float spheres_dist = std::numeric_limits<float>::max();
for (const Sphere &s : spheres) {
float dist_i;
if (s.ray_intersect(orig, dir, s, dist_i) && dist_i < spheres_dist) {
spheres_dist = dist_i;
hit = orig + dir*dist_i;
N = (hit - s.center).normalize();
material = s.material;
}
}
float checkerboard_dist = std::numeric_limits<float>::max();
if (std::abs(dir.y)>1e-3) { // avoid division by zero
float d = -(orig.y+4)/dir.y; // the checkerboard plane has equation y = -4
vec3 pt = orig + dir*d;
if (d>1e-3 && fabs(pt.x)<10 && pt.z<-10 && pt.z>-30 && d<spheres_dist) {
checkerboard_dist = d;
hit = pt;
N = vec3{0,1,0};
material.diffuse_color = (int(.5*hit.x+1000) + int(.5*hit.z)) & 1 ? vec3{.3, .3, .3} : vec3{.3, .2, .1};
}
}
return std::min(spheres_dist, checkerboard_dist)<1000;
}
레이를 원점으로부터 특정 방향으로 진행시킬 때의 hit point와 해당 point의 normal값을 검출하는 함수입니다.
Vec3f refract(const Vec3f &I, const Vec3f &N, const float eta_t, const float eta_i=1.f) { // Snell's law
float cosi = - std::max(-1.f, std::min(1.f, I*N));
if (cosi<0) return refract(I, -N, eta_i, eta_t); // if the ray comes from the inside the object, swap the air and the media
float eta = eta_i / eta_t;
float k = 1 - eta*eta*(1 - cosi*cosi);
return k<0 ? Vec3f(1,0,0) : I*eta + N*(eta*cosi - sqrtf(k)); // k<0 = total reflection, no ray to refract. I refract it anyways, this has no physical meaning
}
refract는 snell's law에 의해 구현되었습니다.
우선 \(cos i\)를 확인하여 광선이 매질 안에서 발사되는 것인지를 확인합니다. 만약 \(cos i\)가 0보다 작으면 매질 안에서 발사되는 것이므로 eta_i와 eta_t를 교환해 다시 refract를 진행해줍니다.
매개변수 중 eta_t는 투과를 진행할 물체의 투과율로 영어로는 ior(index of refraction)라고도 부릅니다. eta_i는 광선이 있는 곳의 투과율입니다.
