SHPC 18 - CUDA Streams

CUDA Streams

• Stream
  • GPU에서 순서대로 실행되는 operation sequence
  • OpenCL의 커맨드 큐와 유사
• 서로 다른 stream의 CUDA operation은 concurrent하게 돌아감
• Default stream (a.k.a stream 0)
  • 특정하지 않으면 사용 (실제로는 kernel<<<grid, block, 0>>> 을 실행하는 거나 마찬가지)
  • 완전히 synchronous하다
    • host, device 관점에서
    • cudaDeviceSynchronize 는 쿠다 op 전후로 넣어 sync를 하는 명령어
    • 마치 이걸 쓴거같이 동작
  • 호스트에 대해 async하게 동작하는 특정 함수들이 존재
    • 호스트가 해당 함수가 돌아가는 동안 다른 동작을 할 수 있는 것
    • cudaMemcpy*Async
    • cudaMemset*Async
    • cudaMemcpy within the same device
    • H2D cudaMemcpy of 64kB or less

CUDA Concurrency

• 여러 CUDA operation을 동시에 실행할 수 있는 능력
  • 여기서 CUDA operation이란
    • CUDA kernel
    • cudaMemcpyAsync - HostToDevice
    • cudaMemcpyAsync - DeviceToHost
    • CPU 연산
• NVIDIA Ferimi
  • 16개의 서로 다른 CUDA kernel을 돌릴 수 있다
  • GPU에 SM이 16개 들어있으면 각각 다른 kernel을 돌릴 수 있는것
    • 사실 느려서 잘 사용하지 않는다
    • 선전용이다
  • 두개의 cudaMemcpyAsyncs
    • 하드웨어의 특성을 반영한 결과
    • 디바이스와 통신을 하게 되면 GPU가 보통 PCI slot에 꽂혀있는데, 이것과 통신하는 버스(PCI express bus)의 특성이 in과 out이 bidirectional한 것이기 때문 이라고 설명하신듯 하다
  • CPU와 독립적으로 GPU가 동작하는데, 당연한 것

Requirements for Concurrency

• 서로 다른 stream의 operation이어야 concurrency를 달성할 수 있다
• cudaMemcpyAsync with host from ‘pinned’ memory
• 충분한 resource가 있어야 한다
  • bidirection 메모리 카피
  • Device resources (shared memory, registers, blocks, etc…)

Pinned Memory

• 페이징을 하지 않는 메인 메모리의 부분 (== page-locked memory)
  • cudaMallocHost() or cudaHostAlloc() 으로 생성
  • 해제는 cudaFreeHost()
• 메인 메모리에 pinned memory는 따로 표시가 되어 있음
• cudaMemcpyAsync를 할 때 호스트 메모리가 꼭 pinned memory로 설정되어야 함
  • DMA를 하는 중에 Paging을 해버리면 엉뚱한 위치를 복사할 수 있기 때문
  • 비동기 메모리 복사에 필수

Synchronous Operations

• Default stream을 이용하여 synchronous하게 동작하는 예시

Asynchronous Operations

• GPU kernels are asynchronous with host by default
• 비슷한데 CPU_function_call() 이 들어가고 있다.
  • 이 코드는 CPU operations을 말하는 것
• 즉 커널코드와 CPU function을 overlapping 하고 있는 것이다
• 참고
  • 그 아래는 cudaMemcpy가 implicit sync operation이기 때문에 얘는 앞에 kernel이 끝나기를 기다린다
  • 즉 표시한 코드 외에는 모두 sync하게 동작하는 예시임

Asynchronous with Streams

• stream을 이용해 비동기적으로 동작하는 예제로, 가장 비동기적임
• Fully asynchronous / concurrent • Data used by concurrent operations should be independent
• stream 1, 2, 3, 4에게 각각 다른 명령을 던져서 모두 overlapping 될 수 있다
• 서로 다른 디바이스가 동시에 돌아감

Explicit Synchronization

• 모든 걸 동기화하기 : cudaDeviceSynchronize()
  • 모든 CUDA call이 끝날 때까지 블록된다
  • cl_finish()와 기능적으로 같음
• 특정 stream을 동기화하기 : cudaStreamSynchronize(streamid)
  • 특정 스트림의 CUDA call이 끝날 때까지 블록
  • cl_finish()의 아규먼트로 특정 command queue를 지정하는 것과 같은 기능
• 이벤트 동기화
  • Create specific ‘Events,’ within streams, to use for synchronization
  • cudaEventRecord(event, streamid) : 이벤트 캡처.
    • 호출 이전에 스트림에 큐잉된 모든 커널이나 작업이 완료되면 해당 이벤트를 기록하라
  • cudaEventSynchronize(event) : 이벤트를 기다리는 것
  • cudaStreamWaitEvent(stream, event) : 특정 이벤트가 기록될 때까지 stream 실행을 중지시키는 것
    • 호출 이후 스트림에 큐잉되는 명령을 대기시키는 것
    • 호출 이전에 큐잉된 명령은 계속 실행됨
  • cudaEventQuery(event) : 이벤트가 기록 완료 상태인지 확인하는 것
• 예제
• 이벤트에 H2D 카피가 캡처되어 있으니, cudaStreamWaitEvent(stream2, event)는 stream 2에 들어 있는 커맨드들은 event가 끝날 때까지 시작되지 말라는 뜻
• Kernel에 stream2를 던지고 있으니, stream2의 op들이 끝날 때까지 기다려야 한다
• 따라서 H2D & D2H → kernel & CPU 순으로 실행될 것

Implicit Synchronization

