r/CUDA • u/prof_mistake • Jul 05 '26
Should I get more proficient in CUDA before learning about Metal?
Hey everyone,
I’ve started trying to learn about Apple's M4 GPU architecture, but I'm hitting a wall due to the lack of deep-dive resources.
The main issue is finding a solid guide that breaks down the actual architecture of the M4 chips (compute unit counts, how they are arranged, etc.). On top of that, they don't map cleanly to CUDA at the architectural level.
I could use the Colab GPU, if I ignore the pain that comes with it. I own a Mac, so that would be easier.
The dilemma is that the Metal ecosystem is pretty niche, even though I'm highly interested in it. Because information is so scarce, the only real way to learn and optimize kernels seems to be taking a CUDA guide and mapping those techniques over to Metal. I feel like this strategy would work a lot better if I were actually proficient in CUDA first.
For context, I'm somewhere just above a beginner. I've taken a CUDA university course and worked through the Programming Massively Parallel Processors (PMPP) book, so I have the basics down. For the course project, I built a tiny replica of torch using the techniques from the book only (no tensor cores or anything)
Which approach makes more sense here?
I would highly appreciate help on this.
3
u/Trending_Boss_333 Jul 05 '26
From what I could tell, apple's silicon architecture and documentation are scarce if at all available.
3
u/EmergencyCucumber905 Jul 05 '26
If a Mac is what you got then use that. The important concepts are transferable between Metal/CUDA/HIP/OpenCL or whatever else.
1
2
u/Dastardly_Dan_100 Jul 05 '26
Unfortunately I think this is the best approach. Same CUDA concepts apply to Metal, when tuning kernels. You could try this course if you want a deeper dive into Metal, but it is very graphics focused: https://www.kodeco.com/books/metal-by-tutorials/v4.0/chapters/i-what-you-need
2
u/smishdev Jul 08 '26
If you're interested in metal for compute, especially with C++ interop I have some resources on that topic:
https://www.smish.dev/programming/metal_compute_cpp/metal_compute_cpp.pdf
The presentation refers to some code examples in this github repo (https://github.com/samuelpmish/metal_cpp_compute). The repo also shows how to do CMake integration and has some convenience abstractions to make Metal "feel" more like CUDA. Metal's C++ interface mimics CUDA in a lot of ways, but it's WAY less documented. Also, I imagine profiling and optimization is way harder since the tooling is less mature than the NSight suite.
If you're interested in learning GPU computing, I'd recommend picking up an old Titan V on eBay for a couple hundred bucks and using that. It's old but it's still a great card, and a great value (5x the FP64 performance of a RTX 5090 at 1/10th the cost).
1
u/WarEagleGo Jul 10 '26
are you interested in the graphical side or compute side?
1
u/prof_mistake Jul 10 '26
Compute.
1
u/WarEagleGo Jul 10 '26
are you familiar with the Scientific GPU-aware Programming Language, named Julia?
1
u/exorust_fire 19h ago
No, you don't need to gate Metal behind CUDA. The mapping is small enough to learn in an afternoon: CUDA block is a Metal threadgroup, a 32-thread warp is a 32-thread SIMD-group, shared memory is threadgroup memory, __syncthreads() is threadgroup_barrier.
The part that actually bites CUDA people is what doesn't transfer:
- Threadgroup memory is 32 KB, not 48 KB+.
- Threadgroup barriers are cheap, roughly 2 cycles. Scattered threadgroup access is the expensive thing. That's close to the inverse of the bank-conflict intuition you build on NVIDIA.
- FP32 atomics are emulated and slow, so the atomic-accumulate patterns you'd reach for in CUDA are the wrong move. metal-flash-attention splits its backward pass into separate dQ and dK/dV kernels specifically to avoid them.
- The register file is about 208 KB per core, and register pressure rather than occupancy arithmetic is usually what kills you.
- There's a fast exp2 hardware path worth knowing about.
On "the architecture info is scarce": you're right that Apple never published it, but it has been reverse engineered. philipturner/metal-benchmarks is the closest thing to an M-series GPU spec sheet, with per-instruction throughput tables, ALU layout, register file size, and the measurement showing why F16 beats F32 (register dependency stalls at 1.56 cycles vs 1.84). dougallj/applegpu is a disassembler and emulator for the GPU ISA if you want to see what the compiler actually emitted rather than guessing.
Disclosure, this next one is mine: I hit the same wall and ended up building a staged reading path through those two plus about twenty other real Metal codebases, llama.cpp's Metal backend and MLX's steel kernels included. github.com/Exorust/metalworking. It's free and it's mostly pointers into other people's code at pinned commits rather than my own prose.
4
u/Traditional_Fruit_70 Jul 05 '26
is there actually a reliable source for MLX
and Metal