게임엔진에서 사용할 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만 올라간다.
yuchi's development 