Global Memory Coalescing ( 연속 접근 vs Stride 접근) test

#include <cstdio>
#include <cuda_runtime.h>

constexpr int N                 = 1 << 24;  // 16M elements
constexpr int THREADS_PER_BLOCK = 256;
constexpr int STRIDE            = 32;       // warp size와 동일 → non-coalesced 유도

// ------------------------------------------------------------
// Coalesced: thread i -> in[i]
// ------------------------------------------------------------
__global__
void coalesced_read_kernel(const float* __restrict__ in,
                           float* __restrict__ out,
                           int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float acc = 0.0f;

    // grid-stride loop
    for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
        acc += in[i];     // 연속 접근
    }
    if (tid < n) {
        out[tid] = acc;   // 최적화 방지용
    }
}

// ------------------------------------------------------------
// Strided: thread i -> in[i * STRIDE]
//  - warp 내 인접 thread들이 서로 멀리 떨어진 주소를 읽도록 만들어
//    transaction 수를 늘림
// ------------------------------------------------------------
__global__
void strided_read_kernel(const float* __restrict__ in,
                         float* __restrict__ out,
                         int n, int stride)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float acc = 0.0f;

    for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
        int idx = (i * stride) % n;  // 범위 fold
        acc += in[idx];              // 비-coalesced 접근
    }
    if (tid < n) {
        out[tid] = acc;
    }
}

int main()
{
    printf("== Global Memory Coalescing Test ==\n");

    size_t bytes = N * sizeof(float);

    // Host 메모리
    float* h_in  = (float*)malloc(bytes);
    float* h_out = (float*)malloc(bytes);

    for (int i = 0; i < N; ++i)
        h_in[i] = 1.0f;

    // Device 메모리
    float* d_in  = nullptr;
    float* d_out = nullptr;
    cudaMalloc(&d_in,  bytes);
    cudaMalloc(&d_out, bytes);

    cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice);

    // 런치 파라미터
    int blocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
    blocks = min(blocks, 256);  // 너무 큰 grid는 적당히 제한

    dim3 grid(blocks);
    dim3 block(THREADS_PER_BLOCK);

    // 타이밍용 cudaEvent
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    // warm-up
    for (int i = 0; i < 3; ++i) {
        coalesced_read_kernel<<<grid, block>>>(d_in, d_out, N);
        strided_read_kernel<<<grid, block>>>(d_in, d_out, N, STRIDE);
    }
    cudaDeviceSynchronize();

    // =============================
    // 1) Coalesced
    // =============================
    cudaEventRecord(start);
    for (int i = 0; i < 10; ++i) {
        coalesced_read_kernel<<<grid, block>>>(d_in, d_out, N);
    }
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

    float ms_coalesced = 0.0f;
    cudaEventElapsedTime(&ms_coalesced, start, stop);
    ms_coalesced /= 10.0f;

    double gbytes = (double)bytes / 1e9;
    double bw_coalesced = gbytes / (ms_coalesced / 1e3);

    printf("[Coalesced]\n");
    printf("  N            = %d\n", N);
    printf("  Time (ms)    = %.3f\n", ms_coalesced);
    printf("  Bandwidth    = %.2f GB/s\n\n", bw_coalesced);

    // =============================
    // 2) Strided
    // =============================
    cudaEventRecord(start);
    for (int i = 0; i < 10; ++i) {
        strided_read_kernel<<<grid, block>>>(d_in, d_out, N, STRIDE);
    }
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

    float ms_strided = 0.0f;
    cudaEventElapsedTime(&ms_strided, start, stop);
    ms_strided /= 10.0f;

    double bw_strided = gbytes / (ms_strided / 1e3);

    printf("[Strided] (stride = %d)\n", STRIDE);
    printf("  N            = %d\n", N);
    printf("  Time (ms)    = %.3f\n", ms_strided);
    printf("  Bandwidth    = %.2f GB/s\n\n", bw_strided);

    cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost);
    printf("Sample output h_out[0] = %f\n", h_out[0]);

    // 정리
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    cudaFree(d_in);
    cudaFree(d_out);
    free(h_in);
    free(h_out);

    return 0;
}

// 빌드 예시:
// nvcc -O3 -arch=sm_86 global_coalescing_test.cu -o global_coalescing_test.exe
//
// Nsight Compute 예시(요약만 보고 싶을 때):
// ncu --set speedOfLight --kernel-name regex:.*coalesced.* ./global_coalescing_test.exe
// ncu --set speedOfLight --kernel-name regex:.*strided.*   ./global_coalescing_test.exe

 

실행 결과

== Global Memory Coalescing Test ==
[Coalesced]
  N            = 16777216
  Time (ms)    = 0.166
  Bandwidth    = 404.79 GB/s

[Strided] (stride = 32)
  N            = 16777216
  Time (ms)    = 1.245
  Bandwidth    = 53.89 GB/s

Sample output h_out[0] = 256.000000

 

Coalesced 패턴

• warp 내 32 threads 가 연속된 주소를 읽음
• 하드웨어 입장에서 
  • 하나의 큰 연속 블록 - 소수의  DRAM transaction 으로 해결
• 결과
  • DRAM bus 가 꽉 차게 활용됨

 

Strided 패턴

• 같은 warp 안에서조차, 서로 다른 cache line / segment 를 찍고 있어서 
  • transaction 이 1개가 아니라 여러 개로 쪼개짐

 

GEMM / Conv / Attention 최적화에서

“먼저 global coalescing 맞추고 → 그 다음 shared tiling” 이라는 순서를 가져가는 이유가

이 실험 하나로 아주 명확하게 드러난다.