공유 메모리

공유 메모리 사용 방법은 크게 세 가지 케이스로 구분한다.

  • L1 캐시: 자주 사용되는 데이터를 직접 분류, 관리하기 어려운 경우
  • 사용자 관리 캐시 1: 개발자가 커널 내 알고리즘의 데이터 접근 패턴을 파악 후, 직접 제어
  • 사용자 관리 캐시 2: 자주 사용하는 데이터의 전역 메모리 접근을 줄이기 위함

스레드 간 공유 메모리와 L1 캐시 활용 방법

  1. 공유 메모리 (Shared Memory)
  • 역할 및 특징: 공유 메모리는 각 블록 내 모든 스레드가 접근할 수 있는 고속 메모리 공간이다. 스레드 간 데이터를 공유하기 위해 중요한 메커니즘으로 활용된다.
  • 동기화의 필요성: 여러 스레드가 공유 메모리에서 데이터를 읽고 쓰는 과정에서 데이터의 일관성을 유지하기 위해 동기화가 필요하다. 이를 위해 __syncthreads() 함수가 사용되며, 이 함수는 모든 스레드가 특정 지점에 도달할 때까지 대기하게 한다.
  1. L1 캐시
  • 역할 및 특징: L1 캐시는 GPU가 자동으로 관리하는 고속 메모리로, 최근 접근한 데이터를 저장해 이후 접근 시 메모리 접근 시간을 줄인다.
  • 캐시 히트와 미스:
    • 캐시 히트: 필요한 데이터가 캐시에 있을 경우, 신속하게 데이터를 가져올 수 있어 성능이 향상된다.
    • 캐시 미스: 필요한 데이터가 캐시에 없을 경우, 메모리에서 데이터를 가져와야 하므로 레이턴시가 증가하고 성능이 저하된다.
  • 최적화 고려사항: 공간적 및 시간적 지역성을 고려한 데이터 접근 패턴을 설계하면, 캐시 히트율을 높여 성능을 최적화할 수 있다.
  1. SIMT 구조와 메모리 접근 패턴
  • SIMT 구조: GPU는 SIMT (Single Instruction, Multiple Threads) 구조로, 워프(warp) 단위로 스레드들이 동시에 실행된다.
  • 메모리 접근 최적화:
    • 효율적인 접근: 워프 내의 스레드들이 동일한 데이터 블록에 접근할 경우, 메모리 트랜잭션 수가 최소화되어 성능이 향상된다.
    • 비효율적 접근: 스레드들이 여러 데이터 블록에 분산되어 접근하면, 메모리 트랜잭션 수가 증가해 성능이 저하될 수 있다.
  1. 사용자 관리 캐시 (User-managed Cache)
  • 역할 및 특징: 개발자가 공유 메모리나 레지스터를 활용해 데이터를 수동으로 관리하는 것을 의미한다. 이는 자동으로 관리되는 L1 캐시와 달리, 알고리즘의 특성에 맞게 데이터를 직접 제어하는 방식이다.
  • 활용 예시: 예를 들어, 행렬 곱셈에서 공유 메모리를 활용하여 데이터를 재사용함으로써 메모리 대역폭을 효율적으로 사용할 수 있다. 이를 통해 커널의 성능을 크게 최적화할 수 있다.

공유 메모리 사용 예제 1: 작은 행렬

하나의 블록만 사용는 행렬 곱셉 커널은 아래와 같다.
각 스레드의 인덱스는 행렬 C에서 동일 인덱스의 원소를 연산한다.
int index는 행렬 C의 원소들을 1차원 형태 인덱스로 표현한 변수이다.

__global__ void matMul_kernel(float* A, float* B, float* C) {
    int row = threadIdx.x;
    int col = threadIdx.y;
    int index = row * blockDim.y + col;
    
    float result = 0;
    for (int k = 0; k < K_SIZE; k++) {
        result += _A[row * K_SIZE + k] * _B[col + k * COL_SIZE];
    }
    _C[index] = result;
}

공유 메모리를 사용자 정의 캐시로 활용하는 핵심은 블록 내 스레드들이 자주 사용하는 데이터들을 공유 메모리에 저장하는 것이다. 이를 통해 전역 메모리의 접근 수를 줄인다. 위 코드에서 행렬 C의 (row, col) 원소를 연산하기 위해서는, 행렬 A의 row와 행렬 B의 col을 한 번씩 읽기가 있다. 그리고 그 결과값을 저장하는 1번의 쓰기가 발생한다.

