[Parallel Computing] (22) Optimization for GPUs

GPU에서 스케쥴링을 수행하는 기본 단위가 존재하고, 해당 기본 단위 안에서 모든 스레드들은 SIMD(Single Instruction Multiple Data) 형태로 처리되며, 모든 스레드가 실행되고 다음 Instruction이 실행되는 Lock-step으로 수행된다.

 NVIDIA에서는 이 기본 단위를 Warp이라 부르며, 하나의 Warp 안에는 32개의 Hardware threads(for OpenCL -> work-times)들이 존재한다. AMD에서는 Wavefront라 부르며, 하나의 Wavefront 안에는 64개의 Hardware threads들이 존재한다.

 

 NVIDIA Fermi Architecture에서는 여러 개의 SM(Steam multiprocessor)가 존재한다. 

1. 각 SM은 여러개의 스레드들을 동시에(concurrent) 처리할 수 있다. 

2. 하나의 GPU 안에는 여러 개의 SM이 존재하며, 수천 개의 스레드를 동시에 수행할 수 있다.

3. 여러 스레드 블록이 하나의 SM에 할당되어 처리될 수 있다. 

> 추가적으로 하나의 스레드 블록은 하나의 SM에만 할당될 수 있다.

> SM의 리소스에 따라 하나 혹은 이상의 스레드 블록이 할당될 수 있다.

4. SM 내부에서 수행 스케쥴링은 Warp 단위(32-Thread)로 수행된다.

> 그러므로 스레드 블록 안에 스레드가 다른 Warp으로 할당되었다면, 물리적으로 실행되는 순서가 다를 수 있다. 

 

5. SM 내부에는 스레드 블록들 간에 공유할 수 있는 Shared Memory가 존재한다.

6. SM 내부에 존재하는 Register들도 스레드에 따라 구분되어 사용된다.

> Register가 부족할 경우, Spill을 사용할 수도 있다.

 

7. 병렬로 수행되고 있는 스레드들이 데이터를 공유하고 있다면, Race Condition이 발생할 수 있다.

> Race Condition은 여러 스레드가 동일한 데이터를 접근할 때, 순서가 정해지지 않았으므로, 예측할 수 없는 결과가 발생한다. 이런 현상을 방지하고자 CUDA는 스레드 블록 내에서 Synchronization을 수행할 수 있는 기능을 제공한다. 주의할 점은 스레드 블록 간에는 Synchronization을 수행하는 Primitives가 존재하지 않는다. 

 

8. 스레드 블록 내 스레드의 threadIdx.x 값이 연속된 32개를 하나의 Warp으로 지정한다. 

> Logical layout이 2D/3D인 경우, 이를 1D Physical layout으로 변경하여 사용.

> 2D :: threadIdx.y * blockDim.x + threadIdx.x

> 3D :: threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x

9. 스레드 블록 내의 warp들은 어떠한 순서로도 스케쥴링이 될 수 있다.

 

10. SM의 리소스에 따라 Active Warps의 수는 달라질 수 있다. 

> Active Warps의 수를 늘리려면, 적은 리소스를 사용하는 많은 Warp를 만들어야 한다. 

 

11. Warp이 어떤 이유(e.g., 메모리에서 데이터 로드를 기다리는 중)로 idle 상태일 경우, SM은 자유롭게 SM 내 존재하는 수행가능한 Warp들 중 하나를 선택하여 수행한다. 

> SM 내부에서 Warp들을 스위칭하는 데(Context Switching) 필요한 Overhead가 없다. 그 이유는 SM 내부에는 각 스레드를 수행하는 데 필요한 하드웨어 자원들이 존재하기 때문이다. 따라서 Context switching을 하기 위해, 데이터를 이동하거나 이런 추가적인 일이 발생하지 않아도 된다.

> Context switch를 하기 위하여 필요한 자원은 Program Counters, Registers, Shared memory(per 스레드 블록)가 있다. 이 자원들이 모두 wrap이 끝날 때까지 on-chip 내부에 존재하므로 context switch 할 때 비용이 없다.

 

