CUDA - 메모리의 계층 구조

그림 - CUDA 메모리 계층 구조

Memory	Location on/off Chip	Cached	Access	Scope	Life Time
Register	On	N / A	R/W	1 thread	Thread
Local	Off	+	R/W	1 thread	Thread
Shared	On	N / A	R/W	All Threads in Block	Block
Global	Off	+	R/W	All Threads + Host	Host Allocation
Constant	Off	Yes	R	All Threads + Host	Host Allocation
Texture	Off	Yes	R	All Threads + Host	Host Allocation

그림 - CUDA 메모리 특성

액세스 속도 (느림 -> 빠름)

Global -> Local -> Shared -> Register

메모리 용량 (작음 -> 큼)

Register -> Shared -> Local -> Global

액세스 범위 (좁음 -> 넓음)

Register -> Local -> Shared -> Global

Register Memory는 온 칩 프로세서에 있는 메모리로 직접 연산을 수행하는 가장 빠른 메모리이다. CUDA C를 이용해 구현할 때도 몇 개의 레지스터를 사용하는 지와 레지스터가 어떻게 재사용 되는지 알 수 없다.

다만 커널 함수 내에서 로컬 변수를 많이 사용할수록 레지스터 사용량이 증가하고, GPU마다 레지스터의 수는 차이가 있다.

sin(), cos() 등 수학 함수를 사용할 때도 레지스터를 사용한다.

로컬 메모리(Local Memory)는 커널 함수 내에서 너무 많은 로컬 변수를 사용하거나, 배열형 변수로 큰 용량을 사용하면 프로세서 밖에 있는 DRAM에 메모리가 할당된다. 로컬 변수가 레지스터로 사용될지 로컬 메모리에 할당 될지 명확하지 않다.

Local Memory로 할당되는 경우는 다음과 같다.

1. 너무 많은 레지스터 변수를 사용했을 때

2. 너무 많은 로컬 변수를 사용했을 때

3. 로컬 변수로 배열을 사용했을 때

쉽게 말하면, 각 Thread는 자신만이 사용하기 위한 메모리이다.

공유 메모리(Shared Memory)는 CUDA의 큰 장점 중 하나이다. CUDA C 컴파일러는 공유 메모리의 변수들을 일반 변수들과는 다르게 취급한다. CUDA C 컴파일러는 GPU에서 실행되는 각 블록마다 변수의 복사본을 하나씩 생성한다. 블록 내의 각 스레드는 메모리를 서로 공유하지만, 다른 블록 내에서 보고 있는 변수의 복사본을 보거나 수정할 수는 없다. 하지만 공유 메모리는 블록 내의 스레드들이 계산하는 데 통신과 협력을 할 수 있게 함으로써 훌륭한 수단을 제공한다. 더욱이 공유메모리 버퍼는 오프칩(off-chip)인 DRAM에 상주하는 것과 달리 물리적으로 GPU 상에 상주한다. 이 때문에 공유 메모리의 접근 지연 시간은 블록당 공유 메모리를 효율적으로 생성하는 일반 버퍼, 소프트웨어로 관리되는 캐시 또는 스크래치패드(Scratchpad)보다 훨씬 짧은 경향이 있다.

GPU의 SM 안에 공유 메모리가 있고, SM은 블록 단위로 프로그램을 처리 하기 때문에 공유 메모리의 데이터는 동일한 블록 안에 있는 스레드 사이에서만 공유할 수 있다. 공유 메모리를 할당하는 방법은 정작할당, 동적할당이 있다.

쉽게 말하면, 각 Thread 간에 data 공유를 위한 메모리이다.

__shared__ int a; 

전역 메모리(Global Memory)는 비디오 카드에 장착된 DRAM 메모리를 의미한다. 전역 메모리는 GPU 칩 외부에 있기 때문에 액세스 속도는 레지스터나 공유 메모리에 비해 많이 느리지만, CPU의 메모리 보다는 빠르다.

쉽게 말하면, Block들의 집합인 Grid간에 data 공유를 위한 메모리이다.

상수 메모리(Constant Memory)는 DRAM에 있는 데이터를 읽기 전용으로 사용하며 캐시를 지원한다. NVIDIA의 그래픽 카드는 Constant Memory의 크기가 64KB이다. 최초 데이터는 DRAM에서 가져오기 때문에 속도가 좀 느리지만, 한번 캐시에 올라온 값을 반복하여 재 사용하기 때문에 속도가 빠르다. Host에서는 값을 쓰고, Device 영역에서는 값을 읽기만 할 수 있다.

쉽게 말하면, Half - Warp내의 모든 스레드가 동일한 주소의 데이터를 상수 메모리로 부터 요청하면 GPU는 단 한번의 읽기만 요청한 후 모든 스레드에게 그 데이터를 알려주는 메모리이다.

__constant__ int a;

텍스처 메모리(Texture Memory)는 상수 메모리처럼 칩에 캐싱된다. 따라서 어떤 경우엔 텍스처 메모리는 오프칩인 DRAM에 대한 메모리 요청을 줄임으로써 매우 효과적인 대역폭을 제공할 수 있다. 텍스처 캐시는 공간 구역성(spatial locality)를 자주 드러내는 메모리 접근 패턴을 가진 그래픽스 어플리케이션들을 위해 설계되었다. 

texture<자료형> a;

출처 : http://blog.naver.com/riverrun17/220420579990

CUDA 메모리에 대해 공부하다가 어떤 분이 잘 정리해 놓으셔서 긁어왔습니다....