9단계 GPU 워크플로우

모든 GPU computation의 여정 — 한 번 외우면 평생 씀

Host 메모리에서 검증된 결과까지 GPU computation은 항상 같은 박스 9개 통과해. GEMM이 새 박스 추가 안 해 — 기존 거 더 크고 바쁘게 만들 뿐.

단계	할 일	CUDA	Metal (unified memory가 단순화)
1	Host 메모리 할당	`malloc` / `std::vector`	Swift `Array` 또는 step 3에서 처리
2	입력 데이터 초기화	loop / `std::iota` / RNG	동일
3	Device 메모리 할당	`cudaMalloc`	`device.makeBuffer(.storageModeShared)`
4	Host → Device 복사	`cudaMemcpyAsync`	no-op (buffer에 바로 write)
5	커널 정의	`__global__ void kernel`	`kernel void kernel`
6	커널 launch	`kernel<<<G,B>>>()`	`encoder.dispatchThreadgroups`
7	Device → Host 복사	`cudaMemcpyAsync` back	no-op (같은 buffer에서 read)
8	검증	host loop / `thrust::equal`	Swift loop
9	Garbage-collect	`cudaFree` + `free`	buffer = nil (ARC 처리)

Apple Silicon에선 unified memory 때문에 step 4랑 7이 no-op (Track 2 다룸). 개념적 단계는 여전히 존재 — 비용이 0일 뿐. 멘털 모델에서 빼지 마; 나중에 CUDA 코드 읽을 때 고마워할 거.

진짜 production 복잡함은 step 6 (launch geometry 선택)이랑 step 8 (어느 검증 전략: memcmp, thrust::equal, CPU reference 대비 numpy.allclose, FP에 대한 numerical-tolerance 비교)에 숨어.

Code

9단계 전부 한 CUDA 프로그램에·cuda

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

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

int main() {
    const size_t N = 1 << 22;             // 4M 원소

    // 1+2: host 할당 + 초기화
    std::vector<float> hA(N, 1.f), hB(N, 2.f), hC(N, 0.f);

    // 3: device 할당
    float *dA, *dB, *dC;
    cudaMalloc(&dA, N*sizeof(float));
    cudaMalloc(&dB, N*sizeof(float));
    cudaMalloc(&dC, N*sizeof(float));

    // 4: H → D
    cudaMemcpy(dA, hA.data(), N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dB, hB.data(), N*sizeof(float), cudaMemcpyHostToDevice);

    // 5+6: 커널 + launch
    int threads = 256;
    int blocks  = (N + threads - 1) / threads;
    vec_add<<<blocks, threads>>>(dA, dB, dC, N);
    cudaDeviceSynchronize();

    // 7: D → H
    cudaMemcpy(hC.data(), dC, N*sizeof(float), cudaMemcpyDeviceToHost);

    // 8: 검증
    bool ok = true;
    for (size_t i = 0; i < N; ++i) if (hC[i] != 3.f) { ok = false; break; }
    printf("Validation: %s\n", ok ? "PASS" : "FAIL");

    // 9: free
    cudaFree(dA); cudaFree(dB); cudaFree(dC);
    return 0;
}

Apple Silicon에서 같은 워크플로우 — 코드론 7단계·swift

import Metal
import Foundation

let device = MTLCreateSystemDefaultDevice()!
let n = 1 << 22
let bytes = n * MemoryLayout<Float>.stride

// 1+2+3+4 합쳐: shared buffer 할당, 바로 write
let bufA = device.makeBuffer(length: bytes, options: .storageModeShared)!
let bufB = device.makeBuffer(length: bytes, options: .storageModeShared)!
let bufC = device.makeBuffer(length: bytes, options: .storageModeShared)!
let pA = bufA.contents().bindMemory(to: Float.self, capacity: n)
let pB = bufB.contents().bindMemory(to: Float.self, capacity: n)
let pC = bufC.contents().bindMemory(to: Float.self, capacity: n)
for i in 0..<n { pA[i] = 1; pB[i] = 2 }

// 5+6: launch (커널 + encoder 코드는 간결성을 위해 생략)
// 7: no-op — pC는 GPU가 막 쓴 같은 RAM
// 8: 검증
let ok = (0..<n).allSatisfy { pC[$0] == 3 }
print("Validation: \(ok ? "PASS" : "FAIL")")
// 9: scope 끝나면 ARC가 buffer drop

Exercise

위 CUDA 9단계 프로그램에서 D→H copy 전 cudaDeviceSynchronize() 일부러 제거. 10번 돌려. 정확한 출력, 부분 출력, 또는 0 — timing 따라 다름. 다음 sync 복원하고 다시 돌려; 출력 일관됨. 이게 CUDA 코드의 가장 흔한 '내 머신엔 되는데 cluster에서 깨지는' 버그 중 하나.

모든 GPU computation의 여정 — 한 번 외우면 평생 씀

Code

External links

Exercise

Progress

댓글 0