글로벌 메모리

글로벌 메모리는 GPU에서 가장 크고 유연한 메모리 영역이다. 이는 모든 스레드가 접근할 수 있는 읽기/쓰기 속성을 가지며, 전역 주소 공간을 사용한다. 글로벌 메모리의 중요한 특징을 정리하면 다음과 같다:

• 접근 범위: 모든 스레드가 접근 가능.
• 사이즈: 매우 큼 (보통 수 기가바이트 수준).
• 속도: 상대적으로 느림 (낮은 대역폭, 높은 지연 시간).
• 영구성: 커널 호출 간에 유지됨. 즉, 여러 커널을 실행하는 동안 데이터를 사용할 수 있다.
• 사용 예시: 다수의 스레드가 데이터에 접근하거나 큰 데이터 세트를 처리할 때 사용.

글로벌 메모리는 CPU의 RAM과 비슷한 역할을 하며, GPU의 모든 연산 유닛이 이에 접근할 수 있다. 이를 통해 대용량 데이터 연산을 효율적으로 처리할 수 있다. 그러나, 글로벌 메모리 접근이 느리기 때문에 자주 접근하는 데이터를 캐시하거나 공유 메모리로 옮기는 등의 최적화가 필요하다.

---

공유 메모리

공유 메모리는 GPU의 SM(Streaming Multiprocessor) 내에서 공유되는 고속 메모리이다. 이는 특정 블록 내의 스레드들이 데이터에 빠르게 접근할 수 있게 해주며, 블록 내에서 데이터를 효율적으로 교환하기 위한 용도로 사용된다. 글로벌 메모리에 비해 속도가 빠르고 지연이 적다.

• 접근 범위: 같은 블록 내의 스레드들만 접근 가능.
• 사이즈: 제한적(일반적으로 몇 십 킬로바이트 수준).
• 속도: 매우 빠름(높은 대역폭, 낮은 지연 시간).
• 영구성: 블록 내에서만 유지됨. 블록이 종료되면 해당 블록의 공유 메모리 데이터는 사라짐.
• 사용 예시: 블록 내의 스레드들이 자주 접근하거나 반복적으로 사용하는 데이터로, 글로벌 메모리에 대한 접근을 최소화하고 싶을 때 사용.

공유 메모리는 데이터를 블록 내에서 효율적으로 공유하고, 글로벌 메모리를 접근하는 횟수를 줄이기 위해 사용된다. 이를 통해 전체 퍼포먼스를 향상시킬 수 있다. 다만, 공유 메모리의 크기가 제한적이므로 신중하게 배열하고 사용하는 것이 중요하다.

접근 패턴에 따른 성능 차이

글로벌 메모리와 공유 메모리의 접근 패턴에 따라 성능 차이가 크게 발생할 수 있다. 적절한 메모리 사용을 위해 접근 패턴과 최적화를 이해하는 것이 중요하다.

• 공유 메모리: 코히런시를 고려하여 사용할 때 성능이 크게 향상됨.
• 글로벌 메모리:
• Coalesced 접근: 여러 스레드가 한 번에 연속된 주소에 접근할 때 성능이 좋다.
• 비 coalesced 접근: 여러 스레드가 불연속된 주소에 접근할 때 성능이 떨어질 수 있다.

메모리 접근 패턴을 최적화하면 GPU의 전체 연산 성능을 최대로 끌어올릴 수 있다.

메모리 최적화 기법

GPU 프로그래밍에서 성능을 극대화하기 위해 다양한 메모리 최적화 기법을 사용할 수 있다. 특히 글로벌 메모리와 공유 메모리의 효율적인 사용이 중요하다.

1. 글로벌 메모리 최적화
2. Coalesced 접근: 여러 스레드가 연속된 글로벌 메모리 주소에 접근하도록 배열하여 메모리 대역폭을 최적화한다.
3. 메모리 정렬: 메모리 접근 주소를 정렬하여 접근 시간을 단축시킨다.

4. 방대한 캐시 사용: 글로벌 메모리에 자주 접근하는 변수나 데이터를 캐시로 옮겨 빠른 접근을 가능하게 한다.

5. 공유 메모리 최적화

6. Bank Conflict 최소화: 공유 메모리는 여러 메모리 뱅크로 구성되어 있으며, 여러 스레드가 동일한 뱅크에 접근할 때 뱅크 충돌이 발생하여 성능이 저하될 수 있음. 이를 피하기 위해 주소를 잘 설계한다.
7. 데이터 재사용: 반복적으로 접근해야 하는 데이터를 공유 메모리에 저장하여 글로벌 메모리 접근을 최소화한다.
8. 적절한 데이터 크기: 공유 메모리 크기를 효율적으로 사용하기 위해 필요 이상의 데이터를 넣지 않도록 설계한다.

코드 예시

공유 메모리를 사용한 예제를 통해 성능 향상을 이해할 수 있다. 예를 들어, 행렬 곱셈에서는 각 블록 내의 데이터를 공유 메모리에 저장하여 글로벌 메모리 접근을 최소화할 수 있다:

__global__ void matrixMulShared(float* A, float* B, float* C, int N) {
    // 블록 및 스레드 인덱스 계산
    int blockRow = blockIdx.y;
    int blockCol = blockIdx.x;
    int row = threadIdx.y;
    int col = threadIdx.x;

    // 공유 메모리 선언
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

    // 결과 초기화
    float Cvalue = 0.0;

    // 각 블록에서의 계산
    for (int i = 0; i < (N / BLOCK_SIZE); ++i) {
        As[row][col] = A[row * N + (i * BLOCK_SIZE + col)];
        Bs[row][col] = B[(i * BLOCK_SIZE + row) * N + col];
        __syncthreads();

        for (int k = 0; k < BLOCK_SIZE; ++k) {
            Cvalue += As[row][k] * Bs[k][col];
        }
        __syncthreads();
    }

    // 결과 저장
    C[row * N + col] = Cvalue;
}

이 코드에서는 행렬 곱셈을 수행하면서 각 블록 내의 데이터들을 공유 메모리에 저장하여 글로벌 메모리 접근을 최소화하고 성능을 향상시킨다.

---

글로벌 메모리와 공유 메모리는 GPU 프로그래밍에서 각기 다른 성능 특성을 가지고 있다. 글로벌 메모리는 큰 용량을 제공하지만 속도가 느리며, 공유 메모리는 빠르지만 제한된 용량을 가지고 있다. 적절한 메모리 사용과 최적화 기법을 통해 GPU의 성능을 최적화할 수 있다.