Shared Memory (1)

References

• Professional CUDA C Programming

Contents

• Shared Memory (SMEM)
• Shared Memory Banks and Access Mode
• Configuring the Amount of Shared Memory
• Synchronization
• Volatile Qualifier

---

이번 포스팅에서는 공유 메모리에 대해서 다시 한 번 알아보는 시간을 갖도록 하려고 합니다 !

 

Intro

GPU 디바이스에는 다음과 같이 두 종류의 메모리 타입이 있습니다.

• On-board memory
• On-chip memory

Global Memory는 크고, on-board 메모리이며 비교적 높은 latency를 갖고 있습니다. Shared Memory는 Global Memory보다는 작고, 더 작은 latency를 갖는 on-chip 메모리입니다. 따라서 Global Memory보다 더 높은 bandwidth를 제공합니다.

Shared Memory를 program-managed cache라고 보면 되고, 일반적으로 다음과 같은 용도로 유용하게 사용할 수 있습니다.

• An intra-block thread communication channel (블록 내 스레드 간의 통신 채널)
• A program-managed cache for global memory data
• Scratch pad memory for transforming data to improve global memory access patterns

---

Shared Memory

Shared Memory (SMEM)은 GPU에서 key 컴포넌트 중의 하나입니다. 물리적으로, 각 SM은 하나의 작고 low-latency 메모리 pool을 포함하는데, 이는 그 SM에서 현재 실행 중인 스레드 블록 내의 모든 스레드가 공유합니다. Shared Memory는 동일한 스레드 블록 내에서 스레드들이 협력할 수 있도록 해주고, on-chip 데이터의 재사용성을 높혀주며 global memory bandwidth를 크게 줄일 수 있도록 해줍니다. Shared Memory의 데이터는 어플리케이션에서 명시적으로 관리되기 때문에 종종 program-managed cache라고도 합니다.

GPU의 메모리 계층은 위와 같습니다. (read-only data를 위한 cache는 Kepler 아키텍처에만 있습니다.)

그림에서 보여주듯이, Shared Memory와 L1 cache는 물리적으로 L2 cache와 Global Memory보다 SM에 더 가깝습니다. 그 결과, Shared Memory의 latency는 Global Memory보다 20~30배 낮으며, bandwidth도 10배 가량 높습니다.

 

스레드 블록들이 실행될 때, 고정된 크기의 shared memory가 각 스레드 블록에 할당됩니다. 이 shared memory 주소 공간은 동일한 스레드 블록의 모든 스레드들에게 공유됩니다. 따라서, shared memory의 데이터들은 스레드 블록에서 동일한 lifetime을 가집니다. 

Share memory의 액세스는 warp 단위로 실행됩니다. warp에서 shared memory에 액세스하기 위한 각 요청은 한 번의 transaction으로 처리되는 것이 이상적입니다. 최악의 경우는 shared memory에 대한 요청이 각각의 32개의 transaction에서 순차적으로 수행되는 것입니다.

 

Shared Memory는 SM에 현재 존재하는 스레드 블록들 사이에서 분할되므로, Shared Memory는 device 병렬화를 제한하는 중요한 리소스입니다. 만약 커널의 더 많은 shared memory를 사용한다면, 동시에 활성화 되는 스레드 블록은 더 적어집니다.

 

Shared Memory Allocation

Shared Memory 변수를 할당하는 방법은 여러 가지가 있습니다. 정적으로나 동적으로 할당할 수 있고, CUDA 커널의 local이나 소스 코드에서 global로 선언될 수도 있습니다.

Share Memory 변수는 선언할 때 다음의 qualifier를 붙여줍니다.

__shared__

아래 코드는 shared memory에 2D float 배열을 선언하는 코드입니다. 만약 커널 함수 내에서 선언한다면, 이 변수의 scope는 커널에 대해 local입니다. 만약 파일에서 커널 외부에서 선언한다면, 이 변수의 scope는 모든 커널에 global 합니다.

__shared__ float tile[size_y][size_x];

만약 컴파일 타임에 shared memory의 사이즈를 알 수 없다면, extern 키워드를 붙여서 size가 없는 배열로 선언할 수 있습니다. 예를 들어, 다음의 코드는 1차원의 사이즈가 결정되지 않은 int 배열을 선언합니다. 이 경우에도 커널 내부나 외부에서 모두 가능합니다.

