ChangHyeon Nam's Blog notes and thoughts

CUDA Crash Course | Vector Addition

•

Comments

CUDA Crash Course: Vector Addition를 듣고 정리한 내용이다. 해당 강의에서는 다음의 내용을 다룰 예정이다.

  1. Learn how programs execute in a GPU
  2. Learn the different granularities of threads (granularities of parallelism) in GPUs
  3. Go through the most basic example of Vector Addition in CUDA

C++이나 python에서 2개의 벡터를 더한다고 생각해보자.

Untitled 보통의 경우 Successive addition을 수행한다. 위의 경우 같은 더하기 연산을 네번하게 된다. SIMT 모델에서는 array + array 방식으로, Parralel하게 더하기 연산을 한번하게 된다. SIMT model이 GPU의 model이다.

Untitled 1 Thread들이 모여 warps라는 것을 이루고 있다. Warps는 lockstep에서 명령어를 수행한다. 이는 어떠한 시점 이라도 thread들이 같은 명령어를 동시에 실행한다는 것을 의미한다. (Mask off되는 경우도 있다.)

Untitled 2이러한 warps들이 thread block으로 합쳐진다. 256개의 threads를 프로그래밍 한다고 하면 자동적으로 warps로 해석하고, thread block이 single shader core에 할당된다. 그리고 이것은 three dimensional 구조이다. (X,Y,Z direction으로 이뤄져 있다.) Thread block들이 grid를 이루고 있고, 이 grid는 problem이 GPU로 어떻게 map되는지 나타낸다. 우리는 grid size에 대한 parameter들이 있는데, 각 grid에 몇개의 thread block을 할당할지 혹은 각 차원에는 thread block이 얼마나 있는지에 대한 값들이다.

Untitled 3 Psacal GP 100 GPU이다. SM(Streaming Multiprocessor)에 thread block을 스케줄링 한다.

Untitled 4 각 코어에는 Warp scheduler, Texture caches, l1 caches, shared memory, double precision unit 등이 들어 있다.

Codes

이제 두개의 벡터를 더하는 코드를 작성해보자. Addition of every single element는 완전히 independent하고, vector의 element는 모두 independent하다고 하자. 이러한 경우를 Embarrassingly parallel task라고 한다. (colab에서 작성 코드 링크)

//
// Created by changhyeonnam on 2023/01/10.
//

#include <algorithm>
#include <cassert>
#include <iostream>
#include <vector>

// CUDA kernel for vector addition
// __global__ means this is called from CPU, and runs on the GPU
__global__ void vectorAdd(const int *__restrict a, const int *__restrict b,
                          int *__restrict c, int N){
    // Calculate global thread ID
    // blockDim = 1 dim (just integer)
    int tid = (blockIdx.x * blockDim.x) + threadIdx.x;

    // Boundary check
    if (tid<N)
        // Each thread adds a single element
        c[tid] = a[tid] + b[tid];
}

// Initialize vector of size n to int between 0~99
void matrix_init(int* a, int n){
    for(int i=0; i<n; i++){
        a[i] = rand() % 100;
    }
}
// Check vector add result
void error_check(int* a, int* b, int* c, int n){
    for(int i=0; i<n; i++){
        assert(c[i] == a[i] + b[i]);
    }
}

// print vector add result
void print_result(int* a, int* b, int* c, int n){
    for(int i=0; i<n; i++){
        if(i%100==0)
            std::cout<<"c["<<i<<"]="<<c[i]<<" = "<<"a["<<i<<"]="<<a[i]<<" + " <<"b["<<i<<"]="<<b[i]<<'\n';
    }
}


int main(){
    // Vector size of 2^16 (65536 elements)
    int n = 1<<16;

    // Host vector pointers
    int *h_a, *h_b, *h_c;

    // Device vector pointers
    int *d_a, *d_b, *d_c;

    // Allocation size for all vectors
    size_t bytes = sizeof(int) * n;

    // Allocate host memory
    h_a = (int*)malloc(bytes);
    h_b = (int*)malloc(bytes);
    h_c = (int*)malloc(bytes);

    // Allocate device(gpu) memory
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    /* There is something called unified memory.
     * one set of memory that gets migrated between the GPU and CPU viceversa.
     * [next lecture]
     */

    // Initialize vectors a and b with random values between 0 and 99
    matrix_init(h_a, n);
    matrix_init(h_b, n);

    // Copy data from the CPU(HOST) to the GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // Threadblock size
    // it's generally good to do this a size of 32 because these have to translate it to warps.
    // which are of size 32.
    int NUM_THREADS = 256;

    // Grid size
    // NUM_THREAD * NUM_BLOCKS = NUMBER of Elements.
    int NUM_BLOCKS = (int)ceil(n/NUM_THREADS);

    // Launch kernel on default strem w/o
    vectorAdd<<<NUM_BLOCKS, NUM_THREADS>>>(d_a, d_b, d_c, n);

    // Copy sum vector from device to host
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // Check result for errors
    error_check(h_a, h_b, h_c, n);
    print_result(h_a, h_b, h_c, n);
    printf("COMPLETED SUCCESFULLY\n");
    return 0;
}