GPU Memory Architecture 문서 원본 보기

[[분류: GPU]]

== 개요 ==
GPU는 높은 memory bandwidth와 대규모 parallelism을 제공하기 위해 CPU와는 다른 형태의 memory hierarchy를 가진다. NVIDIA GPU의 memory는 크게 SM 내부의 on-chip memory와 GPU DRAM에 위치하는 off-chip memory로 나눌 수 있다.

일반적으로 Register와 Shared Memory는 SM 내부에 존재하며 매우 낮은 latency를 제공한다. 반면 Global Memory와 Local Memory는 GPU의 device memory(HBM 또는 GDDR 등)에 위치하며 훨씬 큰 용량을 제공하지만 상대적으로 높은 access latency를 가진다.

== Memory ==
=== Register Memory ===
<syntaxhighlight lang="cpp">
__global__ void kernel() {
    int value = 10; // 일반적으로 register에 저장
}
</syntaxhighlight>
* Register memory는 각 SM에 존재하는 on-chip register file이다. 
* 각 thread는 자신만의 register state를 가지며 다른 thread가 직접 접근할 수 없다.
* Register는 GPU에서 가장 빠른 memory resource 중 하나이지만 크기가 제한되어 있다. 
* 하나의 kernel이 thread당 너무 많은 register를 사용하면 한 SM에서 동시에 실행할 수 있는 warp 또는 thread block의 수가 감소하여 occupancy가 낮아질 수 있다.

필요한 register의 수가 hardware가 제공할 수 있는 양보다 많아지면 compiler는 일부 값을 Local Memory로 spill한다. 이를 '''register spilling'''이라고 한다.

=== Shared Memory ===
<syntaxhighlight lang="cpp">
__global__ void kernel() {
    __shared__ int buffer[256];

    buffer[threadIdx.x] = threadIdx.x;
    __syncthreads();
}
</syntaxhighlight>
* Shared Memory는 SM 내부에 위치하는 빠른 on-chip memory이다.
* 이름 때문에 한 SM에서 실행되는 모든 thread가 공유하는 memory로 오해하기 쉽지만, CUDA programming model에서 Shared Memory의 기본 sharing scope는 '''thread block'''이다. 같은 thread block에 속한 thread들은 Shared Memory를 이용해 데이터를 공유할 수 있다.
* Shared Memory는 programmer가 직접 관리하는 scratchpad memory와 유사하다. Global Memory보다 latency가 작고 bandwidth가 높기 때문에 matrix multiplication, reduction, tiling 등의 연산에서 자주 사용된다.
* Modern NVIDIA GPU에서는 Shared Memory와 L1 cache가 동일한 '''unified data cache'''의 물리적인 resource를 공유한다. 따라서 architecture와 kernel configuration에 따라 L1과 Shared Memory 사이의 용량 배분이 달라질 수 있다. 보통 수십~수백KB의 크기를 per-SM마다 제공한다.
* Shared Memory는 여러 개의 '''memory bank'''로 구성된다. 일반적인 NVIDIA GPU는 32개의 bank를 사용한다. 같은 warp의 여러 thread가 같은 bank의 서로 다른 address에 동시에 접근하면 '''bank conflict'''가 발생하여 access가 serialize될 수 있다.

=== Local Memory ===
<syntaxhighlight lang="cpp">
__global__ void kernel() {
    int large_array[1024];
}
</syntaxhighlight>
* Local Memory는 각 thread가 독립적으로 사용하는 private memory이다.
* 그러나 "Local"이라는 이름과 달리 물리적으로 SM 내부에 존재하는 memory가 아니다. Local Memory는 일반적으로 Global Memory와 동일한 device DRAM에 위치하며 L1/L2 cache를 통해 접근한다. 
* 보통은 Register에 저장하지 못한 내부 Stack memory들이 Local memory로 빠진다.
* Local Memory는 각 thread마다 logical address space가 분리되어 있다. 한 thread의 Local Memory에 다른 thread가 일반적인 CUDA instruction을 이용하여 직접 접근할 수 없다.
* Local Memory의 lifetime은 function (e.g., kernel)이 종료되기 전까지이다.
* 일반적인 CUDA API를 통해 programmer가 Local Memory의 virtual-to-physical mapping을 직접 제어할 수는 없다.

=== Global Memory ===
<syntaxhighlight lang="cpp">
int *ptr;
cudaMalloc(&ptr, sizeof(int) * 1024);
</syntaxhighlight>
* Global Memory는 GPU에 연결된 off-chip device DRAM(HBM 또는 GDDR)에 존재하는 주요 memory 영역이다.
* Global Memory는 GPU의 모든 SM에서 접근할 수 있으며, 일반적으로 CUDA API를 통해 allocation한다.
* Global Memory는 Shared Memory나 Register보다 latency가 훨씬 크기 때문에 cache와 memory coalescing이 performance에서 매우 중요하다.

Warp의 thread들이 인접한 memory address를 접근하면 memory request들을 적은 수의 memory transaction으로 합칠 수 있다. 이를 '''memory coalescing'''이라고 한다.

=== Host Memory ===
<syntaxhighlight lang="cpp">
int *ptr = (int *)malloc(size);
cudaMemcpy(device_ptr, host_ptr, size,
           cudaMemcpyHostToDevice);
