[MPC] 3. CUDA Thread (1)

3. CUDA Thread (1)

3.1 What is Thread?

3.1.1 Process and Thread

• Process
  • 컴퓨터 프로그램을 실행했을 때, 실행된 프로그램의 인스턴스
  • 프로그램 코드 + 실행 상태 (다음 명령어의 주소, 레지스터 상태, 메모리 컨텐츠)
  • 스토리지 (프로그램) > 메모리 (여러 프로세스 생성) > PC는 프로세스에서 명령어 주소를 읽어오고 데이터는 컴퓨팅 엔진과 레지스터, 메모리 사이에서 주고 받는다.
• Thread
  • 프로세스에 대한 실행 흐름 (execution flow)
  • 실행 상태 (다음 명령어의 주소, 레지스터 상태, 스택)
  • 모든 스레드는 프로그램 코드를 공유한다.
    • CPU의 멀티코어와 달리 GPU의 SM은 각 코어가 레지스터만 갖고 PC를 공유한다.
    • SM의 각 코어는 동일한 명령어를 실행한다.
    • 서로 다른 SM간에는 서로 다른 명령어를 실행한다. 하지만, 프로세스는 공유한다.
    • CPU의 멀티코어는 서로 다른 코어가 서로 다른 PC와 레지스터를 갖는다.

3.1.2 Thread Execution

• On Single Core Processors
  • single thread -> serial processing
  • multiple thread -> time sharing (concurrent processing)
• On Multicore Processors
  • single thread, multiple process -> parallel processing
  • multiple thread -> parallel processing

3.2 CUDA Thread Organization

3.2.1 Hierarchy of Threads

• 커널은 여러 스레드들로 실행된다.
  • 하나의 커널을 런치할 때 생성된 스레드들은 동일한 프로그램을 실행한다.
  • 여러 데이터를 sequential loop이 아닌 병렬로 계산한다.
• 스레드들은 스레드 블록들로 구성된다.
  • 블록 내의 스레드들은 shared memory를 통해 데이터를 공유하는 것이 가능하다.
  • shared memory는 SM에 위치한다.
  • 즉, 하나의 블록은 하나의 SM에 할당해서 실행되니까 그 블록에 포함되어 있는 스레드들은 shared memory를 공유하는 것이 가능하다.
• 스레드 블록은 그리드로 grouping된다.
• 왜 스레드에 계층 구조를 만들었는가?
  • 현실의 많은 데이터들이 계층 구조를 갖기 때문이다.
    • Video file
      • Pixel
      • Image: multiple pixels
      • Video: A sequence of frames

3.2.2 IDs and Dimension

• 각 스레드와 각 블록들은 unique ID를 부여받는다.
• Thread
  • thradIdx: 1D, 2D, or 3D
  • unique within a block
• Blocks
  • blockIdx: 1D, 2D, or 3D
  • a unique within a grid
• 차원은 커널 런치시에 정해진다.
  • blockDim: dimension of block
  • gridDim: dimension of grid
  • 둘 다 pre-defined variable

3.2.3 CUDA pre-defined variables

• Pre-defined variables
  • dim3 gridDim: dimension of grid
  • dim3 blockDim: dimension of block
  • uint3 blockIdx: block index within grid
  • uint3 threadIdx: thread index within block
  • int warpSize: number of threads in warp
• dim3 can take 1,2, or 3 arguments (x,y,z)

3.2.4 Configuring Thread Organization

• 커널은 반드시 execution configuration을 호출해주어야 한다.

3.2.4 Executing threads on GPU

• 각 스레드 블록은 하나의 SM에 할당된다.
  • 동일한 SM에서 실행되는 스레드들은 shared memory를 통해 데이터를 공유한다.
    • shared memory는 global memory보다 훨씬 빠르다.
  • Grid 내에 여러 개의 스레드 블록이 있고, 각 스레드 블록은 여러개의 스레드가 있고,
    각 스레드 블록은 각 SM에 할당되고, 하나의 스레드 블록에 있는 스레드들은 SM내의 각 코어에서 실행된다. 스레드 블록들은 동일한 명령어를 실행하게 되고, 그리드에 포함되어 있는 모든 스레드들은 동일한 커널을 실행하게 된다.

3.3 Extending Vector Addition Program for scalable parallel execution

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<<<ceil(n/256.0), 256>>>(A_d, B_d, C_d,n);

  // 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);
}

• 데이터 사이즈가 하나의 블록에 포함될 수 있는 최대 스레드 갯수를 초과하면 커널이 실행될 수 없다.
• 하나의 블록에서 실행되는 스레드 갯수를 고정하고 블록 갯수를 여러 개 만들어주면 이를 해결할 수 있다.

