SHPC 19 - Optimization for GPUs

Hardware Scheduling Units in GPUs

• Basic unit of GPUs for scheduling
  • All threads in it processes a single instruction at the same time in SIMD fashion
  • Lock-step으로 실행
    • 모든 instruction이 실행된 후 그 다음 instruction이 실행되는 방식
• NVIDIA - warp단위 스케쥴링
  • 32 hardware threads (work-items)
• AMD - wavefront단위 스케쥴링
  • 64 hardware threads (work-items)
• 서로 다른 벤더마다 이름이 다름

An SM in NVIDIA Fermi Architecture

• GPU 내에 여러 SM이 존재
• 각 SM이 여러개의 쓰레드를 실행
• 하나의 SM에 여러 쓰레드 블록이 동시에 배정될 수 있음
  • 이런 경우 SM의 리소스에 따라 스케쥴링
  • 그림에 Warp Scheduler 가 스케쥴링을 수행함
    • 이것이 fetch & decode unit이라고 생각하면 됨
    • Dispatch unit이 fetch & decode 결과를 dispatch하여 실행
• 그림 설명
  • 레지스터는 알아서 리저브
  • SFU: sin, cos 빠르게 계산
  • 쿠다 코어 32개
    • 워프가 32개라면, 워프 스케쥴러가 두개니 반반 나눠서 실제로는 16개씩 나눠서 돌아가게 된다
    • 게다가 모두 cos, sin을 계산한다면 네개밖에 못돌아간다.
  • 이런 상황들을 고려해 알맞게 instruction을 issue하는게 warp scheduler의 역할
• 만약 쓰레드 블록이 128개 쓰레드라면, 모두 돌리기 위해서는 1/4씩 나눠서 네번 돌려야 한다.
  • SM에게 나눠지는 작업 단위는 블록이지만 블록 하나를 동시에 돌릴 수 없는 경우도 있다는 것
  • 워프로 나누어 돌아간다

SIMT (NVIDIA Terminology)

• 워프 내의 모든 쓰레드가 동일한 instruction을 동시에 실행하는 것
  • 32개 쓰레드를 묶어서 워프라고 부름
• 각 SM은 런타임에게 할당받은 쓰레드 블록을 warp단위로 나눔
  • 여러 warp를 스케쥴링
  • Logical 한 측면에서는 블록들이 모두 병렬처리됨
    • Physical 한 측면에서는 모든 쓰레드가 동시에 돌아가는 것은 앙님
    • 하나의 쓰레드 블록 내의 서로 다른 쓰레드는 실제로는 서로 다른 속도로 수행될 수 있음

Thread Block scheduling

• 쓰레드 블록은 하나의 SM에 할당됨
  • 하나의 SM에는 여러개의 쓰레드 블록이 할당될 수 있음
• Shared memory는 SM에 resident한 쓰레드 블록들 사이에 partition이 됨
  • resident: 리소스를 잡고 있는 것
    • 커널 코드에 의해 shared memory, register 사용량이 결정됨
    • 이 사용량에 따라 최대 resident한 블록의 갯수가 결정됨
    • 따라서 많이 사용하면 한번에 올라갈 수 있는 블록의 수도 줄어둠
      • concurrency가 떨어진다
    • occupancy도 중요한데, 나중에 설명
  • 따라서 쓰레드 블록 네개가 resident한 경우, shared memory를 1/4로 나눠 사용한다
• 데이터를 공유하는 것은 race를 발생시킬 수 있음
  • 각 쓰레드가 동일한 데이터에 서로 다른 순서로 접근하는 경우, UB 발생
  • 따라서 블록 내 배리어를 사용 가능
  • 그러나 쓰레드 블록 간 synchronization 방법은 없음
    • 소프트웨어적으로 구현 가능하나 너무 느림
• 블록 내의 쓰레드는 워프 단위로 나눠짐
  • 워프는 32개의 (threadId.x 가) 연속적인 쓰레드를 하나로 묶는다
  • 각 쓰레드는 각자의 private data에 대해 동일한 operation을 수행 (SIMT)
• 따라서 128개의 쓰레드가 든 블록 하나는 여러 워프로 나눠서 돌아간다
  • Warp 0: thread 0, thread 1, … thread 31
  • Warp 1: thread 32, thread 33, …, thread 63
  • …
  • 이 과정에서 레지스터, shared memory를 블록 전체가 쓰는 만큼 충분히 잡아야 한다
