CUDA ↔ Metal Cheat Sheet

Same colony, different uniforms

Once you internalize this table, you can read either codebase in five minutes. The hardware concepts are identical; only the spellings differ.

CUDA	Metal	Notes
`__global__ void kernel(...)`	`kernel void kernel(...)`	GPU entry point launched by CPU.
`__device__ float helper()`	(no qualifier — default for non-kernel)	GPU-only helper.
`kernel<<<grid, block>>>()`	`encoder.dispatchThreadgroups(grid, threadsPerTG)`	Same two-level launch geometry.
`blockIdx.x` / `threadIdx.x`	`[[threadgroup_position_in_grid]]` / `[[thread_index_in_threadgroup]]`	SIMT indices.
warp (32 threads)	SIMD-group (32 threads)	Divergence + shuffle math identical.
`__shared__` memory	`threadgroup` memory	On-chip SRAM scratchpad shared by a (thread)group.
`__constant__` memory	`constant` address space	Small read-only LUTs.
`__syncthreads()`	`threadgroup_barrier(mem_flags::mem_threadgroup)`	Local sync + optional memory fence.
`cudaMalloc` / `cudaMemcpy`	`device.makeBuffer(.storageModeShared)`	Apple Silicon: no copy needed.
cuBLAS	MPSMatrixMultiplication	Vendor BLAS, both peak ~80% of theoretical.
nvcc	xcrun metal + metallib	Two-step build (compile, link to .metallib).

The few real differences:

Build pipeline. CUDA: one nvcc command emits an executable. Metal: metal compiles to .air, metallib links to .metallib, swiftc builds the host driver, and the .metallib is loaded at runtime.
printf vs buffer. CUDA has GPU printf. Metal forces a buffer round-trip. The buffer pattern is more production-shaped anyway.
Tensor units. Both have them — Tensor Cores on NVIDIA (mma.sync), matrix coprocessor on Apple. Both are wrapped by vendor BLAS.
Memory model. Discrete VRAM vs unified RAM (lesson 2 of this track).

Code

Same vector add, two languages — read them side by side·cuda

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

Same vector add, Metal version·metal

// Metal
kernel void vec_add(
    const device float *A [[buffer(0)]],
    const device float *B [[buffer(1)]],
    device       float *C [[buffer(2)]],
    constant     uint  &N [[buffer(3)]],
    uint i                [[thread_position_in_grid]])
{
    if (i < N) C[i] = A[i] + B[i];
}

Same colony, different uniforms

Code

External links

Exercise

Progress

Comments 0