junyeokk
Blog
Parallel Computing·2025. 10. 23

Shared memory란

Shared Memory란 무엇인가

Shared memory는 GPU 칩 내부에 있는 빠른 메모리다. 같은 블록 내의 모든 스레드가 공유한다.

하드웨어 관점에서 보면 shared memory는 SM(Streaming Multiprocessor) 내부에 위치한다. CPU의 L1 캐시처럼 칩 위에 직접 올라가 있어서 접근 속도가 매우 빠르다. 1 clock cycle에 데이터를 읽을 수 있다.

반면 global memory는 GPU 칩 외부의 DRAM에 있다. 칩과 DRAM 사이를 왔다 갔다 해야 하므로 100+ clock cycle이 걸린다. 거리로 비유하면 shared memory는 책상 위의 노트고, global memory는 도서관까지 가서 빌려오는 책이다.

하지만 shared memory는 용량이 작다. 블록당 48KB 정도만 사용할 수 있다. Global memory는 GB 단위지만 느리고, shared memory는 KB 단위지만 빠르다.

Shared Memory를 사용하는 이유

데이터 접근 패턴에 따라 적절한 메모리를 선택하면 성능을 크게 향상시킬 수 있다. 스레드 내에서만 읽고 쓰는 데이터는 레지스터를 사용한다. 블록 내에서 공유하는 데이터는 shared memory를 사용한다. 읽기 전용 데이터는 constant memory를 사용한다. 스레드별로 인덱싱된 읽기/쓰기는 local memory를 사용하지만 느리다. 입출력 데이터는 cudaMalloc으로 할당한 global memory를 사용한다.

핵심은 데이터 재사용이다. 같은 데이터를 여러 번 읽어야 한다면 global memory에서 한 번 가져와서 shared memory에 넣어둔다. 그 다음부터는 빠른 shared memory에서 읽는다.

이는 일종의 분할 정복 접근법이다. 큰 데이터를 shared memory에서 처리할 수 있는 작은 단위로 나눈다.

Shared Memory 사용 3단계

Shared memory를 효과적으로 사용하려면 3단계 과정을 거친다.

첫 번째, 협력적으로 데이터를 shared memory로 복사한다. global memory에 있는 원본 데이터를 블록 내 모든 스레드가 협력하여 shared memory로 가져온다.

두 번째, 같은 블록 내 스레드들이 shared memory의 데이터를 사용하고 업데이트한다. 이 단계에서 실제 계산이 일어난다. Shared memory는 빠르기 때문에 반복적인 접근에도 성능 저하가 없다.

세 번째, 협력적으로 결과를 global memory로 복사한다. 계산이 끝난 데이터를 다시 global memory에 쓴다.

Race Condition 문제

Shared memory를 사용할 때 race condition이 발생할 수 있다. 한 스레드가 데이터를 쓰기 전에 다른 스레드가 읽으려고 하면 잘못된 값을 얻는다.

plain
__global__ void kernel_func(void) {
    __shared__ float shared[SIZE];

    shared[i] = value;  // 모든 스레드가 실행

    another = shared[j];  // 다시 모든 스레드가 읽기
}

위 코드의 문제를 구체적인 예로 살펴보자. 블록에 1024개 스레드가 있다고 가정한다.

스레드 0번이 shared[0] = value를 실행하고 바로 다음 줄로 넘어간다. 그런데 스레드 500번은 아직 shared[500] = value를 실행하지 못했다. 이 상태에서 스레드 0번이 another = shared[500]을 실행하면 어떻게 될까? 아직 쓰지 않은 쓰레기 값을 읽게 된다.

문제가 더 복잡한 이유는 GPU가 warp 단위로 실행하기 때문이다. Warp는 32개 스레드가 묶인 단위다. Warp #0 (스레드 0-31)이 값을 업데이트하는 동안, Warp #2 (스레드 64-95)가 이미 읽기를 시도할 수 있다. 하지만 Warp #31 (스레드 992-1023)은 아직 업데이트조차 시작하지 않았을 수 있다.

즉, 각 스레드의 실행 속도가 다르다. 어떤 스레드는 빠르게 진행하고, 어떤 스레드는 느리게 진행한다. Shared memory를 통해 데이터를 공유하려면 모든 스레드가 같은 단계에 있다는 보장이 필요하다.

Barrier Synchronization

__syncthreads() 함수로 race condition을 해결한다.

