GPU 연산 시 메모리 주소 설정을 통한 빠른 접근, 행렬 곱으로의 확장

메모리 접근 정렬 Coalesced Access

GPU 는 32개의 쓰레드가 함께 움직이는 워프 warp 단위로 메모리를 읽는다. 이 때,

• 쓰레드들이 연속적인 주소를 동시에 읽을 때 
  • 하나의 메모리 트랜잭션으로 모두 처리 가능 -> 빠름
• 멀리 떨어진 주소를 접근할 때
  • 여러 개의 트랜잭션 발생 -> 느림

주소 차이가 작을수록 ( 연속적일수록 ) 메모리 접근이 coalesced 되며 매우 빠르게 처리된다.

 

테스트 코드

#include <iostream>
#include <cuda_runtime.h>

#define N (1024 * 1024 * 10)  // 10M 요소
#define THREADS_PER_BLOCK 256
#define REPEATS 1000

__global__ void coalesced_read(float* input, float* output) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N)
        output[idx] = input[idx];
}

__global__ void non_coalesced_read(float* input, float* output) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N)
        output[idx] = input[(idx * 128) % N];  // 큰 stride로 non-coalesced 접근
}

template<typename Kernel>
void benchmark(const char* label, Kernel kernel, float* input, float* output) {
    cudaEvent_t start, end;
    cudaEventCreate(&start);
    cudaEventCreate(&end);

    int blocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;

    cudaEventRecord(start);
    for (int i = 0; i < REPEATS; ++i) {
        kernel<<<blocks, THREADS_PER_BLOCK>>>(input, output);
    }
    cudaEventRecord(end);
    cudaEventSynchronize(end);

    float elapsed;
    cudaEventElapsedTime(&elapsed, start, end);

    // 최적화 방지를 위한 출력 체크
    float* host_output = new float[10];
    cudaMemcpy(host_output, output, sizeof(float) * 10, cudaMemcpyDeviceToHost);
    float checksum = 0;
    for (int i = 0; i < 10; ++i)
        checksum += host_output[i];
    delete[] host_output;

    std::cout << label << " Time over " << REPEATS << " runs: " << elapsed << " ms, Checksum: " << checksum << std::endl;

    cudaEventDestroy(start);
    cudaEventDestroy(end);
}

int main() {
    float *input, *output;
    cudaMalloc(&input, N * sizeof(float));
    cudaMalloc(&output, N * sizeof(float));

    // 간단한 초기화 (모든 값 1.0f)
    float* host_input = new float[N];
    for (int i = 0; i < N; ++i) host_input[i] = 1.0f;
    cudaMemcpy(input, host_input, N * sizeof(float), cudaMemcpyHostToDevice);
    delete[] host_input;

    benchmark("Coalesced", coalesced_read, input, output);
    benchmark("Non-Coalesced", non_coalesced_read, input, output);

    cudaFree(input);
    cudaFree(output);
    return 0;
}

 

Coalesced vs Non-Coalesced GPU 메모리 접근 속도 비교 코드

 

coalesced_read : 연속 주소 접근

non_coalesced_read : 분산 주소 접근

이를 통해 메모리 접근 패턴 차이만으로도 수 배의 성능 차이가 발생함을 확인할 수 있다.

 

#define N (1024 * 1024 * 10)

총 10M 개의 float 요소를 GPU 에서 처리한다.

float 는 4 바이트 -> 총 40MB 데이터

 

cudaMalloc(&input, N * sizeof(float));
cudaMalloc(&output, N * sizeof(float));

GPU 전용 global memory 를 2개 확보한다. 입력용, 출력용

 

float* host_input = new float[N];
for (int i = 0; i < N; ++i) host_input[i] = 1.0f;
cudaMemcpy(input, host_input, N * sizeof(float), cudaMemcpyHostToDevice);

CPU 에서 값을 설정한 후 GPU 로 복사 (cudaMemcpyHostToDevice)

이렇게 하면 GPU 에서의 연산 결과가 예측 가능 (모두 더하면 항상 일정)

 

Matrix Multiplication 구조에선 메모리 접근 방식이 전체 연산 성능에 미치는 영향이 더 극명하게 드러난다.

 

두 행렬 A, B 를 곱할 때, A 는 행 기준 접근, B 는 열 기준 접근한다.

C 언어나 CUDA 에서 2차원 배열은 일반적으로 row-major 로 저장되어 A 에는 연속적 접근으로 빠르고 B 는 비연속 접근으로 느림

 

때문에 CUDA 에서 고성능 GEMM 구현은 반드시 B 를 Transpose 하거나 shared memory 를 사용해서 메모리 접근 패턴을 바꿔야 한다.

 

행을 먼저 읽으면서 그 값을 누적시켜 연산 최적화 수행