[MPC] 2. Fundamentals of CUDA (2)

2. Fundamentals of CUDA (2)

2.1 A Vector Addition Kernel

Review: Data-level Parallelism

• 서로 다른 데이터에 동일한 명령어 흐름을 병렬적으로 실행하는 것
  • Image processing: 이미지 내의 독립적인 픽셀을 다룬다.
  • Molecular dynamics: 여러 원자간의 상호작용에 대해 시뮬레이션한다.
  • Airline scheduling: 수천개의 비행기, 승무원, 게이트들을 다뤄야한다.

2.1.1 Vector Addition

• Scalar vs Vector
  • Scalar: a single number
  • Vector: an array of numbers
    • Vector: 1D array로 표현되는 것
    • Vector addition: 1d array끼리 더하는 것

CPU-based Vecotr Addition

• 순차적으로 덧셈을 처리한다.

CPU-based Vector Addition

CUDA-based Vector Addition

• GPU cores가 동시에 덧셈을 한다. => 병렬 실행

CUDA-based Vector Addition

Review: CUDA Programming Model

• Host: CPU + main memory (host memory)
• Device: GPU + video memory (device memory)
• GNU gcc: linux c compiler
• nvcc: NVIDA CUDA compiler

Review: Execution of CUDA Program

• CUDA Program: host/device code로 구성돼있다.
  • host code에는 순차적으로 실행되거나 약간의 병렬성이 있는 것을 실행시킨다.
  • device code (kernel) 에는 병렬성이 높은 것을 실행시킨다.
• 커널은 스레드 수에 의해 런치된다.
  • grid: 커널 런치에 의해 만들어진 모든 스레드의 그룹

CUDA-based Vector Addition Code

#include <cuda.h>
void vecAdd(int* A, int* B, int* C, int n)
{
  int size = n * sizeof(int);
  int* A_d=0;
  int* B_d=0;
  int* C_d=0;
  // Allocate device memory cudaMalloc((void **) &A_d, size); cudaMalloc((void **) &B_d, size); cudaMalloc((void **) &C_d, size);
     // Transfer A and B to device memory
  cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
  cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);
  // Kernel invocation code – to be shown later
  ...
     // Transfer C from device to host
  cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);
     // Free device memory for A, B, C
  cudaFree(A_d); cudaFree(B_d); cudaFree (C_d);
}

2.2 Kernel Functions and Threading

2.2.1 CUDA Function Declarations

• Compilation unit: Functions - 특정 함수는 CPU에 특정 함수는 GPU에.. > 이를 잘 표시해야됨

  • 각 함수는 CPU나 GPU 혹은 둘 다에 할당될 것이다.

• How to distinguish them?

  • 각 함수에 대한 PREFIX를 쓴다.

    __device__ float DeviceFunc() // 다른 gpu에서 호출되는 함수, 즉 cpu에선 호출될 수 없음.
    __global__ void KernelFunc() // cpu에서 호출되고 gpu에서 실행되는 함수. 반드시 void를 반환해야 된다.
    __host__ float HostFunc() // cpu에서 호출되고 실행되는 함수, 적지 않아도 상관없음.
    
    // __device__ : device에서 실행되고 호출되는 함수
    // __global__ : host에서 호출되고 device에서 실행되는 함수
    // __host__ : host 실행되고 호출되는 함수

Example of a Kernel: Vector Addition Kernel

// Compute vector sum C = A+B
// Each thread performs one pair-wise addition
__global__ // cpu에서 호출되는 함수, 하지만 gpu에서 실행된다.
void addKernel(int* A_d, int* B_d, int* C_d)
{
     // each thread knows its own index
    int i = threadIdx.x;
    C_d[i] = A_d[i] + B_d[i]; 
  // cpu-based code에서는 idx를 구하기 위해 반복문을 사용했지만, gpu kernel에서는 이 부분이 없다.
  // threadIdx를 array에 대한 index로 사용했다.
}

• threadIdx: 각 스레드에 고유한 coordinate를 주는 빌트인 variable

2.2.2 CUDA Kernel Function and Threading

• 커널을 실행하게 되면 여러 개의 스레드가 생성되는데, 모든 스레드는 서로 다른 코어에서 동일한 커널을 병렬적으로 실행할 것이다.

  • SPMD: Single Program Multiple Data

• Warp내의 모든 스레드는 SM의 서로 다른 코어에서 같은 명령어를 병렬적으로 실행할 것이다.

  • SIMT: Single Instruction Multiple Thread

• SPMD vs SIMT
  : 하나의 SM에 있으면 SM에 할당된 블록은 워프 단위로 나뉘어져서, 워프 단위로 하나의 SM에서 실행된다. 워프에 속해있는 스레드들은 SM에서 실행될 때, 그 스레드는 동일한 명령어를 실행한다. 서로 다른 SM에 할당된 블록에 속해있는 스레드는 동일한 명령어를 실행하지만 특정 순간에 다른 명령어를 실행할 수 있다.

  예) SM1에 속해 있는 스레드는 동일한 명령어를 실행해야 한다. SM1, SM2가 동일한 커널을 실행하지만 (SPMD), SM1,SM2가 특정 시점에 서로 다른 명령어를 실행할 수 있다.

• host code가 커널을 런치할 때 grid of threads가 생성된다.

  • 각 grid는 스레드 블록의 array로 구성된다. (aka. block)
  • gird의 모든 블록은 같은 사이즈를 갖는다.
  • 각 스레드 블록에서 스레드의 총 갯수는 커널을 런칭하는 host code에서 특정된다.
    • 동일한 커널은 서로 다른 스레드 갯수를 갖고 런치될 수 있다.

• 커널을 실행하면, 하나의 grid가 생성되고, grid는 여러개의 스레드블록으로 구성되고, 각 스레드 블록은 여러개의 스레드로 구성된다. 커널을 실행할 때 몇 개의 스레드 블록을 생성할 건지, 각 스레드 블록에 몇개의 스레드 블록을 사용할 것인지는 커널을 런치할 때 정해줄 수 있다.

2.3 Kernel Launch

2.3.1 CUDA Kernel Launch

• Kernel launch syntax

  • addKernel <<<1, SIZE>>>(A_d, B_d, C_d);

• CUDA view

  • a thread executes addKernel() with threadIdx.x = 0
  • a thread executes addKernel() with threadIdx.x = 1
  • a thread executes addKernel() with threadIdx.x = 2
    ...
  • a thread executes addKernel() with threadIdx.x = SIZE - 1

2.3.2 A Complete Version: vecAdd.cu

Device Code (Kernel)

#include <cuda.h>
#include <iostream>

// Compute vector sum C = A+B
// Each thread performs one pair-wise addition
__global__
void addKernel(int* A_d, int* B_d, int* C_d)
{
  // each thread knows its own index
   int i = threadIdx.x;
  C_d[i] = A_d[i] + B_d[i];
}

Host Code

void vecAdd(int* A, int* B, int* C, int n)
{
  int size = n * sizeof(int);
  int* A_d=0;
  int* B_d=0;
  int* C_d=0;
  // Allocate device memory 
  cudaMalloc((void **) &A_d, size); 
  cudaMalloc((void **) &B_d, size); 
  cudaMalloc((void **) &C_d, size);

  // Transfer A and B to device memory
  cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
  cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);

  addKernel<<<1, n>>>(A_d, B_d, C_d);

  // Transfer C from device to host
  cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);
  // Free device memory for A, B, C
  cudaFree(A_d); cudaFree(B_d); cudaFree (C_d);
}