12. SM 내부에는 2개의 Warp Schedulers가 존재한다. 각 Warp scheduler 마다 instruction dispatch units이 있다. 예를 들어 SM 내부에 32개의 Warp들이 존재한다면, 16개씩 나눠서 Warp Scheduler가 담당하여 스케쥴링을 수행한다. 즉, 동시에 수행될 수 있는 Warp의 개수는 2개이다.
> Warp의 단위는 32개의 Thread인데 CUDA Core가 32개 단위로 묶인 것이 아니라 16개로 묶인 이유는 CUDA Cores와 LD&ST, SFU가 Warp Scheduler의 2배 클럭으로 동작하기 때문이다. 그렇기 때문에 하나의 Warp 내의 32개의 스레드는 두 배 빠르게 동작하는 16개의 CUDA Cores에 할당되어 동작된다.

 

13. GPU에서는 CPU와 다르게 복잡한 Branch prediction 방식을 사용하지 않는다. 하나의 Warp 내에 존재하는 모든 스레드들은 동시에(같은 사이클에) 정확히 같은 Instruction을 수행해야 한다. 그런데 Branch인 경우, True와 Flase에 따라 수행하는 Instruction이 달라진다. 하나의 warp 내 스레드들이 서로 다른  branch 결과를 가진다면, 모든 스레드들이 True/False의 Instruction을 수행하는 시간을 가진다. 물론 True를 실행할 경우, False를 실행해야 할 스레드들은 동작하지 않는다. 반대도 마찬가지이다. 예를 들어 100개 중에 하나만 False이고 나머지가 True라 해도, 100개의 스레드는 2번(False, True) 동작을 수행해야 한다. False를 수행할 때는 99개의 스레드는 수행하지 않고 놀게 되는 것이다. 그렇기 때문에 하나의 Warp안에서는 동일한 Branch 결과를 만들도록 코딩을 하는 것이 중요하다. 

> 실행단위가 Warp이므로 Warp 간에는 branch 결과 달라도 문제 없다.

 

14. Thread 개수를 적게 사용하고, 하나의 Thread에서 수행하는 Task의 양을 늘리면 아래와 같은 효과가 발생할 수 있다.

1) 여러 thread에 kernel을 lanuching 할 필요가 없으므로 Kernel을 launching overhead가 줄어든다. 

2) 서로 다른 스레드를 사용할 때 발생하는 추가적인 계산이 없어진다.

3) 하나의 스레드가 사용하는 레지스터의 양이 증가한다. 

> 이는 하나의 SM 안에 여러 Warp이 존재할 수 없으므로 Low occupancy를 발생시킨다.

4) 스레드 블록의 수가 줄게 되고, 블록 수에 따른 사용되는 SM이 줄어들어 SM utilization이 낮아진다.

> 동시에 수행할 수 있는 SM들을 충분히 활용하지 못한다.

즉, SM들을 충분히 활용하기 위해서는 스레드 블록의 수가 많은 것이 좋다. (스레드 블록의 수가 총 SM 개수보다 많은 것이 좋다). 또한, 하나의 SM안에 많은 Warp이 존재하여, latency를 충분히 감출 수 있도록 하는 것이 좋다. 많은 warp을 하나의 SM에 할당하려면 하나의 Warp이 필요하는 리소스 자원들이 작을수록 좋다. 하나의 Warp은 32개의 Thread들로 구성되어 있으므로, 스레드 블록 내의 스레드의 수는 32의 배수가 좋다. 또한, Warp stalling으로 인한 latency를 감추기 위해 SM 당 여러 Warp이 할당될 수 있도록 하는 것이 좋다. 여러 가지 경우의 수가 존재하므로 실험을 통해 찾는 것이 좋다.

 

15. Synchonization은 System level / Block level 두 가지의 기능을 제공한다. 

1) 'cudaDeviceSynchronize()' : 모든 CUDA 동작이 끝날 때까지 Host를 막는다.