Intrinsic Function이란?
컴파일러가 특정 위치에 미리 정의된 코드를 생성하는 함수다. 일반 함수 호출과 달리 호출 오버헤드가 없다. CUDA의 __syncthreads()는 하드웨어 동기화 명령으로 컴파일된다.

__syncthreads()는 barrier synchronization을 수행한다. 같은 블록 내 모든 스레드가 이 지점에 도달할 때까지 기다린다. 모든 스레드가 도착한 후에야 다음 명령을 실행한다.

plain
__global__ void kernel_func(void) {
    __shared__ float shared[SIZE];

    shared[i] = value;
    __syncthreads();  // 모두 쓰기 완료될 때까지 대기

    another = shared[j];  // 이제 안전하게 읽기 가능
}

위 코드의 실행 흐름을 보면 다음과 같다.

  1. 각 스레드가 shared[i] = value를 실행한다.
  2. __syncthreads() 지점에 도달한 스레드는 멈춘다.
  3. 빠른 스레드들은 여기서 대기한다.
  4. 느린 스레드들이 하나씩 도착한다.
  5. 마지막 스레드가 도착하면 barrier가 해제된다.
  6. 이제 모든 스레드가 동시에 다음 줄로 진행한다.
  7. another = shared[j]를 실행할 때는 shared 배열이 완전히 채워진 상태가 보장된다.

__syncthreads()는 두 가지 목적으로 사용된다. 첫째, 모든 원소가 업데이트되었음을 보장한다 (write 보장). 둘째, 모든 원소가 사용되었음을 보장한다 (read 보장).

하지만 __syncthreads()는 heavy operation이다. 1024개 스레드가 모두 도착할 때까지 기다려야 하므로 시간이 걸린다. 필요할 때만 사용해야 한다.

Adjacent Difference 예제

배열의 인접한 원소 사이의 차이를 계산하는 문제로 shared memory 사용법을 확인한다.

문제 정의

각 원소에 대해 인접한 원소 사이의 차이를 계산한다.

plain
vecB[i] = vecA[i] - vecA[i - 1]

예외 케이스로 i = 0일 때는 다음과 같이 처리한다.

plain
vecB[i] = vecA[i] - 0.0f

i번째 스레드는 vecA[i-1]과 vecA[i]를 읽어야 한다. 즉, 모든 스레드가 n개의 데이터에 접근한다.

Host 버전

c
unsigned num = 16 * 1024 * 1024;

int main(const int argc, const char* argv[]) {
    for (register unsigned i = 0; i < num; ++i) {
        if (i == 0) {
            vecB[i] = vecA[i] - 0.0f;
        } else {
            vecB[i] = vecA[i] - vecA[i - 1];
        }
    }
}

Host 버전은 단순하다. 모든 원소를 순회하며 이전 값과의 차이를 계산한다. 실행 시간은 31,810 usec다.

CUDA Global Memory 버전

plain
__global__ void kernelAdjDiff(float* b, const float* a, unsigned num) {
    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i == 0) {
        b[i] = a[i] - 0.0f;
    } else if (i < num) {
        b[i] = a[i] - a[i - 1];
    }
}

각 스레드가 global memory에서 a[i]와 a[i-1]을 직접 읽는다. Global memory 접근이 2번 발생한다. 하나는 a[i] 읽기고, 다른 하나는 a[i-1] 읽기다. 결과를 b[i]에 쓰는 것까지 합쳐서 총 3번의 global memory 접근이 일어난다.

실행 시간은 473 usec다. Host 버전보다 약 67배 빠르다.

CUDA Shared Memory 버전

Tiled approach를 사용하여 성능을 더 개선한다.

Tiled Approach란

큰 배열을 작은 조각(tile)으로 나누어 처리하는 방법이다.

예를 들어 16M개 원소를 가진 배열이 있다고 하자. 한 번에 처리하기에는 너무 크다. 그래서 1024개씩 묶어서 16,384개의 타일로 나눈다. 각 블록이 하나의 타일을 담당한다.

Global memory 관점에서는 하나의 긴 배열이다.

plain
[0][1][2][3]...[1023][1024][1025]...[2047][2048]...

하지만 논리적으로는 이렇게 나눈다.

plain
Tile 0:    [0]...[1023]      <- Block 0 담당
Tile 1:    [1024]...[2047]   <- Block 1 담당
Tile 2:    [2048]...[3071]   <- Block 2 담당
...

각 블록은 자기가 담당하는 타일만 shared memory로 복사한다. Shared memory는 작기 때문에 (48KB) 전체를 담을 수 없다. 하지만 1024개 원소 (4KB)는 충분히 담을 수 있다.

