ಣ ᴄᴏᴍᴘᴜᴛᴇʀ ᴄɪᴛʏ ഒ

Search

SearchSearch

SHPC 14 - CUDA

Dec 31, 20249 min read

Heterogeneous Computing

  • Host
    • The CPU and its memory (host memory)
  • Device
    • The GPU and its memory (device memory)
    • 요즘에는 ARM, RISC-V를 GPU에 넣는 경우도 있는데, 그럴때는 GPU가 마치 peripheral device처럼 동작하는 것이다.

Programs

  • Host program
    • 커널 실행 담당
  • Kernels
    • Compute device에서 실행가능한 코드의 최소단위
    • 한번 실행되면 많은 동일한 인스턴스들이 생성되어 data parallelism을 exploit할 수 있다
  • Host program과 kernel은 아래 그림처럼 병렬로 실행된다

Execution process

  1. CPU 메모리에서 input data를 복사하여 GPU 메모리로 옮긴다
  2. GPU 코드를 로드해서 실행한다 (요때는 On-chip cache를 쓴다)
  3. GPU memory에서 CPU memory로 결과를 복사한다

“Hello World! in CUDA”

  • A program with no device code
  • NVIDIA compiler (nvcc) 를 이용해 컴파일
    • Host code와 Device code를 분리하여 Host code는 gcc, Device code는 NVIDIA 컴파일러를 사용해 컴파일한다

“Hello World!” with Device Code

  • __global__ 은 CUDA C/C++ 에서 사용하는 키워드로, host code 에서 부르는 device의 함수
  • <<<>>> 은 host code에서 device code로의 call을 의미
    • Kernel launch 라고도 함
    • 1, 1은 블록 갯수로 1개, 블록 크기로 스레드 1개를 사용한다는 의미
  • 위 코드에서는 device가 하는건 아무것도 없다

A Simple Kernel to Add Two Integers

Device kernel code

  • Device에서 동작

Host code

  • cudaMalloc GPU를 위한 메모리(*a, *b, *c) 할당
    • 메모리를 할당하여 세개의 포인터를 CUDA의 디바이스에게 넘겨줌
    • 중요한 것은 실제로 주소공간 자체가 호스트와 디바이스가 따로 있기 때문에, 메모리 allocation을 할 때 양쪽을 할당한 다음 호스트 쪽에 값을 할당하고 디바이스쪽으로 copy(cudaMemcpy())를 해야 한다.
      • 여기서 alloc된 메모리는 pagable한데, sync 방식이라 괜찮다
      • 아니라면 pinned memory 할당을 위해 cudaMallocHost 혹은 cudaHostAlloc을 쓰고, 해제를 위해 cudaFreeHost()를 써야 함
    • *a, *b, *c를 GPU에게 넘겨줄 때는 디바이스쪽 주소를 넘겨야 한다 (d_a, d_b, d_c)
  • 디바이스 포인터 vs 호스트 포인터
    • 디바이스 포인터
      • GPU 메모리를 가리킴
      • 호스트 코드에서 전달 가능하지만, 직접 참조(dereference는 불가능하다)
    • 호스트 포인터
      • CPU 메모리를 가리킴
      • 장치 코드에서 전달 가능하지만, 직접 참조는 불가능
  • 메모리 관리 함수: cudaMalloc(),cudaFree(),cudaMemcpy()
    • cudaMemcpy의 플래그
      • cudaMemcpyHostToDevice
        • 호스트 → 디바이스 메모리 복사
      • cudaMemcpyDeviceToHost
        • 디바이스 → 호스트 메모리 복사
    • add<<1, 1>> : 그리드(커널 인덱스 스페이스) 내에 블록 및 쓰레드가 각각 하나
      • 그리드가 전체 인덱스 스페이스
      • 블록은 쓰레드의 집합으로, 워크 그룹과 같은 것
      • 쓰레드는 하나의 워크 아이템 같은 것
      • 즉 본 예제에서는 쓰레드가 하나인 하나의 블록만 실행하는 예제

Running add() in Parallel

  • add() 를 N개 병렬적으로 돌린다고 생각해 보자
  • add()의 parallel invocation은 block이라고 불림

  • blockIdx 는 약속된 변수로, 1차원 변수의 경우 x, 2차원 변수의 경우 x, y, 3차원 변수의 경우 x, y, z … 를 사용한다

  • blockIdx == get_group_id()

  • threadIdx == get_local_id()

  • blockIdx.x * blockDim.x + threadIdx.x == get_global_id()

  • 여기서는 그냥 group id 랑 global id가 같아서 관계 없음

  • SPMD를 사용한다

  • 실제 코드

