연산자 변주 & 용량 지대

커널 한 개, 맛 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	일어나는 일	탈출구
🟢 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은 지대 자체가 옮겨가. 분리된 VRAM이 없으니까 'Green'이 OS랑 다른 앱이 잡고 있는 몫만 빼고 unified pool 전체로 넓어져. Yellow 지대 (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 용량 지대 — 너무 많이 요청하면 무슨 일

Code

External links

Exercise

Progress

댓글 0