2) '__syncthreads()' : 스레드 블록 내 스레드들의 동작이 모두 끝날 때까지 막는다.

16. Occupancy는 SM(Streaming Multiprocessor) 내에서 active warp의 수를 말한다. 이는 Compile time에 알 수 있다.

우선, 스레드 블록이 active 하다는 의미는 실행하기 위해 필요한 모든 자원들을 할당받았다는 것이다. 이 Active 스레드 블록 내에 존재하는 warp을 Active warp이라고 한다. Warp scheduler는 active warp들 중에 하나를 선택하여 실행한다. Active warp은 스케쥴링 상태에 따라 3가지로 구분할 수 있다.

1) Selected warp :: 현재 수행 중인 warp

2) Stalled warp    :: Active wrap이지만 실행이 준비되지 않은 상태

3) Eligible warp    :: 실행할 수 있는 상태이지만 현재 실행하지 않은 Active Wrap

warps stalling으로 인항 latency를 감추기 위해서는 active warp 수를 크게 가져가는 것이 좋으며, 이 occupancy는 active warp 수를 SM이 가질 수 있는 최대 warp수로 나눈 것을 의미한다. 

스레드 블록 사이즈를 조절하여 Occupancy를 조절할 수 있다. 

1) 스레드 블록 사이즈(스레드 per 스레드 블록)가 작으면, 모든 리소스를 완전히 활용하기 전에 하드웨어가 SM당 warps의 수를 제한한다.

2) 스레드 블록 사이즈 (스레드 per 스레드 블록)이 커지면, 각 스레드가 이용가능한 하드웨어 리소스를 줄인다.

17.Global memory의 loads/stores/는 모두 Caching 된다. L2/L1 Cache가 존재한다. Compile time에 L1 Cache 사용여부를 선택가능하며, L1의 Cacheline size는 128B이다. 32개의 스레드가 4B씩 처리한다면, 32*8 = 128B를 한번에 처리하는 것으로 필요 하는 데이터 사이즈와 Cacheline size가 일치한다. 그렇기 때문에 한번의 Cache Read로 32개의 스레드가 동시에 수행가능하다. 

18. Aligend memory accesses는 Cacheline 단위로 메모리 접근이 발생하는 것이다. mis-aligned 되어있을 경우, 불필요한 데이터를 읽어오면서 발생하는 bandwidth 손해가 있다. 

19. Coalesced memory accesses는 하나의 warp안에 있는 32개의 스레드 모두가 연속된 메모리 chunk에 접근하는 것을 말한다. 

20. Host와 Device 간의 데이터 전송은 global memory에 접근하는 것보다 낮은 bandwidth를 가진다. 그러므로 host와 device 간의 통신은 최소화하는 것이 좋다. 전송하는 오버헤드가 존재하므로 작은 데이터를 여러 번 전송하는 것보다 큰 데이터 한번을 전송하는 것이 더 효율적이다.

21. 데이터를 전송 시 소모되는 시간을 감추기 위해 데이터 전송과 Kernel 수행을 동시에 수행하도록 할 수 있다. asysnchronous Copy를 이용하면 가능하다. 전송과 Kernel 수행을 동시에 하기 위해서 전송용 Stream과 Kernel 수행용 stream을 만들어 두고, Data 전송이 끝나면 바로 kernel을 수행하도록 할 수 있다. Kernel를 수행하는 동안 다음 Data 전송을 수행할 수 있다. 

예를 들어, A + B = C를 수행한다고 하자. 일반적으로는 A와 B 데이터 모두를 GPU에 올려두고 Kernel를 수행한다. 여기서 데이터 전송을 감추려면 전체 데이터를 연산하는 양만큼 나눈 다음, 순차적으로 데이터를 전송한다. 데이터를 전송이 끝나면 바로 해당 데이터 영역에 대한 Kernel를 수행한다. 이렇듯 Kernel를 수행하는 동안 다음 데이터를 준비한다. 이런 방식을 double buffering이라 한다. 데이터 읽기와 쓰기를 동시에 수행한다는 관점에 명명되었다.