The 9-Step GPU Workflow

The journey every GPU computation makes — memorize once, write forever

From host memory to verified result, a GPU computation always passes through the same nine boxes. GEMM doesn't add new boxes; it just makes the existing ones bigger and busier.

Step	What you do	CUDA	Metal (unified memory simplifies)
1	Allocate host memory	`malloc` / `std::vector`	Swift `Array` or step 3 covers it
2	Initialize input data	loop / `std::iota` / RNG	same
3	Allocate device memory	`cudaMalloc`	`device.makeBuffer(.storageModeShared)`
4	Copy host → device	`cudaMemcpyAsync`	no-op (write directly into the buffer)
5	Define kernel	`__global__ void kernel`	`kernel void kernel`
6	Launch kernel	`kernel<<<G,B>>>()`	`encoder.dispatchThreadgroups`
7	Copy device → host	`cudaMemcpyAsync` back	no-op (read from same buffer)
8	Validate	host loop / `thrust::equal`	Swift loop
9	Garbage-collect	`cudaFree` + `free`	buffer = nil (ARC handles it)

On Apple Silicon, steps 4 and 7 are no-ops because of unified memory (we covered this in Track 2). The conceptual steps still exist — they're just zero-cost. Don't skip them in your mental model; you'll thank yourself when you read CUDA code later.

The real production complexity hides in step 6 (launch geometry choices) and step 8 (which validation strategy: memcmp, thrust::equal, numpy.allclose against a CPU reference, or numerical-tolerance comparison for FP).

Code

All 9 steps in one CUDA program·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 elements

    // 1+2: host alloc + init
    std::vector<float> hA(N, 1.f), hB(N, 2.f), hC(N, 0.f);

    // 3: device alloc
    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: kernel + 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: validate
    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;
}

Same workflow on Apple Silicon — 7 steps in code·swift

import Metal
import Foundation

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

// 1+2+3+4 collapse: allocate shared buffer, write directly
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 (kernel + encoder code omitted for brevity)
// 7: no-op — pC is the same RAM the GPU just wrote
// 8: validate
let ok = (0..<n).allSatisfy { pC[$0] == 3 }
print("Validation: \(ok ? "PASS" : "FAIL")")
// 9: ARC drops buffers when scope ends

Exercise

Take the CUDA 9-step program above and intentionally remove the cudaDeviceSynchronize() before the D→H copy. Run it 10 times. You may see correct output, partial output, or zeros — depending on timing. Now restore the sync and run again; output is consistent. This is one of the most common 'works on my machine, breaks on the cluster' bugs in CUDA code.

The journey every GPU computation makes — memorize once, write forever

Code

External links

Exercise

Progress

Comments 0