Global Index vs Local Index

여기서 두 가지 인덱스가 필요하다.

plain
global index (i) = blockIdx.x * blockDim.x + threadIdx.x
local index (tx) = threadIdx.x
  • *Global index (i)**는 전체 배열에서의 위치다. 예를 들어 Block 2의 Thread 5라면 i = 2 * 1024 + 5 = 2053이다.
  • *Local index (tx)**는 타일 내에서의 위치다. 같은 Thread 5라면 tx = 5다.

왜 둘 다 필요할까?

  • Global memory 접근: a[i] - 전체 배열에서 내 위치
  • Shared memory 접근: s_data[tx] - 타일 내에서 내 위치

Block 2의 Thread 5는:

  • a[2053]을 읽어서
  • s_data[5]에 저장한다

모든 블록이 같은 코드를 실행하지만, 서로 다른 타일을 처리한다.

plain
__global__ void kernelAdjDiff(float* b, const float* a, unsigned num) {
    __shared__ float s_data[1024];
    register unsigned tx = threadIdx.x;
    register unsigned i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < num) {
        s_data[tx] = a[i];
    }
    __syncthreads();

    if (tx > 0) {
        b[i] = s_data[tx] - s_data[tx - 1];
    } else if (i > 0) {
        b[i] = s_data[tx] - a[i - 1];
    } else {
        b[i] = s_data[tx] - 0.0f;
    }
}

코드 실행 흐름

위 커널의 동작을 구체적인 예로 살펴본다. Block 2를 가정한다.

1단계: Shared memory 선언

__shared__ float s_data[1024]로 블록당 1024개 float를 담을 공간을 만든다. 이 공간은 Block 2의 모든 스레드가 공유한다.

2단계: 인덱스 계산

Thread 5 입장에서 보면:

  • tx = 5 (타일 내 위치)
  • i = 2 * 1024 + 5 = 2053 (전체 배열 위치)

3단계: Global → Shared 복사

plain
if (i < num) {
    s_data[tx] = a[i];
}

Thread 5는 a[2053]을 읽어서 s_data[5]에 저장한다. 동시에:

  • Thread 0은 a[2048] → s_data[0]
  • Thread 1은 a[2049] → s_data[1]
  • ...
  • Thread 1023은 a[3071] → s_data[1023]

모든 스레드가 협력하여 타일을 채운다.

4단계: 동기화

__syncthreads()로 대기한다. 모든 스레드가 복사를 완료할 때까지 기다린다. 이제 s_data[0]부터 s_data[1023]까지 완전히 채워진 상태다.

5단계: 계산 수행

세 가지 경우로 나뉜다.

경우 1: 타일 중간 스레드 (tx > 0)

Thread 5는 tx = 5 > 0이므로:

plain
b[i] = s_data[tx] - s_data[tx - 1];

b[2053] = s_data[5] - s_data[4]를 계산한다. 두 값 모두 shared memory에서 읽는다. Global memory 접근이 없다.

경우 2: 타일 첫 번째 스레드, 첫 블록 아님 (tx == 0, i > 0)

Thread 0은 tx = 0이지만 i = 2048 > 0이므로:

plain
b[i] = s_data[tx] - a[i - 1];

b[2048] = s_data[0] - a[2047]을 계산한다. 이전 타일의 마지막 원소가 필요하므로 global memory를 읽는다.

경우 3: 전체 배열의 첫 원소 (i == 0)

Block 0의 Thread 0만 해당한다.

plain
b[i] = s_data[tx] - 0.0f;

b[0] = s_data[0] - 0.0f를 계산한다.

메모리 접근 패턴을 정리하면 다음과 같다. 일반적인 경우 global memory 읽기 1번 (a[i] → s_data[tx]), shared memory 읽기 2번 (s_data[tx], s_data[tx-1]), global memory 쓰기 1번 (b[i])이다.

블록의 첫 스레드 (tx == 0)인 경우 global memory 읽기 2번, shared memory 읽기 1번, global memory 쓰기 1번이다.

성능 비교와 분석

실행 시간을 비교하면 다음과 같다.

plain
Host version:         31,810 usec
CUDA global memory:      473 usec (67배 향상)
CUDA shared memory:      587 usec (오히려 더 느림!)

Shared memory 버전이 global memory 버전보다 느리다. 왜 그럴까?

데이터 재사용이 없기 때문이다.

