GPU Memory Architecture

개요

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

__global__ void kernel() {
    int value = 10; // 일반적으로 register에 저장
}

• 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

__global__ void kernel() {
    __shared__ int buffer[256];

    buffer[threadIdx.x] = threadIdx.x;
    __syncthreads();
}

• 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

__global__ void kernel() {
    int large_array[1024];
}

• 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

int *ptr;
cudaMalloc(&ptr, sizeof(int) * 1024);

• 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

int *ptr = (int *)malloc(size);
cudaMemcpy(device_ptr, host_ptr, size,
           cudaMemcpyHostToDevice);

• Host Memory는 CPU에 연결된 system DRAM이다.
• 기본적으로 malloc()이나 new로 할당된다.
• GPU와 CPU 사이에서 데이터를 전송하기 위해 일반적으로 cudaMemcpy()를 사용한다.
• 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

int *ptr;
cudaMallocManaged(&ptr, size);

• 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)을 통해 일반 malloc()으로 할당된 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은 이를 이용하여 cudaMemcpyDefault와 같은 기능을 제공한다.

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에 따라 다음과 같은 개념적인 구조가 사용된다.

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

L2 Cache

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

        SM0 ─ L1
          \
        SM1 ─ L1
           \
            ─── L2 Cache ─── GPU DRAM
           /
        SM2 ─ L1

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

GPU MMU

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을 사용할 수 있다.

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

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

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