extern __shared__ int tile[];

이 배열의 사이즈는 컴파일 시간에 알 수 없기 때문에, 각 커널을 시작할 때 원하는 크기를 byte단위로 execution configuration의 3번째 파라미터에 지정해주어서 동적으로 할당할 수 있습니다.

kernel<<<grid, block, size * sizeof(int)>>>(...)

동적으로 할당하는 경우에는 오직 1차원의 배열만 선언할 수 있습니다.

 

---

Shared Memory Banks and Access Mode

Latency와 Bandwidth는 메모리 성능을 최적화할 때 측정해야하는 두 가지 속성입니다. Global Memory를 사용할 때에는 메모리 액세스 패턴에 따라서 성능의 차이를 보입니다. Shared Memory는 Global Memory의 latency와 Bandwidth 성능 영향을 숨기는데 사용될 수 있습니다. 이를 이해하기 위해서 shared memory가 어떻게 정렬되는지 알아보도록 하겠습니다.

 

Memory Banks

높은 메모리 bandwidth를 달성하기 위해서 shared memory는 banks(뱅크)라는 32개의 동일한 사이즈의 메모리 모듈로 나뉘며 이는 동시에 액세스할 수 있습니다. 32개의 bank인 이유는 하나의 warp에 32개의 스레드가 있기 때문입니다. Shared Memory는 1차원 주소 공간입니다. GPU의 compute capability에 따라서 shared memory의 주소는 다른 패턴으로 다른 bank에 매핑됩니다. 만약 warp에 의해 발생한 shared memory load 또는 store 명령이 bank 당 둘 이상의 메모리에 액세스하지 않는다면, 그 명령은 한 번의 memory transaction으로 처리될 수 있습니다. bank 당 둘 이상의 메모리에 액세스한다면, 그 명령은 여러 memory transaction에 의해서 처리되고, 결국 메모리 bandwidth 활용을 감소시킵니다.

 

Bank Conflict

shared memory의 여러 주소들에 대한 요청이 동일한 메모리 bank에 발생하면, bank conflict가 발생하고 이 요청은 반복됩니다. 하드웨어는 bank conflict가 발생한 요청을 conflict가 발생하지 않는 여러 transaction으로 필요한만큼 분할합니다. 따라서, 분할된 memory transaction의 비율만큼 effiective bandwidth를 감소시킵니다.

 

Warp에 의해서 shared memory에 대한 요청이 있을 때, 일반적으로 아래와 같은 3가지의 상황이 발생합니다.

• Parallel access: multiple addresses accessed across multiple banks
• Serial access: multiple addresses accessed within the same bank
• Broadcast access: a sing address read in a single bank

Parallel access는 가장 일반적인 패턴입니다. 이 패턴은 전부는 아니지만 addresses의 일부가 하나의 memory transaction으로 처리될 수 있다는 것을 의미합니다. 모든 address가 별도의 bank에 있을 때, conflict-free인 shared memory 액세스가 수행됩니다.

 

Serial access는 worst 패턴입니다. 여러 addresses가 동일한 bank에 속할 때, 그 메모리 요청은 serial로 처리됩니다. 만약 warp 내의 모든 스레드가 단일 bank 내에 존재하는 다른 메모리 위치에 액세스할 때, 32번의 memory transaction이 필요하고, 병렬로 한 번의 요청으로 처리될 때보다 32배나 오래 걸리게 됩니다.

 

Broadcast access는 한 warp의 모든 스레드가 하나의 bank 내의 동일한 주소를 읽습니다. 한 번의 memory transaction이 수행되고, 액세스된 word는 모든 스레드로 broadcast 됩니다. broadcast access에서 오직 한 번의 memory transaction만이 필요하지만, 아주 적은 양의 bytes만 읽으므로 bandwidth 활용도는 낮습니다.

 

아래의 그림은 최적의 parallel access 패턴을 보여줍니다.

Optimal parallel access pattern

각 스레드는 하나의 32-bit word에 액세스합니다. 각 스레드는 다른 bank의 주소에 액세스하기 때문에 bank conflict가 발생하지 않습니다.

 

아래 그림은 불규칙적이고 랜덤한 액세스 패턴을 보여줍니다.

Irregular, random access pattern

