소프트웨어 최적화 기법

개요

GPU 드라이버 및 소프트웨어 인터페이스는 그래픽 처리 장치(GPU)의 성능을 극대화하기 위해 다양한 소프트웨어 최적화 기법을 사용한다. 소프트웨어 최적화는 프로그래머가 차용할 수 있는 기법으로, GPU의 처리 능력을 최대한 활용할 수 있도록 코드를 작성하는 것이다.

1. 병렬 프로그래밍

GPU는 다수의 코어를 가지고 있어 병렬 프로그래밍에 적합한다. 이는 대규모 연산을 동시에 수행할 수 있도록 해준다. 병렬 프로그래밍은 대규모 데이터를 빠르게 처리하는 데 큰 도움을 준다.

__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N) 
        C[i] = A[i] + B[i];
}

2. 메모리 최적화

GPU 메모리는 크게 글로벌 메모리, 쉐어드 메모리, 레지스터 메모리로 나뉜다. 각 메모리의 접근 속도와 초기화 방법에 따라 최적화가 필요하다.

글로벌 메모리 접근 패턴 최적화: GPU에서 글로벌 메모리는 큰 레이턴시를 가지므로, 연속적인 메모리 접근 패턴을 사용하여 이를 최소화해야 한다.
공유 메모리 사용: 글로벌 메모리보다 빠른 접근 속도를 가지는 쉐어드 메모리를 사용하여, 동일 블록 내에서 데이터를 공유한다.

__global__ void redKernel(float *input, float *output) {
    __shared__ float s_data[BLOCK_SIZE];
    int thIdx = threadIdx.x;
    int gIdx = blockIdx.x * blockDim.x + thIdx;

    s_data[thIdx] = input[gIdx];
    __syncthreads();

    for(unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
        if(thIdx % (2 * stride) == 0) {
            s_data[thIdx] += s_data[thIdx + stride];
        }
        __syncthreads();
    }

    if (thIdx == 0) 
        output[blockIdx.x] = s_data[0];
}

3. 쓰레드 활용

효율적인 쓰레드 관리와 동기화는 GPU 성능에 중요한 영향을 미친다.

워프 다이버전스 최소화: 워프 내의 쓰레드가 서로 다른 코드를 실행하면 성능 저하가 발생하기 때문에, 조건문 처리를 최소화해야 한다.
동기화: 쓰레드 간의 데이터 정합성을 유지하기 위해 동기화를 적절히 사용해야 한다. 그러나 과도한 동기화는 성능을 저하시킬 수 있으므로, 최소한으로 사용하는 것이 좋다.

__global__ void kernel() {
    int tid = threadIdx.x;
    if (tid < 32) { // Avoid warp divergence
        // Do something
    }
}

4. 연산 최적화

연산 병합: 다수의 반복문에서 이루어지는 연산을 하나의 명령으로 병합하여 성능을 개선할 수 있다.
루프 펼치기(loop unrolling): 반복문이 많은 경우, 컴파일러가 루프를 펼칠 수 있도록 코드를 작성하여 오버헤드를 줄일 수 있다.

__global__ void add() {
    int idx = threadIdx.x;

    // Without loop unrolling
    for (int i = 0; i < N; i++)
        C[idx * N + i] = A[idx * N + i] + B[idx * N + i];

    // With loop unrolling
    C[idx * N]     = A[idx * N]     + B[idx * N];
    C[idx * N + 1] = A[idx * N + 1] + B[idx * N + 1];
    // Continue for the loop unrolling
}

5. 데이터 로컬리티

데이터 로컬리티는 데이터가 메모리에서 서로 인접하게 배치되어 있으면 메모리 접근 성능이 개선되는 것을 의미한다.

__global__ void compute(float *data, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        float local_data = data[idx];       // Load data to local variable
        // Perform computation
        data[idx] = local_data * 2.0f;      // Store result back to global memory
    }
}

6. 매트릭스 최적화

매트릭스 곱셈과 같은 벡터 연산에서 타일링 기법을 사용하여, 메모리 접근 패턴을 최적화 할 수 있다.

__global__ void matMul(float *A, float *B, float *C, int n) {
    __shared__ float s_A[TILE_SIZE][TILE_SIZE];
    __shared__ float s_B[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float temp = 0.0;

    for (int i = 0; i < (n - 1) / TILE_SIZE + 1; ++i) {
        if (row < n && (i * TILE_SIZE + threadIdx.x) < n) 
            s_A[threadIdx.y][threadIdx.x] = A[row * n + i * TILE_SIZE + threadIdx.x];
        else 
            s_A[threadIdx.y][threadIdx.x] = 0.0;

        if (col < n && (i * TILE_SIZE + threadIdx.y) < n) 
            s_B[threadIdx.y][threadIdx.x] = B[(i * TILE_SIZE + threadIdx.y) * n + col];
        else 
            s_B[threadIdx.y][threadIdx.x] = 0.0;

        __syncthreads();

        for (int j = 0; j < TILE_SIZE; ++j) 
            temp += s_A[threadIdx.y][j] * s_B[j][threadIdx.x];

        __syncthreads();
    }

    if (row < n && col < n) 
        C[row * n + col] = temp;
}