CUDA Pitch, cudaMallocPitch() - 공부하는 도비

2025.01.09 - [CUDA] - CUDA 디바이스 메모리 할당, 초기화 및 해제 - 공부하는 도비

CUDA 디바이스 메모리 할당, 초기화 및 해제 - 공부하는 도비

 

저번 피드에서 호스트에서 디바이스로 메모리를 할당하고, 간단히 출력해 보는 걸 해보았다. 

그때 사용된 함수는 cudaMalloc()인데, 이 함수는 1D처럼 단순한 데이터일 때 사용이 된다. 

 

2D, 3D 구조에 메모리 할당을 하고자 한다면 피드 제목에 적힌 Pitch라고 하는 개념을 알고 있어야 한다. 

 

CUDA에서 Pitch 란? 

Pitch : 각 행(row) 간의 간격을 의미하고, 이 값은 각 행의 "실제 메모리 크기"를 나타낸다.  

 

CUDA에서 2D 메모리를 다룰 때, 행렬의 각 행이 메모리에서 연속적으로 저장되지 않을 수가 있는데, 

이를 해결하기 위해 CUDA는 각 행의 시작 주소가 정렬(alignement)된 상태를 유지하도록 padding을 추가하는데, 이 padding을 Pitch 또는 Stride라고 한다. 

 

cudaMallocPitch() 란? 

이 함수는 pitch를 고려하여 메모리를 할당해준다. 

이 함수의 원형은 다음과 같다. 

 

cudaError_t cudaMallocPitch(
    void **devPtr,       // 할당된 메모리의 시작 주소(포인터)
    size_t *pitch,       // 행 간의 실제 메모리 간격(바이트 단위)
    size_t width,        // 메모리의 논리적 너비(바이트 단위)
    size_t height        // 메모리의 높이(행의 개수)
);

 

단위를 잘 봐야 한다. pitch 또한 바이트 단위기 때문에, 코드를 보면 size_t 변수를 통해 선언된 걸 확인할 수 있다.  

 

아래 예제는 cudaMallocPitch() 함수에 대한 코드이다. 

#include <cuda_runtime.h>
#include <stdio.h>

#define WIDTH  5
#define HEIGHT 5
/*
cudaError_t cudaMallocPitch(
    void **devPtr,       // 할당된 메모리의 시작 주소(포인터)
    size_t *pitch,       // 행 간의 실제 메모리 간격(바이트 단위)
    size_t width,        // 메모리의 논리적 너비(바이트 단위)
    size_t height        // 메모리의 높이(행의 개수)
);
*/

int main() {
    int *d_matrix;
    size_t pitch;

    // cudaMallocPitch()는 pitch를 고려하여 메모리를 할당함
    cudaMallocPitch((void**)&d_matrix, &pitch, WIDTH * sizeof(int), HEIGHT);
   
    printf("Pitch (row size in bytes): %zu\n", pitch);

    // 메모리 해제
    cudaFree(d_matrix);

    return 0;
}

출력 값

메모리를 할당한 후 그 때 사용된 Pitch를 출력해 보았다. 

CUDA 장치 메모리는 일반적으로 32, 64, 128 바이트로 정렬이 되어 있다.

NVIDIA GPU의 경우, 32바이트 정렬이 기본 값이며, 이는 pitch가 32의 배수가 됨을 의미한다. 

 

즉, Pitch 값이 512로 나오는 이유는 입력된 width와 데이터 타입의 크기를 곱한 값이 32의 배수가 되도록 padding이 추가되었기 때문이다. 

당연 논리적인 너비는 5(width) * 4(sizeof(int) = 4) = 20바이트이다. 

 

Pitch 값이 클 경우,

즉 논리적인 너비 보다 크면, 메모리 사용량이 증가하지만, 정렬 덕분에 메모리 접근 속도가 빨라진다는 장점이 있다.