여기에도 각 스레드가 다른 bank의 주소에 액세스하므로, bank conflict는 발생하지 않습니다.

 

다음 그림은 몇몇의 스레드가 동일한 bank의 주소에 액세스하는 또 다른 불규칙 액세스 패턴을 보여줍니다.

Irregular access pattern where several threads access the same bank

이러한 메모리 요청에는 아래의 두 가지 종류가 있습니다.

• Conflic-free broadcast access if threads access the same address within a bank
• Bank conflict access if threads access different addresses within a bank

 

Access Mode

Shared Memory Bank width는 shared memory bank가 있는 shared memory addresses를 정의합니다. Memory bank width는 compute capability에 따라 다릅니다.

• 4 bytes (32-bits) for devices of compute capability 2.x
• 8 bytes (64-bits) for devices of compute capability 3.x 이상
  (Maxwell 부터는 다시 Fermi의 style로 다시 복구)

Fermi에서 bank width는 32비트이며 32개의 bank가 있습니다. 각 bank는 2 클럭 사이클 당 32비트의 bandwidth를 갖습니다. 연속적인 32비트 word는 연속적인 bank를 매핑합니다. 따라서, shared memory address에서 bank index의 매핑은 다음의 식으로 계산될 수 있습니다.

bank index = (byte address / 4 bytes/bank) % 32 banks

아래의 위쪽 그림은 byte address와 word index의 매핑을 보여주고, 아래 그림은 word index와 bank index의 매핑을 보여줍니다.

warp에서 가능한 동시 액세스를 최대화하기 위해서 인접한 word는 다른 bank에 분류됩니다.

 

Kepler 디바이스부터 shared memory는 다음의 2가지 address mode의 32 bank를 갖습니다.

• 64-bit mode
• 32-bit mode

64-bit 모드에서 연속적인 64-bit words가 연속적인 bank에 매핑됩니다. 각 bank는 클럭 사이클당 64 bits의 bandwidth를 갖습니다. shared memory address에 매핑되는 bank index는 다음과 같이 계산할 수 있습니다.

bank index = (byte address / 8 bytes/bank) % 32 banks

 

Memory Padding

Memory Padding은 bank conflict를 피하는 한 가지 방법입니다. 아래 그림은 간단한 예시를 보여줍니다.

5개의 shared memory bank를 갖는 디바이스라고 가정해보겠습니다. 만약 모든 스레드가 bank 0의 다른 위치에 액세스한다면, 5-way의 bank conflict가 발생합니다. 이러한 종류의 bank conflict를 피하는 방법은 마지막 원소 뒤에 하나의 padding word를 추가하는 것입니다. 이렇게 padding을 추가하면, word들이 bank에 매핑되는 것이 위 그림의 오른쪽처럼 변경됩니다. 따라서, padding 때문에 기존에 bank 0에 포함되었던 word들이 다른 bank에 속하게 됩니다.

 

Access Mode Configuration

Kepler부터 4-byte와 8-byte의 shared memory access mode를 지원합니다. default mode는 4-byte mode입니다. 이 access mode는 아래의 CUDA 런타임 API 함수로부터 쿼리할 수 있습니다.

cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig* pConfig);

결과가 pConfig에 반환되며, 이 값은 아래의 값들이 될 수 있습니다.

• cudaSharedMemBankSizeFourByte
• cudaSharedMemBankSizeEightByte

그리고, 다음의 API 함수를 통해서 디바이스의 bank size를 설정할 수 있습니다.

cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);

가능한 파라미터 값은 다음과 같습니다.

• cudaSharedMemBankSizeDefault
• cudaSharedMemBankSizeFourByte
• cudaSharedMemBankSizeEightByte

커널 수행 간의 shared memory configuration을 변경하는 것은 암시적으로 device 동기화를 요청할 수 있습니다. shared memory의 bank size를 변경하는 것은 shared memory의 사용량이 증가하거나 커널의 occupancy에 영향을 미치지는 않지만 성능에는 큰 영향을 미칠 수 있습니다. 큰 bank size는 shared memory 액세스에 대한 bandwidth가 높아질 수 있지만, 어플리케이션의 shared memory access 패턴에 따라 더 많은 bank conflict가 발생할 수 있습니다.

---

Configuring the Amount of Shared Memory