Extended Vector Addition Kernel

__global__
void addKernel(int* A_d, int* B_d, int* C_d, int n)
{
  // each thread knows its own index
  // int i = threadIdx.x;

    int i= blockDim.x*blockIdx.x + threadIdx.x;
    if(i<n) C_d[i] = A_d[i] + B_d[i];
}

Main Code

int main(void) {
  const int SIZE = 2048;
  int a[SIZE];
  int b[SIZE];
  int c[SIZE];

  for(int i=0;i<SIZE;i++)
  {
    a[i]=i;
    b[i]=i;
  }
  vecAdd(a,b,c, SIZE );

  // print the result
  for(int i=0;i<SIZE;i++)
    printf(“%d\n”, c[i]);

  // done
    return 0;
}

3.4 Matrix Addition

• We want tod add matrix A and matrix B
• We need a two-dimensional kernel

3.4.1 Memory Layout of Matrix

• row-major matrix storage

  • logical layout:
    $$
    \begin{bmatrix}
    a_{00} & a_{01} & a_{02} \\
    a_{10} & a_{11} & a_{12} \\
    a_{20} & a_{21} & a_{22}
    \end{bmatrix}
    $$

  • physical layout: 1D-array
    $$
    \begin{bmatrix}
    a_{00} & a_{01} &a_{02} &a_{10} &a_{11} &a_{12} &a_{20} &a_{21} &a_{22}
    \end{bmatrix}
    $$

  • re-interpret:
    $$
    \begin{bmatrix}
    a_{0} & a_{1} &a_{2} &a_{3} &a_{4} &a_{5} &a_{6} &a_{7} &a_{8}
    \end{bmatrix}
    $$

• index change:

  • idx = y * WIDTH + x (WIDTH = row length of matrix)

3.4.2 matadd-dev.cu

• Thread Orgianization

  • Matrix -> 2D
  • Small size matrix -> a single block

• CUDA kernel

  #include <iostream>
  // kernel program for the device (GPU): compiled by NVCC
  __global__ void addKernel(int* c, const int* a, const int* b) {
    int x = threadIdx.x;
    int y = threadIdx.y;
      int i = y * (blockDim.x) + x; // [y][x] = y * WIDTH + x; 
    c[i] = a[i] + b[i];
  }

• Host code

  int main(void) {
    // host-side data
    const int WIDTH = 5;
    int a[WIDTH][WIDTH];
    int b[WIDTH][WIDTH];
    int c[WIDTH][WIDTH] = { 0 };
  
    // make a, b matrices
    for (int y = 0; y < WIDTH; ++y) {
        for (int x = 0; x < WIDTH; ++x) {
        a[y][x] = y * 10 + x;
        b[y][x] = (y * 10 + x) * 100;
          } 
    }
    // device-side data 
    int* dev_a = 0;
    int* dev_b = 0;
    int* dev_c = 0;
  
    // allocate device memory
    CUDA_CHECK( cudaMalloc((void**)&dev_a, WIDTH * WIDTH * sizeof(int)) );
    CUDA_CHECK( cudaMalloc((void**)&dev_b, WIDTH * WIDTH * sizeof(int)) );
    CUDA_CHECK( cudaMalloc((void**)&dev_c, WIDTH * WIDTH * sizeof(int)) );
    // copy from host to device
    CUDA_CHECK( cudaMemcpy(dev_a, a, WIDTH * WIDTH * sizeof(int), cudaMemcpyHostToDevice) ); CUDA_CHECK( cudaMemcpy(dev_b, b, WIDTH * WIDTH * sizeof(int), cudaMemcpyHostToDevice) ); // launch a kernel on the GPU with one thread for each element.
    dim3 dimBlock(WIDTH, WIDTH, 1); // x, y, z
    addKernel <<< 1, dimBlock>>>(dev_c, dev_a, dev_b); // dev_c = dev_a + dev_b;
    CUDA_CHECK( cudaPeekAtLastError() );
    // copy from device to host
    CUDA_CHECK( cudaMemcpy(c, dev_c, WIDTH * WIDTH * sizeof(int), cudaMemcpyDeviceToHost) ); // free device memory
    CUDA_CHECK( cudaFree(dev_c) );
    CUDA_CHECK( cudaFree(dev_a) );
    CUDA_CHECK( cudaFree(dev_b) );
    // print the result
    for (int y = 0; y < WIDTH; ++y) {
        for (int x = 0; x < WIDTH; ++x) {
      printf("%5d", c[y][x]);
        }
    printf("\n");
    }
    // done
    return 0; 
  }