• 아래 operation들은 암묵적으로 동기화를 수행함
  • Page-locked memory allocation
    • cudaMallocHost
    • cudaHostMalloc
  • Device memory allocation
    • cudaMalloc
  • Non-Asycnc version of memory operations
    • cudaMemcpy
    • cudaMemset
  • L1/shared memory configuration 변경
    • cudaDeviceSetCacheConfig

cudaDeviceSetCacheConfig

L1 cache와 shared memory가 동일한 영역을 나눠 사용하기 때문에, L1 cache를 얼마나 사용할 것인지 설정해야 한다. 해당 역할을 cudaDeviceSetCacheConfig로 수행함

Q. stream의 갯수는 어떻게 할당? GPU 개수에 맞게? A. 상관없음. 더 많이 만들어도 되고 그럼

Stream Scheduling

• OpenCL과 다른 부분
• Fermi hardware에는 GPU마다 세개의 엔진 큐가 있다
  • 스트림마다 있는 것이 아니라 GPU마다 하드웨어적으로 구현되어 있음
  • command queue와는 다른 개념
• stream으로 들어간 cuda op들이 dispatch되는 곳
  • Issue vs Dispatch
    • Issue: Reservation station에 operation을 넣는 것
    • Dispatch: Functional unit으로 operation을 가져가는 것
  • 규칙
    • 하나의 stream 내에서는 앞의 op들이 전부 dispatch 되어야 dispatch됨
    • Dispatch 후 실행 중에 resource가 남아있으면 그 다음 op도 dispatch
      • 하는게 규칙이지만 실제로는 남아있지 않는 경우가 많아서 잘 사용되지 않는다
• 엔진 큐 사이의 stream dependency는 유지되나, 엔진 큐 내에서는 유지되지 않는다
  • Stream dependencies between engine queues are maintained but lost within an engine queue
  • 이건 설명을 나중에 해주시겠다고 했는데 제대로 안해주셨고 아직 헷갈린다
  • 특히 하나의 엔진 큐라고 해도 서로 다른 스트림이 아니면 동시실행이 안되기 때문에(스트림 자체가 blocking, in-order 전제) within an engine queue에서 dependency가 사라진다는게 더 이해가 안간다
  • 아마 서로 다른 엔진 큐에 같은 스트림의 op가 들어갈 경우 유지되고, 같은 엔진 큐에 서로 다른 스트림의 op가 들어가면 유지가 안된다는 이야기 같다
• 만약 stream이 다르다면 커널들을 동시에 실행 가능하다
  • 서로 다른 stream에서 H2D, Compute, D2H 각각을 실행하는 경우 이 세가지를 병렬처리 할 수 있다는 뜻
  • 중요
  • 더블 버퍼링 같은 것을 할 수 있기 때문
    • 더블 버퍼링이란? : 데이터를 잘게 잘라서 보내고 계산하고 가져오는 세단계 과정을 반복하는 것.
      • 데이터를 계산하는 동안 그 다음 데이터를 보낼 수 있어서 파이프라이닝이 가능
      • 교수님이 꼭 해보라 하셔서… Double buffering
  • 엔진이 세개있다 (H2D, D2H, Compute) → 따라서 스트림이 서로 다르면 이 세가지를 하나의 GPU에 대해 동시에 할 수 있다.
    • 하드웨어(PCI express bus)가 bidirectional하기 때문에 가능
    • 문제는 메인 메모리의 읽고 쓰기 포트가 따로 되어있어야 성능이 잘 나온다는 것. 보통은 따로 해둔다.

Blocked Queue

Issue order의 중요성

• Issue order == program order
• 빨간 화살표는 signal
  • Signal: 같은 stream의 다음 Op를 실행시켜도 된다는 것을 알려주는 용도
• 예제 1
  • Stream 1: HDa1, HDb1, K1, DH1 (issued first)
  • Stream 2: DH2 (completely independent of stream 1)
• 별표한 곳을 보면 DH1, DH2 간의 dependency는 없어지고, 엔진에 들어온 대로 수행된다!
  • 이것이 within queue, stream dependencies are lost의 뜻
• 큐 자체는 블로킹 큐이기 때문에 순서대로 실행된다
• 예제 2
  • Issue 순서만 바꾼 것
  • 아래와 같이 issue하면 더 concurrent한 실행이 가능하다

Blocked Kernel

여전히 Issue order matters

• 일의 양도 고려해야 한다
  • 
  • 근데 복잡하기만 하고 성능이 더 떨어진다

Concurrent Kernels and Blocking

• Compute engine queue 한정으로 sequentially issue된 마지막 kernel의 종료 지점까지 signal이 delay된다
• K1이 끝나고 시그널을 주는게 아니라 모든 큐의 실행이 끝나고 시그널을 준다
  • Compute queue에 K1, K2, K3가 시퀀셜하게 issue가 되었기 때문
• 더블 버퍼링 할때 이런식으로 잘라서 넣게 되는데 이렇게 하면 안되는거다
• 시그널이 저렇게되는 CUDA(NVIDIA) 에 한정적인 거다. OpenCL(AMD GPU)에도 똑같이 저런 큐가 있는데 시그널링이 다르다. 하드웨어가 그렇게 생겨먹어서 어쩔수가 없다

• issue order를 바꾸면 해결된다
• 그래서 issue order가 매우 중요하다!

Simultaneous Execution in Kernels

• 전체 GPU를 이용하되 서로 다른 커널을 돌릴 수 있게 하는 것
• 그러나 실제로 서로 다른 커널을 동시에 돌리는 경우가 별로 없다. 별로 효과가 없기 때문에
• 그냥 하나의 kernel에 일을 많이 줘서 디바이스 하나를 통째로 하나의 커널로 쓰는게 낫다