그런데 행렬 A의 row는 C(row, 0)에서 C(row, col)까지 연산하는 동안 계속 접근이 필요하다. 따라서 행렬 A의 row는 col번 만큼 접근이 필요하다. 마찬가지로 B의 col은 row번 만큼 접근이 필요하다. 따라서 두 가지 결론이 나온다.

  • 행렬 A, B의 데이터는 반복해서 접근한다.
  • 행렬 C는 데이터를 1번만 저장하면 된다.

자주 사용하는 데이터를 공유 메모리에 저장하는 것이 중요하다고 하였다. 따라서 행렬 A, B를 공유 메모리에 저장하는 것이 성능에 유리하다. (현재 문제는 공유 메모리 크기보다 작은 행렬이므로 공유 메모리 크기 자체는 고려하지 않는다.)

공유 메모리를 사용하는 행렬 곱셉 커널은 아래와 같다.

__global__ void matMul_kernel(float* A, float* B, float* C) {
    int row = threadIdx.x;
    int col = threadIdx.y;
    int index = row * blockDim.y + col;

    __shared__ float sA[ROW_SIZE][K_SIZE];
    __shared__ float sB[K_SIZE][COL_SIZE];

    if (threadIdx.x == 0 && threadIdx.y == 0) {
        for (int r = 0; k < ROW_SIZE; k++) {
            for (int k = 0; k < K_SIZE; k++) {
                sA[r][k] = _A[r * K_SIZE + k];
            }
        }

        for (int c = 0; k < ROW_SIZE; k++) {
            for (int k = 0; k < K_SIZE; k++) {
                sB[k][c] = _B[c + k * K_SIZE];
            }
        }
    }

    __syncthreads();
    ...
    
    float result = 0;
    for (int k = 0; k < K_SIZE; k++) {
        result += _A[row * K_SIZE + k] * _B[col + k * COL_SIZE];
    }
    _C[index] = result;
}

코드에서는 블록 내에서 threadIdx.x == 0 && threadIdx.y == 0 인 스레드 하나만이 공유 메모리에 데이터를 복사하도록 설정되어 있다. 이 경우, 다른 스레드들은 이 복사가 끝날 때까지 __syncthreads() 에서 대기하게 된다. 이는 비효율적인 방법이다. 복사 작업을 단 하나의 스레드에게만 맡기면, 전체 블록이 이 작업이 끝날 때까지 기다리게 되어 병렬 처리의 장점을 제대로 살리지 못한다.

더 많은 스레드가 복사 작업에 참여하도록 하여 병렬로 데이터를 복사하게 하면, 복사 시간을 단축할 수 있다. 하지만 모든 스레드가 자신과 관련된 (row, col) 데이터를 복사하게 되면 중복 복사 문제 가 발생할 수 있다. 예를 들어, 행렬 A의 동일한 행(row)을 복사할 때 여러 스레드가 동일한 데이터를 공유 메모리에 중복으로 복사하면 메모리 사용이 비효율적이 된다. 이러한 중복 복사는 메모리 대역폭을 낭비하고, 연산 지연을 초래한다.

이 문제를 해결하기 위해서는 각 스레드가 서로 다른 부분의 데이터를 복사하도록 작업을 분배해야 한다. 특정 스레드들끼리 복사할 부분을 지정하는 방식으로 중복 복사를 피할 수 있다. 이렇게 하면 모든 스레드가 병렬로 데이터를 복사해 공유 메모리 초기화 시간을 단축할 수 있다. 아래 코드가 그 예시이다.

__global__ void matMul_kernel(float* A, float* B, float* C) {
    ...

    __shared__ float sA[ROW_SIZE][K_SIZE];
    __shared__ float sB[K_SIZE][COL_SIZE];

    if (col == 0) { // read matrix A
        for (int k = 0; k < K_SIZE; k++) {
            sA[row][k] = _A[row * K_SIZE + k];
        }
    }

    if (row == 0) { // read matrix B
        for (int k = 0; k < K_SIZE; k++) {
            sB[k][col] = _B[col + k * COL_SIZE];
        }
    }

    __syncthreads();
    ...
}

공유 메모리보다 작은 크기의 데이터를 활용할 때에는 데이터 전체를 공유 메모리에 올려놓아도 문제가 없다. 하지만 공유 메모리보다 큰 데이터를 처리하기 위해서는 또 다른 전략이 필요하다.

공유 메모리 사용 예제 2: 큰 행렬

작성 예정