amd (7)
![Brian Piercy](data:image/svg+xml;utf8,<svg xmlns="http://www.w3.org/2000/svg" viewBox="0 0 24 24" fill="%230A66C2"><path d="M20.447 20.452H16.89v-5.569c0-1.328-.027-3.037-1.852-3.037-1.853 0-2.136 1.445-2.136 2.939v5.667H9.351V9h3.414v1.561h.046c.477-.9 1.637-1.85 3.37-1.85 3.601 0 4.267 2.37 4.267 5.455v6.286zM5.337 7.433a1.982 1.982 0 1 1 0-3.964 1.982 1.982 0 0 1 0 3.964zm1.782 13.019H3.555V9h3.564v11.452zM22.225 0H1.771C.792 0 0 .774 0 1.729v20.542C0 23.227.792 24 1.771 24h20.451C23.2 24 24 23.227 24 22.271V1.729C24 .774 23.2 0 22.222 0h.003z"/></svg>)
let's talk more technical details about AMD's Helios! Anthropic just signed a $5B deal with AMD. and AMD's own version of GTC, Advancing AI, is happening today and tomorrow. the high level details are: 72 MI455X GPUs, 432GB of HBM4 each, and 2.9 exaflops of FP4. but the most interesting part about all this, in my opinion, is that the chip inside Helios is the most CUDA-shaped ISA AMD has ever built. for years AMD ran two separate GPU lines. CDNA for datacenter, RDNA for gaming, different instruction sets. MI300 and MI355 were CDNA, wave64, MFMA matrix instructions, a split register file. but the MI400 shows up in LLVM as gfx1250, and it's not in the gfx9xx CDNA family at all. it jumped to the 12.x line. this means that CDNA and RDNA just merged. every CDNA kernel ever tuned assumed wave64: 64 threads march in lockstep. gfx1250 is wave32 only. this is a re-think of how work maps onto the hardware. MFMA, the matrix instruction at the heart of every fast AMD kernel, is gone. it's WMMA now, the RDNA-style matrix op, with shapes like 16x16 K=128 for FP8. the old split between vector registers and matrix accumulators is gone too, unified into 1024 registers per wave. LDS and the L0 cache merged into one 448KB pool you partition yourself. one of the coolest parts: gfx1250 adds thread-block clusters where one workgroup writes directly into another's shared memory. that's NVIDIA's distributed shared memory. it adds a tensor data mover for bulk global-to-shared copies. that's NVIDIA's TMA. it adds barrier-arrive instructions. that's NVIDIA's mbarrier. So BASICALLY: AMD looked at the Hopper and Blackwell programming model and built the same primitives into their silicon. so if the hardware caught up, what's left? the fabric software. the scale-up network on Helios is Infinity Fabric tunneled over Broadcom's ethernet, and RCCL is still basically a fork of NVIDIA's NCCL. the cross-GPU collectives are the bottleneck. that's exactly why the ISA added cooperative 128-byte atomics and NVIDIA-compatible LL128, hardware built to make those collectives fast. let me know if you'll be at AMD advancing AI, i'd love to chat! emilio@wafer.ai
Examining RDNA 4's out-of-order memory accesses in detail, and investigating with testing
Lenovo, the firm emerging as a driving force behind AI computing, has expressed tremendous optimism about AMD's Instinct MI300X accelerator.
While there have been efforts by AMD over the years to make it easier to port codebases targeting NVIDIA's CUDA API to run atop HIP/ROCm, it still requires work on the part of developers.
AMD recently unveiled 3D V-Cache, their first 3D-stacked technology-based product. Leapfrogging contemporary 3D bonding technologies, AMD jumped directly into advanced packaging with direct bonding and an order of magnitude higher wire density.
AMD is one of the oldest designers of large scale microprocessors and has been the subject of polarizing debate among technology enthusiasts for nearly 50 years. Its...
AMD ROCm documentation