각 SM에는 64KB의 on-chip 메모리가 있습니다. Shared Memory와 L1 cache는 이 하드웨어 리소스를 공유합니다. CUDA는 L1 cache와 shared memory의 크기를 설정하기 위한 두 가지 방법을 제공합니다.

• Per-device configuration
• Per-kernel configuration

다암의 런타임 함수를 통해 주어진 디바이스에서 실행되는 커널이 사용하는 L1 cache와 shared memory 크기를 설정할 수 있습니다.

cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);

cacheConfig 파라미터는 현재 CUDA device에서 L1 cache와 shared memory간의 분할되는 방법을 지정합니다.

• cudaFuncCachePreferNon: no preference(default)
• cudaFuncCachePreferShared: prefer 48KB shared memory and 16KB L1 cache
• cudaFuncCachePreferL1: prefer 48KB L1 cache and 16KB shared memory
• cudaFuncCachePreferEqual: prefer 32KB L1 cache and 32KB shared memory

어떤 mode가 좋은지는 커널에서 얼만큼의 shared memory를 사용하는지에 따라 다릅니다.

일반적으로는 다음과 같습니다.

• 커널에서 더 많은 shared memory를 사용할 때 shared memory를 선호
• 커널에서 register를 더 많이 사용할 때, L1 cache을 선호

 

CUDA 런타임은 요청된 on-chip 메모리 configuration을 사용하지만, 커널 함수를 실행할 때 다른 configuration을 선택하도록 할 수 있습니다. per-kernel configuration은 다음의 런타임 함수를 통해 설정할 수 있습니다.

cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCahceca cheConfig);

func 포인터에 커널 함수를 지정하면 되고, 각 커널당 이 함수는 한 번만 호출하면 됩니다.

 

L1 cache와 shared memory는 동일한 on-chip 하드웨어에 위치하지만, 몇 가지 차이점이 있습니다. Shared Memory는 32 banks를 통해 액세스되는 반면, L1 cache는 cache line을 통해 액세스됩니다. 또한 Shared Memory를 사용하면 저장되는 항목과 위치를 완전히 제어할 수 있지만, L1 cache의 경우에 데이터 제거는 하드웨어에 의해서 수행됩니다.

 

---

Synchronization

메모리 이름에서 암시하듯이, shared memory는 한 스레드 블록 내의 여러 스레드들에 의해서 동시에 액세스될 수 있습니다. 이는 동일한 shared memory 위치에 동기화없이 여러 스레드 블록에 의해서 수정될 때 inter-thread conflict를 일으킵니다. CUDA는 intra-block 동기화를 수행하기 위한 몇 가지 런타임 함수들을 제공합니다.

일반적으로 동기화에는 두 가지 접근이 있습니다.

• Barriers
• Memory fences

Barrier에서 동시에 수행되는 모든 스레드들은 다른 스레드가 이 barrier point에 도달할 때까지 대기합니다. Memory Fence에서는 모든 스레드가 메모리에 대한 수정이 다른 모든 스레드에 visible할 때까지 중지합니다.

 

CUDA의 intra-block barriers와 memory fences에 대해 살펴보기 전에, 먼저 CUDA에서 채택된 weakly-ordered memory model을 살펴보겠습니다.

 

Weakly-Ordered Memory Model

현대 메모리 아키텍처는 relaxed memory model입니다. 이는 memory 액세스가 프로그램 내에서 표현하는 순서로 수행될 필요가 없다는 것을 의미합니다. CUDA는 weakly-ordered memory model을 채택하여 더 공격적인 컴파일 최적화를 가능하게 합니다.

 

GPU 스레드가 데이터를 다른 메모리들(shared memory, global memory, page-locked host memory or the memory of a peer device) 쓰는 순서는 소스 코드에서의 액세스 순서와 동일하지 않습니다. 스레드의 wrtie가 다른 스레드에 보이는 순서는  실제 그 write가 수행되는 순서와 일치하지 않습니다.

스레드가 다른 메모리로부터 데이터를 읽는 순서도 명령어가 서로 독립적이라면 프로그램에서 나타나는 명령어의 순서와 같지 않습니다.

따라서, 프로그램을 정확히 특정 순서로 강제하려면, memory fence와 barrier를 반드시 코드에 삽입해주어야 합니다. 이것이 리소스를 다른 스레드와 공유하는 커널의 올바른 동작을 보장할 수 있는 유일한 방법입니다.

 