• 2-3차원 쓰레드 블록은 physical layout 1차원으로 매핑 가능
• 즉 2차원 쓰레드 블록의 경우, 각 쓰레드의 identifier는 다음과 같이 계산됨
  • threadIdx.y * blockDim.y +threadIdx.x
  • 3차원: threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x

Warp scheduling

• 하나의 쓰레드 블록 내의 warp는 어떤 순서로든 스케쥴링될 수 있음
• Active warp의 수는 SM resource에 의해 제한됨
  • Active warp: SM resource를 필요한 만큼 잡은 warp. 즉, resident한 warp
    • 여기서의 SM resource: pc, shared mem, register …
• 하나의 워프가 idle상태이면(데이터 read를 기다린다거나 해서), 어떤 thread block이든 resident인 쓰레드 블록에서 available한 active warp를 가져오게 됨 (매 사이클마다 체크)
• 워프 간 컨텍스트 스위칭은 오버헤드가 적음
  • 하드웨어 리소스가 할당되어 있기 때문
• 각 SM이 두개의 워프 스케쥴러와 두개의 dispatch unit이 있는 경우
• 쓰레드 블록이 SM에 할당되면, 쓰레드 블록 내의 모든 쓰레드는 워프로 변환됨
  • 두개의 워프 스케쥴러가 각자 하나씩 워프를 골라서 각 워프의 Inst를 16개 CUDA core or 16개 load/store/unit or 네개의 SFU에 issue함
• 페르미 구조에서는 SM마다 48개 워프를 동시에 돌릴 수 있게 구성
  • 48 * 32 = 1536 개 쓰레드가 하나의 SM에 동시에 resident(active)
  • 이유는 Pc가 48개라서
  • 여기서 동시에 돌린다는 것도 logical한 의미 (resident)
    • physical한 의미에서는 워프 스케쥴러의 수에 의해 결정됨

Branch Divergence

• GPU는 복잡한 branch prediction mechanism을 가지고 있지 않음
• 워프 내 모든 쓰레드가 동일한 Instruction을 수행해야 하는 구조
  • 같은 워프가 서로 다른 path를 따르면 문제가 될 수 있음
  • if - else문에서 발생
  • 쉬는 쓰레드가 생기기 때문
• 따라서 레지스터를 적게 쓰기 위해 If를 많이 넣는 것은 트레이드오프가 있음
  • sweet spot을 알아서 찾아야

Resources for warp

• 워프의 local execution context는 주로 다음과 같은 리소스로 구성되어 있음
  • PC
  • Register
  • Shared memory (local memory in OpenCL)
• 할당받은 리소스(execution contetxt)는 워프가 끝날 때까지 SM에 존재
  • 따라서 컨텍스트 스위칭이 비용이 들지 않는 것
• 병렬화 가능성을 늘리기 위해서는 resident 블록, 워프의 수를 늘려야 하고, 이는 shared memory와 register 사용량에 의해 결정
  • pc의 수에 의해서도 결정됨 (upper limit)
  • 많이 사용하면 그만큼 더 적게 resident 하므로 병렬화 X
• 컨텍스트 스위칭 기회를 늘려야 한다

Thread Granularity

• 각 thread(work item)에 많은 일을 투입하고, 적은 쓰레드를 사용하는 것
  • Work item 하나에 많은 일을 할당해 전체 work item의 수를 줄이는 전략
• 장점
  • 커널 런칭 오버헤드 감소
    • 교수님께서는 커널 런칭의 횟수가 동일할 것이기에 아닐 것 같다고 하셨지만, 내 생각에 Input size가 동일하다면 커널 런칭 수도 감소할 것이므로, 해당 환경을 가정하면 맞는 것 같다.
  • 쓰레드 간 필요없는 작업을 없앨 수 있음
• 단점
  • 사용
  • 레지스터의 수를 늘려 occupancy가 낮아질 수 있다
  • 쓰레드 블록의 수가 줄어들어 SM을 충분히 활용할 수 없을 수 있다

Grid and Block Size

• 블록 하나에는 워프 하나정도(32개)의 쓰레드는 들어있어야
  • 아니면 워프가 비어있을 것이기 때문
