Operator 변주 & 용량 zone

커널 한 개, 맛 4가지 — 그리고 사이즈 reality check

Vector add 되면 operator만 바꾸면 sub, mul, div 공짜로 얻어. launch geometry, 메모리 layout, 검증 동일 — ALU op만 바뀜. 근데 비용은 항상 동일 아님:

Op	본문 변경	비고
Add	`C[i] = A[i] + B[i]`	Baseline.
Sub	`+` → `-`	Add랑 같은 throughput — single FP unit.
Mul	`+` → `*`	FMA pipeline 사용, 같은 throughput.
Div	`+` → `/`	5–10× 느림 — non-pipelined reciprocal 때문. NaN/inf 피하려면 `B[i] == 0` 가드.

VRAM 용량 zone — 너무 많이 요청하면 무슨 일

Zone	일어나는 일	탈출구
🟢 Green: VRAM에 들어감	풀 속도, 데이터 GPU에 resident	조치 불필요
🟡 Yellow: UVM paging	Unified Memory가 페이지 migrate — 10–50× 느려짐	청크로 처리
🔴 Red: 하드 alloc 실패	`cudaMalloc`이 `cudaErrorMemoryAllocation` return	Slice-and-stream loop
🟣 Crimson: grid 한계	`gridDim.x`가 2³¹⁻¹ 초과	Outer host loop, base-index arg

Apple Silicon은 zone 자체가 옮겨감: 분리 VRAM 없으니까 'Green'이 OS랑 다른 앱이 잡고 있는 거 빼고 unified pool 전체로 확장. Yellow zone (swap에 paging)은 여전히 존재하는데 훨씬 늦게 발동.

Code

Op별 커널 한 개 generic 패턴 (CUDA)·cuda

// Template + functor 또는 preprocessor.
// Functor 버전이 디버거 친화적:

struct AddOp { __device__ float operator()(float a, float b) const { return a + b; } };
struct SubOp { __device__ float operator()(float a, float b) const { return a - b; } };
struct MulOp { __device__ float operator()(float a, float b) const { return a * b; } };
struct DivOp { __device__ float operator()(float a, float b) const {
    return b == 0.f ? 0.f : a / b;   // 가드
} };

template <typename Op>
__global__ void elementwise(const float* A, const float* B, float* C, size_t N, Op op) {
    size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = op(A[i], B[i]);
}

// 원하는 Op로 launch; nvcc가 호출 inline:
// elementwise<<<blocks, threads>>>(dA, dB, dC, N, MulOp{});

데이터가 VRAM 초과할 때 slice-and-stream·cuda

// VRAM에 들어갈 CHUNK 크기로 N 원소 처리.
const size_t CHUNK = 1 << 24;        // 16M 원소 ≈ buffer당 64 MB
for (size_t off = 0; off < N; off += CHUNK) {
    size_t this_n = std::min(CHUNK, N - off);

    cudaMemcpyAsync(dA, hA + off, this_n*sizeof(float),
                    cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(dB, hB + off, this_n*sizeof(float),
                    cudaMemcpyHostToDevice, stream);

    int threads = 256;
    int blocks  = (this_n + threads - 1) / threads;
    vec_add<<<blocks, threads, 0, stream>>>(dA, dB, dC, this_n);

    cudaMemcpyAsync(hC + off, dC, this_n*sizeof(float),
                    cudaMemcpyDeviceToHost, stream);
}
cudaStreamSynchronize(stream);

Exercise

같은 64M-원소 벡터에 operator 4개 (add, sub, mul, div) 다 timing. add ≈ sub ≈ mul, div가 눈에 띄게 느려야 함 (GPU 세대에 따라 3–7× 자주). 그 다음 VRAM의 2× 크기 vector size 골라서 커널이 alloc error로 crash (CUDA)하거나 paging으로 기어가는 (Apple Silicon) 거 봐. Slice-and-stream 버전 구현해서 in-VRAM 케이스랑 같은 원소당 속도 나오는지 확인.

커널 한 개, 맛 4가지 — 그리고 사이즈 reality check

VRAM 용량 zone — 너무 많이 요청하면 무슨 일

Code

External links

Exercise

Progress

댓글 0