2D 메모리 할당과 Pitched memory
메모리 접근 패턴이 성능을 결정한다
10,000 × 10,000 행렬 덧셈을 구현했을 때 흥미로운 결과가 나타났다. index = row * ncol + col로 접근하면 3,282 usec가 걸렸지만, 행과 열을 바꿔서 index = col * nrow + row로 접근하니 11,046 usec로 3.4배나 느려졌다. 같은 연산인데 메모리 접근 순서만 바꿨을 뿐인데 왜 이런 차이가 날까?
답은 하드웨어가 메모리를 읽는 방식에 있다. GPU는 메모리를 128바이트 단위로 읽는다. 32개 스레드가 연속된 주소를 읽으면 한 번의 트랜잭션으로 처리되지만, 불연속적으로 읽으면 여러 번의 트랜잭션이 필요하다.
cudaMallocPitch를 사용하면 성능이 더 개선된다. 3,282 usec에서 3,235 usec로 약간의 개선이지만, 이는 메모리를 256바이트 경계로 정렬하기 때문이다. 메모리 정렬은 하드웨어 효율성을 높이는 기법이다.
DRAM 하드웨어 이해
DRAM의 2D 구조
컴퓨터 메모리는 두 가지 기술로 만들어진다. TR(Transistor logic)은 빠르지만 비싸서 CPU 캐시에 사용되고, DRAM은 느리지만 저렴해서 메인 메모리로 사용된다. DRAM은 1비트를 저장하는 데 트랜지스터 1개와 커패시터 1개만 필요해서, SRAM보다 6배 적은 공간을 차지한다.
TR(Transistor logic)이란?
트랜지스터만으로 구성된 메모리 회로를 말한다. SRAM(Static RAM)이 대표적이며, 1비트를 저장하는 데 6개의 트랜지스터가 필요하다. 전원이 공급되는 동안 데이터가 유지되며, 매우 빠르지만 면적을 많이 차지하고 비싸다. 그래서 CPU의 L1, L2, L3 캐시처럼 작지만 빠른 메모리에 사용된다.
DRAM이 2D 구조로 설계된 이유는 무엇일까? 1M 비트를 1D로 배치하면 주소선이 1,000,000개 필요하지만, 1,000 × 1,000 2D 배열로 만들면 행 주소선 1,000개와 열 주소선 1,000개만 있으면 된다. 2,000개로 줄어드는 것이다.
메모리 주소 0xABCD를 읽는다고 하자. 상위 8비트(0xAB)로 행을 선택하고, 하위 8비트(0xCD)로 열을 선택한다. 행을 선택하는 데 수십 나노초가 걸리지만, 같은 행의 다른 열을 읽는 것은 몇 나노초면 된다. 이것이 burst 모드의 핵심이다.
Burst 모드와 순차 접근
프로그램은 대부분 순차적으로 메모리를 읽는다. 배열을 순회하거나 명령어를 실행할 때 연속된 주소를 접근한다. DRAM은 이 특성을 활용해서 한 번 행을 선택하면, 그 행의 여러 열을 빠르게 읽는다.
주소 0x0100, 0x0101, 0x0102를 연속으로 읽는다고 하자. 행 0x01을 한 번만 선택하고, 열 0x00, 0x01, 0x02만 바꾸면 된다. 행 선택이 가장 비싼 연산이므로, 이를 재사용하면 전체 속도가 크게 향상된다.
Burst 모드의 타이밍
메모리 읽기는 크게 두 단계로 나뉜다.
- Delay (지연 시간): 행을 선택하고 데이터를 준비하는 시간. 수십 나노초가 걸린다.
- Data Transfer (데이터 전송): 준비된 행에서 열 데이터를 연속으로 읽는 시간. 각각 몇 나노초만 걸린다.
첫 번째 데이터를 읽을 때는 delay + transfer 시간이 모두 필요하지만, 같은 행의 다음 데이터들은 transfer 시간만 필요하다. 이것이 burst 모드가 빠른 이유다. 한 번의 긴 delay 후 여러 개의 짧은 transfer가 연속으로 이루어진다.
현대 DRAM은 항상 burst 모드로 동작하도록 설계되었다. CPU나 GPU가 1바이트만 요청해도, DRAM은 연속된 64바이트(cache line 크기)를 전송한다. 나머지는 캐시에 저장되어 곧 사용될 가능성이 높다.
DDR - 클럭당 2번 전송
DDR(Double Data Rate)은 클럭의 상승 엣지와 하강 엣지 모두에서 데이터를 전송한다. SDR(Single Data Rate)이 100MHz 클럭에서 초당 100M 데이터를 전송한다면, DDR은 같은 클럭에서 200M 데이터를 전송한다.
하드웨어는 prefetch 구조로 구현된다. DDR은 2개 메모리 배열을 읽어서 MUX로 합치고, DDR2는 4개, DDR3와 DDR4는 8개, DDR5는 16개를 합친다. 클럭 주파수를 높이지 않고도 대역폭을 늘릴 수 있는 방법이다.
GPU 메모리 최적화 원리
DRAM은 burst 모드로 같은 행의 데이터를 빠르게 읽는다. GPU는 이 특성을 최대한 활용하도록 설계되었다. GPU에서 32개 스레드(warp)가 동시에 메모리를 읽을 때, 이들이 연속된 주소를 읽으면 DRAM은 한 행에서 여러 열을 순차적으로 전송한다. 이것이 GPU 메모리 최적화의 핵심이다.
하지만 실제로는 어떻게 동작할까? GPU 메모리 컨트롤러는 여러 스레드의 메모리 요청을 어떻게 처리할까?
Memory Coalescing - 여러 접근을 하나로
Warp란?
GPU에서 함께 실행되는 32개 스레드의 그룹이다. 같은 warp의 스레드들은 같은 명령어를 동시에 실행한다. 예를 들어 1024개 스레드로 구성된 블록은 32개의 warp로 나뉜다(Thread 0-31이 Warp 0, Thread 32-63이 Warp 1, ...).
GPU는 오프칩 메모리를 청크 단위로 접근한다. 최신 GPU는 128바이트 청크를 사용한다. float 하나만 읽어도 실제로는 128바이트 전체를 읽어온다. 나머지 124바이트가 낭비된다.
구체적인 예를 보자. Warp 0의 Thread 0이 주소 0x1000에서 float 하나를 읽는다고 하자.
요청: float 1개 (4바이트)
실제 전송: 128바이트 청크 (주소 0x1000 ~ 0x107F)
낭비: 124바이트
효율: 4 / 128 = 3.1%
하지만 warp의 32개 스레드가 연속된 메모리를 읽으면 어떻게 될까?
Thread 0: 주소 0x1000 읽기 (float)
Thread 1: 주소 0x1004 읽기 (float)
Thread 2: 주소 0x1008 읽기 (float)
...
Thread 31: 주소 0x107C 읽기 (float)
요청: float 32개 (128바이트)
실제 전송: 128바이트 청크 (주소 0x1000 ~ 0x107F)
낭비: 0바이트
효율: 128 / 128 = 100%
32개 float는 정확히 128바이트다. 한 번의 트랜잭션으로 모든 스레드가 필요한 데이터를 받을 수 있다. 이것이 memory coalescing이다.
청크는 경계에 정렬되어야 한다. 128바이트 청크는 128의 배수 주소에서 시작한다. 주소 100부터 227까지 읽으면 두 청크에 걸쳐 있어서 두 번의 트랜잭션이 필요하다.
청크 0: 주소 0 ~ 127 (128바이트)
청크 1: 주소 128 ~ 255 (128바이트)
잘못된 접근: 주소 100 ~ 227
- 청크 0에서 28바이트 (100 ~ 127)
- 청크 1에서 100바이트 (128 ~ 227)
- 트랜잭션 2번 필요
올바른 접근: 주소 128 ~ 255
- 청크 1에서 128바이트 (128 ~ 255)
- 트랜잭션 1번만 필요
주소 128부터 255까지로 정렬하면 한 번의 트랜잭션으로 처리된다. 메모리 시작 주소를 청크 경계에 맞추는 것이 정렬(alignment)의 핵심이다.
CPU의 메모리 정렬
CPU도 메모리 정렬이 중요하다. 32비트 CPU에서 int는 4바이트이므로, 4의 배수 주소에 저장되어야 효율적이다. 정렬되지 않은 주소에 접근하면 일부 CPU는 에러를 발생시키고, 일부는 느리게 처리한다.
SSE는 128비트(16바이트) 레지스터를 사용하므로 16바이트 정렬이 필요하다. AVX는 256비트(32바이트)이므로 32바이트 정렬이 필요하다. GPU는 warp가 32개 스레드를 가지므로 128바이트 정렬이 유리하다.
C++11은 alignas와 alignof 키워드를 제공한다.
int main() {
alignas(16) int a[4];
alignas(1024) int b[4];
printf("%p\\n", a);
printf("%p\\n", b);
}
위 코드에서 a는 16바이트 경계에, b는 1024바이트 경계에 정렬된다. 구조체에도 적용할 수 있다.
struct alignas(16) sse_t {
float sse_data[4];
};
Pitched Memory 사용법
Memory coalescing을 위해서는 메모리가 청크 경계에 정렬되어야 한다. 1D 배열은 cudaMalloc이 자동으로 정렬된 주소를 반환한다. 하지만 2D 행렬은 문제가 다르다. 각 행의 시작 주소도 정렬되어야 하는데, 행의 너비가 청크 크기의 배수가 아니면 자동으로 정렬되지 않는다.
여기서 pitched memory가 등장한다. Pitch는 "행의 실제 너비"를 의미한다. 프로그래머가 요청한 너비보다 큰 값으로, 각 행을 청크 경계에 정렬시킨다.
cudaMallocPitch - 자동 정렬 할당
cudaMalloc으로 10,000 × 10,000 행렬을 할당하는 과정을 보자.
float* dev_mat;
unsigned nrow = 10000;
unsigned ncol = 10000;
cudaMalloc((void**)&dev_mat, nrow * ncol * sizeof(float));
각 행이 10,000개 float를 가지므로 40,000바이트를 차지한다. 메모리 레이아웃은 다음과 같다.
행 0: 주소 0 ~ 39,999 (40,000바이트)
행 1: 주소 40,000 ~ 79,999 (40,000바이트)
행 2: 주소 80,000 ~ 119,999 (40,000바이트)
...
40,000을 128로 나누면 312.5다. 행 1의 시작 주소 40,000은 128의 배수가 아니다. 청크 경계에 정렬되지 않았다.
청크 312: 주소 39,936 ~ 40,063
- 행 0의 마지막 64바이트 포함
- 행 1의 처음 64바이트 포함
- 두 행이 하나의 청크에 섞임
행 1의 첫 번째 원소를 읽으려면 청크 312를 읽어야 하는데, 이 청크에는 행 0의 데이터도 포함되어 있다. Memory coalescing 효율이 떨어진다.
cudaMallocPitch는 이 문제를 해결한다. 각 행을 256바이트 경계로 정렬해서 할당한다.
float* dev_mat;
size_t pitch;
unsigned nrow = 10000;
unsigned ncol = 10000;
cudaMallocPitch((void**)&dev_mat, &pitch, ncol * sizeof(float), nrow);
printf("Allocated pitch: %zu bytes\\n", pitch);
함수는 필요한 너비(40,000바이트)를 256의 배수로 올림해서 할당한다.
요청한 너비: 10,000 float = 40,000바이트
256의 배수로 올림: 40,000 → 40,192바이트
pitch = 40,192바이트
행 0: 주소 0 ~ 40,191 (40,192바이트, 실제 사용 40,000)
행 1: 주소 40,192 ~ 80,383 (40,192바이트, 실제 사용 40,000)
행 2: 주소 80,384 ~ 120,575 (40,192바이트, 실제 사용 40,000)
...
40,192를 256으로 나누면 정확히 157이다. 모든 행이 256바이트 경계에서 시작한다. 각 행의 마지막 192바이트는 사용되지 않는 패딩이지만, memory coalescing 효율이 크게 향상된다.
실제 pitch 값은 시스템에 따라 다르지만, 항상 256의 배수로 정렬된다. 함수가 pitch를 반환하므로, 프로그래머가 미리 계산할 필요가 없다.
2D 행렬에서 (row, col) 원소의 주소는 다음과 같이 계산한다.
T* pElem = (T*)((char*)baseAddr + row * pitch) + col;
이 공식이 복잡해 보이는 이유는 pitch가 바이트 단위이고, col은 원소 단위이기 때문이다. 단계별로 분해해보자.
행 2, 열 5의 float 원소를 읽는다고 하자 (pitch = 40192).
1. 해당 행의 시작 주소 계산 (바이트 단위)
행 2는 2 * 40192 = 80384바이트 떨어진 곳에 시작
2. baseAddr을 char*로 캐스팅
char*는 1바이트씩 이동하므로 바이트 단위 연산 가능
(char*)baseAddr + 80384 → 행 2의 시작 주소
3. 다시 T*로 캐스팅
(T*)((char*)baseAddr + 80384) → 행 2의 첫 번째 원소 주소
4. 열 인덱스 추가 (원소 단위)
pElem = ... + 5
float*이므로 5칸 이동 (5 * sizeof(float) = 20바이트)
왜 char*로 캐스팅할까? C++에서 포인터 연산은 타입 크기만큼 이동한다.
float* ptr = baseAddr;
ptr + 1; // 4바이트 이동 (sizeof(float) = 4)
ptr + 2; // 8바이트 이동
char* ptr = (char*)baseAddr;
ptr + 1; // 1바이트 이동
ptr + 40192; // 정확히 40192바이트 이동
pitch는 바이트 단위 값이므로, float* + pitch를 하면 pitch * 4 바이트만큼 이동해서 잘못된 주소가 된다. char*로 캐스팅하면 정확히 pitch 바이트만큼 이동할 수 있다.
cudaMemcpy2D - pitch를 고려한 복사
일반 cudaMemcpy는 pitch를 고려하지 않는다. cudaMemcpy2D를 사용해야 한다.
__host__ cudaError_t cudaMemcpy2D(void* dst, size_t dpitch,
const void* src, size_t spitch,
size_t width, size_t height,
cudaMemcpyKind kind);
dst와 dpitch는 목적지 포인터와 pitch다. src와 spitch는 소스 포인터와 pitch다. width는 각 행에서 복사할 바이트 수이고, height는 행의 개수다.
호스트에서 디바이스로 복사하는 예제를 보자.
cudaMemcpy2D(dev_mat, dev_pitch,
host_mat, ncol * sizeof(float),
ncol * sizeof(float), nrow,
cudaMemcpyHostToDevice);
파라미터를 하나씩 살펴보자.
dst: dev_mat - 목적지 포인터 (device)
dpitch: dev_pitch - 목적지 pitch (예: 40192바이트)
src: host_mat - 소스 포인터 (host)
spitch: ncol * sizeof(float) - 소스 pitch (40,000바이트)
width: ncol * sizeof(float) - 복사할 너비 (40,000바이트)
height: nrow - 복사할 행 개수 (10,000)
host_mat는 일반 malloc으로 할당했으므로 패딩이 없다. 각 행이 정확히 ncol * sizeof(float) 바이트다. 따라서 spitch = ncol * sizeof(float)를 사용한다.
dev_mat는 cudaMallocPitch로 할당했으므로 패딩이 있다. 각 행이 dev_pitch 바이트다 (40,192바이트). 따라서 dpitch = dev_pitch를 사용한다.
width는 실제로 복사할 너비다. 패딩을 제외한 데이터 영역만 복사하므로 ncol * sizeof(float) (40,000바이트)다.
함수는 각 행을 개별적으로 복사한다.
행 0 복사:
src: host_mat + 0 * 40000 = host_mat + 0
dst: dev_mat + 0 * 40192 = dev_mat + 0
크기: 40,000바이트
행 1 복사:
src: host_mat + 1 * 40000 = host_mat + 40000
dst: dev_mat + 1 * 40192 = dev_mat + 40192
크기: 40,000바이트
행 2 복사:
src: host_mat + 2 * 40000 = host_mat + 80000
dst: dev_mat + 2 * 40192 = dev_mat + 80384
크기: 40,000바이트
...
호스트의 연속된 메모리를 디바이스의 pitched 메모리로 변환한다. 각 행의 패딩 영역(192바이트)은 복사되지 않는다.
Pitched Memory 사용 커널
커널 함수는 pitch 매개변수를 받아서 인덱스를 계산한다.
__global__ void kernel_matadd(float* c, const float* a, const float* b,
unsigned nrow, unsigned ncol, size_t dev_pitch) {
register unsigned col = blockIdx.x * blockDim.x + threadIdx.x;
if (col < ncol) {
register unsigned row = blockIdx.y * blockDim.y + threadIdx.y;
if (row < nrow) {
register unsigned offset = row * dev_pitch + col * sizeof(float);
*((float*)((char*)c + offset)) = *((const float*)((const char*)a + offset))
+ *((const float*)((const char*)b + offset));
}
}
}
위 커널 코드를 단계별로 분석해보자. Thread (5, 2)가 실행된다고 하자 (dev_pitch = 40192).
1. 열 인덱스 계산
col = blockIdx.x * 32 + threadIdx.x
block과 thread에 따라 결정됨 (여기서는 5라고 가정)
2. 행 인덱스 계산
row = blockIdx.y * 32 + threadIdx.y
block과 thread에 따라 결정됨 (여기서는 2라고 가정)
3. offset 계산 (바이트 단위)
offset = row * dev_pitch + col * sizeof(float)
offset = 2 * 40192 + 5 * 4
offset = 80384 + 20 = 80404바이트
4. 메모리 접근
(char*)c + 80404 → 바이트 단위로 이동
(float*)(...) → float* 타입으로 해석
*(...) → 값 읽기/쓰기
offset은 바이트 단위로 계산된다. row * dev_pitch는 해당 행의 시작 위치(바이트)를 가리키고, col * sizeof(float)는 그 행에서 col 번째 원소의 오프셋(바이트)이다.
왜 col * sizeof(float)를 더할까? dev_pitch는 행의 시작 주소만 제공한다. 행 안에서 원하는 열로 이동하려면 col * sizeof(float) 바이트를 추가로 더해야 한다.
행 2 시작 주소: 80384바이트
열 5 오프셋: 5 * 4 = 20바이트
최종 주소: 80384 + 20 = 80404바이트
이전 버전은 i = row * ncol + col을 사용했지만, pitched 버전은 dev_pitch를 사용한다. ncol은 원소 개수였지만 dev_pitch는 바이트 개수이므로 단위가 다르다.
이전 방식 (pitched 메모리 아님):
i = row * ncol + col = 2 * 10000 + 5 = 20005 (원소 인덱스)
주소 = baseAddr + 20005 * sizeof(float)
Pitched 메모리 방식:
offset = row * dev_pitch + col * sizeof(float)
= 2 * 40192 + 5 * 4 = 80404 (바이트 오프셋)
주소 = (char*)baseAddr + 80404
메인 함수에서는 다음과 같이 사용한다.
float *dev_matA, *dev_matB, *dev_matC;
size_t dev_pitch;
cudaMallocPitch((void**)&dev_matA, &dev_pitch, ncol * sizeof(float), nrow);
cudaMallocPitch((void**)&dev_matB, &dev_pitch, ncol * sizeof(float), nrow);
cudaMallocPitch((void**)&dev_matC, &dev_pitch, ncol * sizeof(float), nrow);
cudaMemcpy2D(dev_matA, dev_pitch, matA, ncol * sizeof(float),
ncol * sizeof(float), nrow, cudaMemcpyHostToDevice);
dim3 dimBlock(32, 32, 1);
dim3 dimGrid((ncol + dimBlock.x - 1) / dimBlock.x,
(nrow + dimBlock.y - 1) / dimBlock.y, 1);
kernel_matadd <<<dimGrid, dimBlock>>> (dev_matC, dev_matA, dev_matB,
nrow, ncol, dev_pitch);
GeForce RTX 2070에서 3,235 usec가 소요된다. 이전 버전(3,282 usec)보다 47 usec 빠르다. 약간의 개선이지만, 더 큰 행렬이나 복잡한 연산에서는 차이가 커진다.
메모리 정렬은 일부 공간을 낭비하지만 성능 향상이 더 크다. 10,000 × 10,000 행렬에서 240바이트 * 10,000 = 2.4MB가 추가로 필요하지만, 원래 400MB에서 0.6% 증가에 불과하다. 메모리는 저렴하고 성능은 귀중하다.
관련 문서
- CUDA Runtime API - cudaMallocPitch - Pitched Memory 할당 API 레퍼런스
- CUDA Runtime API - cudaMemcpy2D - 2D 메모리 복사 API
- CUDA C Programming Guide - Memory Coalescing - 메모리 병합 접근과 성능 최적화
- CUDA C Best Practices Guide - Memory Optimizations - 메모리 정렬과 최적화 기법
- NVIDIA Blog - How to Access Global Memory Efficiently - Global Memory 효율적 접근 방법