CUDA Execution Model

  • Thread (Work item)
    • Sequential execution unit
    • 병렬로 실행되고, 모든 쓰레드가 동일한 sequential program을 실행함
    • 실제로는 stack이 없기에 아키텍처적인 쓰레드라고 보기는 어렵다
  • Threads Block (Work group)
    • A group of threads
    • 동일한 SM 내에서 실행됨
    • 같은 쓰레드 블록의 경우 협력 가능
      • 경량 동기화: 동기화 배리어를 이용
      • 데이터 교환 (아마도 shared memory를 이용해서)
  • Grid
    • 여러 스레드 블록으로 구성
    • 여러 SM에 걸쳐 실행됨
    • 글로벌 메모리를 통해 블록 간 통신
      • 비용이 많이듬

매핑

  • Grid → GPU : 하나의 그리드가 하나의 GPU에 대응됨
  • Block → SM : 각 블록은 SM에 매핑되어 처리됨 (OpenCL에서 워크그룹 단위로 SM에 일을 assign하는 것과 마찬가지)
  • Thread → SP : 블록 내 스레드는 SM의 SP(Scalar processor, PE인듯)에서 실행됨
  • Warp: 각 블록의 스레드는 Warp 단위로 실행(하나의 SM이 여러 warp를 돌림)
    • Warp는 32개의 스레드로 구성

커널 실행

  • 다음과 같은 구조로 호출
    • myKernel<<<B, T>>>(arg1, ... )
  • 차원에 따라 다르게 호출 가능. 예를 들어 2차원인 경우,
dim3 grid(4, 4);
dim3 block(16, 16);
myKernel<<<grid, block>>>();
  • 이런식으로 호출 가능하다
    • 그리드는 4 * 4 = 16개 블록의 그리드
    • 블록은 16 * 16 = 256 스레드
    • 총 스레드는 16 * 256 = 4096 스레드

Built in variables

  • blockIdx.x, blockIdx.y, blockIdx.z : 그룹 아이디
  • threadIdx.x, threadIdx.y, threadIdx.z : 로컬 아이디 (블록 내에서의 offset)
  • blockDim.x, blockDim.y, blockDim.z : 블록을 구성하는 스레드 수

Running add() in Multiple Threads

  • 커널코드가 동일할때, 아래와 같이 바꿔주면 된다. 아까와 다르게 하나의 SM에서 실행된다.
  • 쓰레드 개수가 적을 때는 성능이 비슷할 텐데, 쓰레드 갯수가 많으면 SM을 놀게하는거기 때문에 별로 좋지 않다
    • 어차피 하나의 블록(SM)이 한번에 실행하는 쓰레드의 수에는 한계(32개)가 있기 때문이다. 다른 SM이 동시실행을 해줘야 함
    • 그러나 블록 간 동기화나 통신이 필요할 때에는 유리할 수 있겠음. 블록 내의 스레드는 shared memory를 공유하기 때문

Multiple Threads and Multiple Blocks

  • CUDA에서는 글로벌 인덱스라는게 없기 때문에, 아래와 같이 계산해야 한다

  • 즉 SM에 공평하게 로드를 나눠주고 싶으면 위 코드에 아래와 같은 호스트 코드를 붙여줌

Handling Arbitrary Vector size

  • vector가 1023 같은 식이면 어케할 것인가 → 실제로 개수가 1023이 아니라 1024가 되면, 0~ 1023 이 index일 것. 이걸 다루는 코드로 1021개 백터를 가진 애를 핸들링한다고 해보자. 1021, 1022, 1023 을 접근할 때는 seg fault가 나니 보호해 줘야 한다. 따라서 n을 아규먼트로 넘겨줘서 보호해 주는 코드
  • 문제는 블록이 가득 차 있고, 마지막에 저 세개를 담당하는 애가 있으면 걔는 전부 if의 false가 나올테니, branch divergence가 생기게 된다. 아주 약간(마지막 블록만 손해보는 거니)의 손해가 발생한다.
    • Branch divergence를 줄이는 것이 참 중요하다
  • 모든 함수들은 기본적으로 blocking으로 동작한다 (memcpy)
    • 이상한데? 왜냐하면 내가 찾아봤을 때는 비동기도 있다고 나왔기 때문
    • 의문 해결 ⇒ SHPC 18 - CUDA Streams

Graph View

Backlinks