• 작은 블록 사이즈를 막을 것
  • 하나의 블록에 최소 128 혹은 256개의 쓰레드를 두고 늘려 나갈 것
    • 블록 갯수가 SM갯수보다 많아서, SM에 여러개 블록이 할당되면 어차피 동시실행되니 상관없지 않을까? 라는 생각을 했는데, 그것도 나름의 스케쥴링/공유 메모리 사용 불가 등의 이슈로 성능이 낮아질 수도 있음…
  • 커널의 resource requirement에 따라 블록의 크기를 키우고 줄일 것
    • 리소스를 적게 사용한다면 블록의 크기를 늘려도 되고, 리소스를 많이 사용한다면 블록의 크기를 줄여야 할 것
    • 리소스를 적게 사용하는 경우: 하나의 쓰레드가 적은 리소스를 사용하므로, 동시에 많은 쓰레드를 실행하는 것이 유리 → 블록 크기 늘리기
    • 리소스를 많이 사용하는 경우
      • 레지스터의 경우 : 하나의 쓰레드가 레지스터를 많이 써서 다른 쓰레드를 대기시킴 → 블록 크기 줄이기
      • 공유 메모리의 경우 : 블록에 할당된 공유 메모리가 커지므로 SM 공유 메모리 총량 초과 가능 → 블록 크기 줄이기
• 블록의 수를 SM의 수보다 더 많이 가져가야 함
  • SM이 놀고 있으면 비효율적이니

Syncrhonization

• 2 level 동기화
  • 시스템 레벨 : cudaDeviceSynchronize()
  • 블록 레벨 : __syncthreads()

Parallel Reduction

• 각 Column이 쿠다 코어이자 쓰레드라고 했을 때, 두 가지의 reduction이 가능
• 네모는 해당 코어가 접근하는 글로벌 메모리라고 생각
• 두 장단점
• 위 그림보다는 이게 더 나아서 가져와봄
  • Naive approach
  • 
    • Compute 측면: 커널코드를 짰을 때 분기가 생김. 예를 들어 두번째 단계에서는 2의배수 쓰레드들만 연산함. 따라서 branch divergence가 발생한다
    • Memory 측면: 결과값 저장이 칸을 띄우고 이뤄지고 있음. 데이터가 큰 경우 서로 다른 청크에 접근하여 메모리 접근이 여러 번 필요할 수 있음
  • Compute 측면 해결
    • 
    • 분기가 생기지 않는다. 워프가 2개 쓰레드 단위라고 하면, 두번째 단계에서는 워프 0, 1만 실행하면 되고 그 내부에서는 분기가 없다
  • Memory 측면도 해결
    • 
    • 결과값 저장이 연속적으로 이뤄짐. 데이터가 클 때도 묶어서 트랜잭션을 보낼 수 있게 됨
    • Coalesced memory access

Occupancy

• SM 하나당 active warp의 수
  • 컴파일 타임에 계산된다
• Active warp: Active block에 포함되는 warp
  • Active block: 계산을 위한 리소스를 할당받은 블록
  • 워프 스케쥴러가 매 사이클마다 active warp를 골라 dispatch한다
• Active warp의 분류
  • Selected warp : 실제로 실행되고 있는 active warp (running)
  • Stalled warp : 실행할 준비가 되지 않은 warp (blocked)
  • Eligible warp : 실행할 준비가 된 상태, 실제로 실행되지는 않고 있는 active warp (ready)
• Warp가 실행되는 조건
  • 32개 CUDA core가 available
  • Current inst의 모든 argument가 ready