이 문제에서 각 데이터는 딱 한 번만 읽힌다. Thread 5 입장에서 보면:

  • a[2053]을 읽는다 → 단 1번
  • a[2052]를 읽는다 → 단 1번 (이웃 스레드가 읽은 값)

Global memory 버전:

  1. a[i] 읽기 (1번)
  2. a[i-1] 읽기 (1번)
  3. b[i] 쓰기 (1번)

총 3번 global memory 접근.

Shared memory 버전:

  1. a[i] 읽기 (1번)
  2. s_data[tx]에 쓰기
  3. __syncthreads() 대기
  4. s_data[tx] 읽기 (shared)
  5. s_data[tx-1] 읽기 (shared)
  6. b[i] 쓰기 (1번)

총 2번 global memory 접근 + shared 복사 오버헤드 + 동기화 오버헤드.

Coalesced Access의 위력

더 중요한 점은 global memory 버전도 이미 최적화되어 있다는 것이다.

Block 2의 모든 스레드가 동시에 a[2048], a[2049], a[2050], ... a[3071]을 읽는다. 이는 연속된 메모리 주소다. GPU는 이런 패턴을 coalesced access로 처리한다. 하나의 메모리 트랜잭션으로 여러 데이터를 한꺼번에 가져온다.

결국 global memory에서 직접 읽는 것과 shared memory로 복사하는 것의 비용이 비슷하다. 오히려 복사 단계와 동기화 단계가 추가되어 느려진다.

Shared Memory가 유리한 경우

같은 데이터를 여러 번 읽을 때 shared memory가 빛을 발한다.

예를 들어 3x3 convolution을 생각해보자. 한 픽셀을 계산하려면 주변 9개 픽셀이 필요하다. 인접한 픽셀들은 겹치는 영역이 많다.

plain
Pixel A 계산: [0][1][2]
              [3][4][5]
              [6][7][8]

Pixel B 계산:   [1][2][3]
                [4][5][6]
                [7][8][9]

[4]번 픽셀은 A와 B 모두에서 사용된다. Global memory에서 매번 읽으면 2번 접근이다. 하지만 shared memory에 한 번 로드하면 여러 스레드가 재사용할 수 있다.

이런 경우 shared memory를 사용하면 10배 이상 빨라진다.

Shared Memory 사용 시기

Shared memory를 사용하면 항상 빨라지는 것은 아니다. 적절한 상황에서 사용해야 한다.

사용해야 할 때

1. 데이터 재사용이 많을 때

같은 데이터를 여러 스레드가 읽거나, 한 스레드가 여러 번 읽는 경우다.

  • 2D/3D Convolution - 겹치는 영역이 많음
  • Matrix Multiplication - 같은 행/열을 여러 번 사용
  • Stencil 연산 - 주변 픽셀을 반복 접근

이런 경우 global memory에서 한 번 로드한 후 shared memory에서 여러 번 읽는다.

2. Coalesced access가 불가능할 때

Global memory 접근 패턴이 불규칙하면 성능이 크게 떨어진다.

예를 들어 행렬의 전치(transpose)를 생각해보자. 행 단위로 읽어서 열 단위로 쓴다. 읽기는 연속이지만 쓰기는 불연속이다. 이때 shared memory로 데이터를 재배치하여 접근 패턴을 개선할 수 있다.

3. 블록 내 스레드 간 통신이 필요할 때

Reduction, prefix sum, scan 같은 알고리즘은 중간 결과를 공유해야 한다.

예를 들어 배열의 합을 구하는 reduction에서는 스레드들이 부분합을 교환한다. Global memory로는 느리고, shared memory가 필수다.

4. 계산 복잡도가 높을 때

메모리 복사 오버헤드보다 계산 시간 절약이 크면 유리하다. 간단한 연산(덧셈, 뺄셈)보다 복잡한 연산(행렬곱, FFT)에서 효과가 크다.

사용하지 말아야 할 때

1. 데이터를 한 번만 읽을 때

Adjacent difference 예제처럼 각 데이터를 딱 한 번만 읽는다면 global memory에서 직접 읽는 것이 낫다.

2. Coalesced access가 이미 가능할 때

연속된 메모리를 순차적으로 읽는다면 이미 최적화되어 있다. Shared memory로 복사하는 비용이 오히려 손해다.

3. Shared memory 용량이 부족할 때

블록당 48KB밖에 없다. 큰 데이터를 처리한다면 타일로 나누어야 하는데, 타일이 너무 작으면 경계 처리 오버헤드가 커진다.

관련 문서