Explicit Barrier

CUDA에서 동일한 스레드 블록 내의 스레드들 간에서만 barrier를 수행할 수 있습니다. 커널에서 다음의 intrinsic function을 호출함으로써 barrier point를 지정할 수 있습니다.

void __syncthreads();

__syncthreads는 블록의 스레드들이 모든 스레드가 해당 포인트에 도달할 때까지 반드시 기다리도록 합니다. 또한, __syncthreads 이전에 그 스레드들에 의한 모든 global과 shared memory 액세스가 동일 블록의 모든 스레드들에게 visible하도록 해줍니다.

 

__syncthreads는 동일 블록의 스레드 간의 커뮤니케이션을 조정하는데 사용됩니다. 블록 내의 어떤 스레드들이 shared or global memory의 동일한 주소에 액세스할 때, 잠재적인 위험(read-after-write, write-after-read, write-after-write)가 있으며 이는 그 메모리 공간에 정의되지 않은 동작이나 상태를 유발합니다. 이러한 의도하지 않은 동작은 conflicting 액세스 사이의 스레드들을 동기화함으로써 피할 수 있습니다.

 

조건문에서 __syncthreads를 사용할 때에는 특히 주의해야합니다. 전체 스레드 블록에서 동일한 조건을 보장할 때에만 __syncthreads를 호출할 수 있습니다. 그렇지 않다면 실행이 중단되고 의도하지 않은 side effect가 발생할 수 있습니다.

예를 들어, 아래의 코드는 블록 내의 모든 스레드가 동일한 barrier point에 도달하지 않으므로 블록 내의 스레드가 다른 스레드를 무한정 기다릴 수 있습니다.

if (threadID % 2 == 0) {
	__syncthreads();
}
else {
	__syncthreads();
}

 

블록 간의 동기화는 허용되지 않기 때문에 스레드 블록은 어떠한 순서로, 병렬이든 순차적이든 어떤 SM에서도 실행될 수 있습니다. 이러한 블록 실행의 독립성은 CUDA 프로그래밍을 더 scalable하게 만들어 주어서, 임의의 core 갯수에서도 실행될 수 있도록 해줍니다.

 

Memory Fence

Memory Fence 함수는 fence 전에 어떠한 memory write가 fence 후의 다른 스레드들에게 visible하도록 보장합니다. 여기에는 scope(block, grid, system)에 따라서 3가지 종류의 memory fence가 있습니다.

 

스레드 블록 내에서 memory fence는 다음의 intrinsic 함수를 사용하여 생성할 수 있습니다.

void __threadfence_block();

__threadfence_block은 호출된 스레드에 의해 fence 이전에 발생한 shared memory와 global memory의 모든 쓰기 작업이 fence 후의 동일한 블록의 다른 스레드들에게 visible 하도록 합니다. memory fences는 어떠한 스레드 동기화를 수행하지 않으므로, 블록의 모든 스레드가 실제로 이 명령을 실행할 필요는 없습니다.

 

grid level의 memory fence는 다음의 intrinsic 함수를 사용하여 생성할 수 있습니다.

void __threadfence();

__threadfence는 global memory에 대한 모든 쓰기 작업이 동일한 grid의 모든 스레드들에게 visible할 때까지 호출된 스레드를 중지시킵니다.

 

system 간(host and device)의 memory fence는 다음의 instrinsic 함수를 사용하여 생성할 수 있습니다.

void __threadfence_system();

__threadfence_system은 global memory, page-locked host memory, 다른 디바이스의 memory에 대한 모든 쓰기 작업이 모든 device와 host threads에게 visible하도록 하기 위해서 호출한 스레드를 중지시킵니다.

 

Volatile Qualifier

global이나 shared memory에 volatile qualifier를 사용하여 변수를 선언하는 것은 registers나 local memory에 data를 캐시하는 것과 같은 컴파일러 최적화를 막습니다. volatile qualifier를 사용하면, 컴파일러는 그 변수의 값이 변하거나 다른 스레드에 의해서 사용된다고 가정합니다. 그러므로, 이 변수에 대한 어떠한 참조는 캐시를 스킵하는 global memory read나 global memory write 명령어로 컴파일 됩니다.