• Occupancy
  • $occupancy=\frac{\#ofactivewarps}{\#ofmaximumwarps}$
  • 1이 가장 이상적이지만 리소스를 많이 사용하는 커널 코드를 쓰면 1보다 적어짐
  • 쓰레드 블록의 크기 조절, 리소스 사용량을 적게 사용하는 것
  • 쓰레드 블록의 크기가 작다: under utilization
    • 하드웨어를 다 쓰지도 못했는데 upper limit에 도달할 수 있음
  • 쓰레드 블록의 크기가 크다: active 쓰레드 블록이 적어진다
    • 블록이 커지면 각 쓰레드에게 줄 수 있는 하드웨어 리소스가 적어짐

Global Memory Access

• 글로벌 메모리 load/store는 캐시를 이용함
  • 모든 global mem에 대한 접근은 L2 cache 지나감
    • L2는 global과 레이턴시 비슷하나 bandwidth 때문에 캐싱을 한다
    • 한꺼번에 모아서 주는게 좋기 때문
    • 캐시라인 사이즈는 32 바이트
  • GPU 구조나 access type에 따라 L1 cache를 지날 때도 있음
    • L1 cache를 global memory caching에 사용하는 구조가 따로 존재
    • 컴파일 타임에 의해 enable, disable됨 찾아볼것
      • Enable: nvcc -Xptxas -dlcm=ca test.cu
      • Disbable: nvcc -Xptxas -dlcm=cg test.cu
      • 이렇게 하면 L1 cache를 안쓸 수 있음
    • L1 cache는 캐시 라인 사이즈가 128 바이트이고, 디바이스 메모리의 128-aligned segment에 매핑됨
      • 따라서 만약 워프 내 각 쓰레드가 4바이트 value를 요청하면, 128 byte이므로 딱 맞음
• 커널 메모리 요청은 보통 DRAM - SM on chip memory (cache) 사이에 128byte or 32 byte memory tx로 이뤄짐
  • L1, L2 캐시가 둘 다사용되면 128
  • L2 캐시만 사용되면 32 바이트

Aligned and Coalesced Memory Accesses

• Aligned memory accesses

  • 디바이스 메모리 트랜잭션의 첫 address가 cache line size의 배수인 경우
  • 하나의 캐시라인에만 접근하면 됨
    • 이것도 L2, L1 둘 다 쓰면 128, L2만 쓰면 32일 듯

• Coalesced memory access

  • 32개 쓰레드가 워프에 존재할 때, 각 쓰레드가 하나의 연속적인 메모리 청크(캐시라인 크기의 메모리)를 접근하면 하나의 메모리 트랜잭션으로 묶을 수 있다
  • 다른 청크를 접근하면 각 청크별로 접근해야 하기 때문에 느려짐

• Global memory throughput을 증가시키기 위해서는 aligned coalesced memory access가 좋음

  • 캐시라인의 배수로 align된 메모리를 하나의 청크에 대해 접근하는 것
  • (Shared memory - bank conflict 최소화. 서로 다른 뱅크에 접근해야 효율적. 접근 대상이 다름)

• 아래와 같이 aligned coalesced memory access를 하면 하나로 묶을 수 있음

• 아래와 같이 misaligned & uncoalesced 의 경우 여러번의 트랜잭션으로 접근해야 함.

  • 0부터 접근하는 것 하나
  • 128부터 접근하는 것 하나
  • 256부터 접근하는 것 하나

• alignment를 만족시키지 않으면 에러가 나거나 느려질 수 있음 글로벌 메모리 접근할 때 캐시라인에 align해야 한다 • Coalesced memory accesses : 32개 쓰레드가 워프에 있을 때, 하나의 청크 안에 있는 데이터를 접근하면 하나의 메모리 트랜잭션이 된다. 다른 청크 접근하면 각 청크별로 접근해야 하기 때문에 느려진다.(일반적으론는 하나에 접근하게 되어있음)

Host-device Data Transfer

• Host와 device 사이의 데이터 전송은 글로벌 메모리 접근보다 더 밴드위쓰가 낮다
  • PCI-E: few GB/s
  • Global memory: a few hundred GB/s
• 따라서 최소화해야 할 필요가 있음
  • 계산 값을 넘겨주는 것보다 GPU에서 직접 계산하는 게 나음
  • 글로벌 메모리를 그 중간값 저장에 이용하는 식이라고
    • 아마도 커널 퓨전을 의미하는 듯 하다.
    • 커널 퓨전도 연속적으로 실행되는 두 커널을 하나로 합치는 과정에서 중간값을 레지스터나 shared memory에 저장하게 됨
    • 이 과정에서 global - host 복사가 필요 없어짐
• 한번에 큰 데이터를 전송하는 것이 좋음
  • DMA 하는 오버헤드 때문
  • 작은 것을 여러번 하면 그만큼 DMA 개입을 많이 해야 함

Async Copy

• Data transfer와 execution을 할 수 있음
  • copy하는 동안 커널 코드 실행
  • Command queue, stream을 두개 사용해야 함
  • Computation - communciation overlapping
• PCI-E를 쓰는 가속기들에 대해서만 적용된다고 함 (?)
• 두개의 큐 사용
  • 하나는 data transfer, 하나는 execution
  • event를 이용한 dependency 체크
• 파이프라이닝 (더블 버퍼링)
  • 아까랑은 다르게, 작게 나누어 transfer - compute - transfer를 여러 스트림에 이슈

Pipelining 더블 버퍼링/트리플 버퍼링 이런식으로 vector add를 생각했을 때… A, B transfer + execution을 작게 나눔 적용해 봐야 하는 중요한 기법

다음시간 shared memory 쓰는 것 배움, vector io 중요 워프 스케쥴러 수에 의해 워프 사용 갯수가 달라진다. 워프 스케쥴러가 네개면 네개의 워프를 돌릴 수 있음