CUDA / Compute Shader / HW대응관계

게임엔진에서 사용할 DirectX Compute Shader 코드 짜다가 생각나서 간단하게 정리 해봤다.
DirectX 기반 게임에서 그래픽스 처리 등을 위해서 사용할거면 당연히 Compute Shader를 사용하는게 맞다.
단 GPGPU자체만으로 놓고 볼때는 CUDA가 성능면에서나 사용범위 면에서나 압도적으로 우월하다.
초심자들이 thread,block,shared memory에서 헷갈려 하는데 다행히 CUDA나 Compute Shader는 (하드웨어에 대응되는 개념이다보니) 이 부분은 사실살 똑같다.

CUDA Compute Shader 대응 H/W
연산 단위 Thread Thread SP or CUDA Core
HW점유 단위 Block Group SM
Shared Memory 48KB, __shared__ 16KB , groupshared L1 Cache (16KB + 48KB)
Barrier __syncthreads() GroupMemoryBarrierWithGroupSync()
언어 C/C++ ,ptr사용 가능 HLSL, similar to C

HW 점유단위라고 쓴 것은 처음에 스케쥴링 단위로 쓰려고 했다가 고친 것이다.
엄밀히 따져 스케쥴링 단위가 아닌데 개념적으로는 그렇게 인지하고 있어도 무방하다고 생각한다. 하지만 역시 오해의 소지가 있어 HW점유단위로 적었다.
분기등 흐름제어가 가능한 HW는 SM(SMX)이다. 이 SM안에 다수의 SP(core)를 가진다. 다수의 스레드 그룹(block/group)을 GPU HW에 올려놓고 실행을 한다. CPU처럼 레지스터 상태를 백업/복구하며 태스크 스위칭을 하지 않는다. 따라서 로컬메모리는 레지스터에, shared memory는 그 상태 그대로 shared memory에 저장해둔 상태로 실행할 코드 어드레스만 바뀐다.
현재 하드웨어 자원(레지스터 개수/shared memory사이즈)에 맞춰서 최대한 많은 block/group을 올려놓고 처리한다. 그래서 HW점유단위라고 적었다.
즉 block당 1024개의 스레드가 한 block일때 block 두개를 실행하려고 하면 1024 + 1024스레드가 자원을 점유하고 실행된다. shared memory나 register가 2개의 bloick을 실행하기엔 부족할 경우 1block, + 0.5block해서 1024 + 512 스레드로 올라가지는 않는다. 이런 경우 깔끔하게 1 block만 올라간다.


답글 남기기

아래 항목을 채우거나 오른쪽 아이콘 중 하나를 클릭하여 로그 인 하세요:

WordPress.com 로고

WordPress.com의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

Google+ photo

Google+의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

Twitter 사진

Twitter의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

Facebook 사진

Facebook의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

%s에 연결하는 중