</syntaxhighlight>
* Host Memory는 CPU에 연결된 system DRAM이다.
* 기본적으로 <code>malloc()</code>이나 <code>new</code>로 할당된다.
* GPU와 CPU 사이에서 데이터를 전송하기 위해 일반적으로 <code>cudaMemcpy()</code>를 사용한다.
* CUDA는 host memory pinning도 지원하며, Pinned Memory는 OS가 page를 다른 physical location으로 이동시키지 않도록 고정하기 때문에 GPU와 DMA transfer를 수행할 때 일반 pageable memory보다 효율적인 경우가 많다.
* 일부 pinned memory는 GPU address space에 mapping하여 GPU가 system memory를 직접 접근하는 '''Zero-Copy''' 형태로 사용할 수도 있다.

=== Unified / Managed Memory ===
<syntaxhighlight lang="cpp">
int *ptr;
cudaMallocManaged(&ptr, size);
</syntaxhighlight>
* CUDA Unified Memory는 CPU와 GPU가 같은 allocation을 접근할 수 있도록 제공하는 memory abstraction이다.
* Discrete GPU에서는 CPU와 GPU가 서로 다른 physical memory를 가지므로, Unified Memory runtime이 page migration과 page-table update를 통해 필요한 processor 쪽으로 memory page를 이동시킬 수 있다. 반대로 CPU가 GPU에 resident한 page를 접근하면 GPU에서 CPU로 migration이 발생할 수도 있다.

최근 Linux에서는 Heterogeneous Memory Management(HMM)을 통해 일반 <code>malloc()</code>으로 할당된 system memory까지 GPU에서 접근할 수 있는 환경이 존재한다.

또한 Grace Hopper와 같이 CPU와 GPU가 hardware-coherent memory system으로 연결된 시스템에서는 CPU와 GPU가 logically combined page table과 cache coherence를 활용할 수 있다. 이러한 환경에서는 단순한 page migration뿐 아니라 remote coherent access가 가능하다.

==== Unified Virtual Addressing ====
CUDA의 '''Unified Virtual Addressing(UVA)'''은 하나의 process에서 CPU와 GPU memory가 하나의 unified virtual address space에 배치되는 기능이다.

따라서 pointer 값 자체를 이용하여 해당 pointer가 CPU memory인지 특정 GPU의 memory인지 구분할 수 있으며, CUDA runtime은 이를 이용하여 <code>cudaMemcpyDefault</code>와 같은 기능을 제공한다.

== Cache ==
Modern NVIDIA GPU의 일반적인 data cache hierarchy는 크게 '''L1 cache'''와 '''L2 cache'''로 나눌 수 있다.

=== L1 Cache ===
각 SM은 자신만의 L1 cache를 가진다.

Modern architecture에서는 L1 cache와 Shared Memory가 동일한 unified data cache의 physical resource를 공유한다.

따라서 architecture에 따라 다음과 같은 개념적인 구조가 사용된다.

<pre>
SM
 ├── Register File
 ├── L1 Cache
 └── Shared Memory
      ↑
      └─ 일부 architecture에서는 같은 unified data cache resource 공유
</pre>

=== L2 Cache ===
L2 cache는 GPU 전체의 SM들이 공유한다.

<pre>
        SM0 ─ L1
          \
        SM1 ─ L1
           \
            ─── L2 Cache ─── GPU DRAM
           /
        SM2 ─ L1
</pre>

Global Memory와 Local Memory access는 L2를 통해 caching될 수 있다.

== GPU MMU ==
[[파일:GPU MMU.png|500px|섬네일|가운데]]

Modern GPU 역시 CPU와 마찬가지로 Virtual Memory를 사용한다.
GPU instruction에서 사용되는 address는 일반적으로 virtual address이며, GPU의 [[MMU]]가 이를 physical address로 translation한다.

=== Page Table ===
Page Table은 Virtual Page Number(VPN)을 Physical Page Number(PPN)에 mapping한다.
CUDA는 GPU에서 여러 physical page size를 지원한다. 정확한 page size와 mapping granularity는 GPU architecture와 platform에 따라 달라질 수 있으며 NVIDIA는 현재 GPU가 여러 page size를 지원하고 2MiB 이상의 큰 physical page를 선호한다고 설명한다.

=== CPU/GPU Separate Page Table ===
일반적인 discrete GPU 기반의 software-coherent system에서는 CPU와 GPU가 서로 다른 logical page table을 가질 수 있다.

Unified Memory subsystem은 두 processor의 page table과 memory residency를 관리하며 필요한 경우 page를 migration한다.

Linux HMM 역시 CPU와 GPU가 서로 다른 page-table structure를 사용하는 환경에서 system memory를 GPU address space와 연결하고 software coherence를 제공할 수 있다.

=== CPU/GPU Combined Page Table ===
Grace Hopper와 같은 hardware-coherent system에서는 CPU와 GPU가 logically combined page table을 사용할 수 있다.

<pre>
             Virtual Address
                   |
           Combined Page Table
             /           \
            /             \
        CPU MMU           GPU MMU
          |                 |
       CPU Memory <----> GPU Memory
             coherent fabric
</pre>

이러한 환경에서는 GPU가 CPU에서 생성한 mapping을 사용할 수 있으며, CPU와 GPU 사이에 hardware cache coherence가 제공될 수 있다.

이 경우 항상 page 전체를 CPU와 GPU 사이에서 migration해야 하는 것은 아니며 remote memory access가 가능하다.