Scrying the AMD GFX1250 LLVM Tea Leaves
In just a few short days, AMD will be showing off their brand new MI400 series of Datacenter Accelerators at their Advancing AI event but before that event comes, we thought it would be fun to attempt to scry the tea leaves that are LLVM commits to see what we can ascertain about this next generation of AMD accelerator.
LLVM currently has preliminary support for two new accelerators from AMD, under the codenames GFX1250 and GFX1251.
GFX1250 is an accelerator chip oriented at the machine learning market called MI455X and will be what powers the Helios rack. While the GFX1251 is oriented at the HPC market, named the MI430X and is projected by AMD to deliver over 200 TFLOPs of native double-precision compute.
We’ll focus on the GFX1250 for now though and save the GFX1251 for another day.
We’ll compare the architecture to the consumer variant, the previous generation of CDNA accelerators, and to some extent Nvidia’s Blackwell.
A curious note here is that while we know that these accelerators are not APUs, they are listed as such in LLVM.
Similarities to RDNA4
The biggest similarity with the consumer architectures is that we now have a WGP (Workgroup Processor) that are built up from two pairs of two SIMD32 arrays, but there is also four SIMDs per CU just like there was in GCN and CDNA. This seeming discrepancy is because there is no longer any need for the software to separate the two CUs in the WGP and AMD seems to use the terms interchangeably.
two pairs of two SIMD32 arrays
This “change” is probably not as structural as it might at first seem and might not even exist in hardware, the changes in the caching hierarchy simply makes the distinction pointless for the compiler as there is now a shared vector L0 cache for the whole WGP.
GFX1250 operates only in Wave32 mode, unlike RDNA that can run in Wave64 in addition to Wave32 mode. Prior CDNA GPUs can only run in Wave64 mode and we can expect to see some oddness there from a performance engineering point of view when porting things to the new accelerators and a lot of kernels will have to be re-evaluated for the new architecture.
It also seems like GFX1250 can run 20 waves per SIMD, which is four more than RDNA4 can do.
We do not know if any of these characteristics are a sign for where AMD goes on their GPUs, but I personally guess that it is a sign of times to come and a future generation of AMD GPUs will be locked to Wave32 and the “WGP” as it currently exist goes the way of the dodo but it seems like this is not happening anytime soon.
Similarly to RDNA4 we now get “WMMA” support, unlike the older “MFMA” featured in CDNA and this is where most of the machine learning capabilities of the GFX1250 will come from and we’ll look into it more later.
The dynamic VGPR allocation feature that was heavily touted in RDNA4 is missing, which is a bit strange to me since I would imagine it is even more important for machine learning workloads that are under much higher register pressure due to working with whole tensors instead of just individual vectors or matrices. The mode-switch is still there but the instruction to allocate new registers is documented as a no-op.
A Change in (C)DNA
One of the largest changes found in GFX125X is that each wave in a SIMD can now address up to 1024 Vector General Purpose Registers (VGPRs) from the Register File. This is a major improvement over prior CDNA generations where you had a split allocation of registers with up to 256 VGPRs and 256 Accumulation VGPRs, which can only be used by the matrix units, for a total of up to 512 registers allocated to a single wave at a time.
This is an even larger improvement over the register allocation behavior of the RDNA series of GPUs, RDNA can only allocate up to 256 registers to a single Wave with no allocation split. This is a major update that might at some point be part of a future generation GPU IP but I wouldn’t bet on it coming soon. We will just have to wait to see if the feature is going to stay locked to the Instinct line or if this ability to allocate up to 1024 registers to a single wave will make an appearance in the Radeon line up as well.
The maximum addressable local memory (LDS) is now 320kB per wavefront which is twice as much as CDNA4 and dwarfs the 64kB addressable in RDNA. But the actual LDS can actually be even larger, a major change in GFX125X is the merging of the LDS and vector L0 caches into what AMD is calling the WGP Cache (WGP$). While both nVidia and Intel have had a single structure for both cache and shared memory for a long time, AMD has been the odd vendor out with a split cache and LDS. GFX1250 now has a single 448KB structure that can be broken up into different allocations depending on what the compiler or programmer requests.
merging of the LDS and vector L0
GFX1250 still supports packed fp32 operations allowing a single lane in the SIMD to process two fp32 operations at once just like CDNA2 and later, this practically means that when using packed instructions it has twice the vector width of RDNA.
It also inherits the more numerous and powerful SDMA units from CDNA.
Graphics Support
It is notable that nVidia and Intel do have some vestigial graphics hardware left on their accelerators and GFX125x differs a little from their offerings on this front.
No export instructions and therefore no rasterizer (like CDNA)
No export instructions and therefore no rasterizer (like CDNA)
No image (texture) instructions (like CDNA)
No image (texture) instructions (like CDNA)
No BVH (raytracing) instructions (like CDNA)
No BVH (raytracing) instructions (like CDNA)
No vector parameter interpolation instructions that are used to access and interpolate vertex data in pixel shaders (like CDNA)
No vector parameter interpolation instructions that are used to access and interpolate vertex data in pixel shaders (like CDNA)
No vector parameter interpolation instructions
No MTBUF (typed buffer) instructions are supported (unlike CDNA and RDNA3)
No MTBUF (typed buffer) instructions are supported (unlike CDNA and RDNA3)
No MUBUF (untyped buffer) instructions are supported (unlike CDNA and RDNA)
No MUBUF (untyped buffer) instructions are supported (unlike CDNA and RDNA)
No LDS parameter and parameter direct loads (like CDNA)
No LDS parameter and parameter direct loads (like CDNA)
We’re seeing that this GFX125x is even more of a pure compute accelerator than even prior CDNA architectures with nearly all graphics features having been removed. This probably saves a little bit of die space and removes features no one will miss, so it makes a lot of sense to me.
Tensor Operations
The headline feature of the MI455X is its AI capabilities, and the tensor units are what powers that. This seems to be a best of both worlds situation where AMD is combining the simpler programming model from RDNA4 with the higher performance and capabilities of CDNA4.
On RDNA4 WMMA operations are M = N = K = 16 (matrices are 16x16) except for the 4-bit ones that have K = 32 (arguments are 16x32 or 32x16 and accumulated into 16x16).
GFX1250 keeps M = N = 16, but K instead is the same as it was on CDNA4. K = 4 for 64-bit and 32-bit floating-point elements, K = 32 for floating-point 16-bit elements, K = 64 for fixed-point 8-bit elements and K = 128 for floating-point 8-bit elements. Interestingly the i4 operations stay at RDNA4 level at K=32 and are no longer a very popular data format in large AI inference scenarios due to the adoption of MX floating point operations.
Unlike on RDNA4 the sparse versions of WMMA seem to do less work per instruction than the dense ones by keeping K the same.
RDNA4 supports the following matrix of formats for WMMA
While CDNA4 supports the following for MFMA
And GFX1250 seems to support all the datatypes that are found on both CDNA4 and RDNA4 except for fp64 which seems to be restricted to GFX1251, including the support of OCP MX-style scaling that was previously available solely on CDNA4.
all the datatypes that are found on both CDNA4 and RDNA4
seems to be restricted to GFX1251
In addition there’s new instructions that allow you to transpose submatrices in LDS to complement the ones already in RNDA4. The fp4/fp6 versions of these have their own feature flag, that might be a sign that these will show up on consumer GPUs or it is just a coincidence.
It’s HIP to be … CUDA compatible?
Clusters
This seems to be AMD’s version of “Thread Block Clusters” that nVidia introduced in their Hopper architecture, and it seems like they allow you to schedule a cluster to a specific shader engine and seems to be a feature that will be available on future graphics hardware too.
In addition to allowing you to access the cluster information, AMD has also introduced cluster-level loads which allow you to write to the LDS of other workgroups within a cluster, similar or equivalent to multicast loads on Nvidia hardware.
access the cluster information
Author’s Note: A claim that wasn’t substantiated about DSMEM was removed here.
Barriers
GFX1250 adds support for cluster-level barriers, which is a feature that is essentially required to make use of clusters and is equivalent to the Nvidia feature of the same name.
In addition we’re getting intrinsics for barrier arrival, the equivalent of mbarrier.arrive in Nvidia PTX. AMD does seem to take this support one step further though by not just implementing barriers in hardware but also full monitors, allowing the wave to be notified if a specific cache line is evicted from the L2 cache.
It’s unclear what these exactly are intended to be useful for since Nvidia doesn’t have a direct equivalent already, but I guess GPU programmers will find a lot of good uses for these intrinsics.
In addition to the split barriers and barrier objects already added in RDNA4, it also supports 16 named barrier objects. These are likely added as a compatibility thing, since they have been around since Nvidia’s Tesla architecture from 2006.
Transcendentals
AMD is adding hardware support for tanh, which is rare in traditional geometry but very common in neural networks as an activation function and in the solution of many differential equations.
But the whole unit seems to have been improved, they added bf16 support and the latency for fp32 operations have gone down from 9 cycles in RDNA4 to 8 cycles in GFX1250.
Cooperative Atomics
With GFX1250 we now have intrinsics to do 32x4B, 16x8B, and 8x16B atomic loads and store to global memory, the funky naming is because these are actually 128B atomic loads and stores, the first number implies how many lanes of each wavefront are involved in the atomic but the result is essentially 1, 2, or 4 complete 128B atomic memory operations.
At the hardware-level they aren’t special instructions but just a sequence of regular loads or stores that the hardware knows how to handle so that they end up on the bus correctly.
The first use-case for this that will pop into certain peoples minds is to use this to accelerate RPCS3’s Cell emulation, and while someone is obviously going to try and use it for that, there’s a bunch of use-cases for this. But a goal is likely to accelerate AMD’s equivalent to Nvidia’s Collective Communications Library (NCCL), creatively named ROCm Collective Communications Library, as explored in the paper Demystifying NCCL.
LL128 support was to some extent already enabled on CDNA but the format has now been changed to be equivalent to the Nvidia one instead of splitting into two 64B halves.
For technical reasons these operations are not available over PCIe but I assume these will be usable over CXL and UALink. Nvidia has a very similar limitation that it is only available in global memory or over NVLink.
Tensor Data Mover
Nvidia’s marketing name for this feature is the “Tensor Memory Accelerator”, and while this feature was already available to some extent in earlier CDNA iterations it has now gained more features especially related to tensors.
The basic version just allows you to load data from global memory into LDS without using the VGPRS.
load data from global memory into LDS
While it isn’t completely clear to me exactly what capabilities are available to the tensor version since the types involved are opaque, the code isn’t totally obvious and the documentation isn’t completely clear. But Triton has some good documentation where the capabilities are enumerated. It is somewhat similar to Nvidia’s tcgen05.cp in capability except they transfer between global memory and LDS instead of between shared memory and tensor memory.
Explicit Prefetching of Vector Memory
GFX1250 adds explicit prefetching of vector memory, and AMD has finally added this feature that is available on many CPUs that are old enough to vote. Nvidia has supported this since Fermi so it really isn’t that exciting.
Some people might think that this has been a curious omission for GPUs, but unlike CPUs we’re rarely starved for registers so we can load data early and we can use the many waves running on the same core to hide the latency.
But I can imagine that inference kernels can gain some minor performance from this so it makes sense to add now I guess.
Data Dependencies
On AMD GPUs and accelerators the hardware can most of the time figure out data dependencies automatically, but in certain cases you need to do it manually.
On GCN and by extension CDNA you had three counters that allowed you to wait on operations to be completed so that you could ensure that previous operations had been completed before continuing.
VM_CNT for vector memory operations.Incremented when a vector-memory load is issued and decremented when the data is written to the VGPRs.Incremented when a vector memory store is issued and decremented when the data is written to L2 cache.Incremented when a sample or image instruction is issued and decremented when the data is returned to VGPRs.
VM_CNT for vector memory operations.
Incremented when a vector-memory load is issued and decremented when the data is written to the VGPRs.
Incremented when a vector-memory load is issued and decremented when the data is written to the VGPRs.
Incremented when a vector memory store is issued and decremented when the data is written to L2 cache.
Incremented when a vector memory store is issued and decremented when the data is written to L2 cache.
Incremented when a sample or image instruction is issued and decremented when the data is returned to VGPRs.
Incremented when a sample or image instruction is issued and decremented when the data is returned to VGPRs.
LGKM_CNT for LDS, GDS, Konstant (scalar memory read), and Message instructions.Incremented when a LDS or GDS load is issued and decremented when the data is written to the VGPRs.Incremented when a LDS or GDS store is issued and decremented when the data is written to the LDS or GDS.Incremented for every DWORD worth of scalar load issued and decremented again for every DWORD read from the scalar data cache (SMEM)Incremented for every message send issued and decremented when the message is actually sent.
LGKM_CNT for LDS, GDS, Konstant (scalar memory read), and Message instructions.
Incremented when a LDS or GDS load is issued and decremented when the data is written to the VGPRs.
Incremented when a LDS or GDS load is issued and decremented when the data is written to the VGPRs.
Incremented when a LDS or GDS store is issued and decremented when the data is written to the LDS or GDS.
Incremented when a LDS or GDS store is issued and decremented when the data is written to the LDS or GDS.
Incremented for every DWORD worth of scalar load issued and decremented again for every DWORD read from the scalar data cache (SMEM)
Incremented for every DWORD worth of scalar load issued and decremented again for every DWORD read from the scalar data cache (SMEM)
Incremented for every message send issued and decremented when the message is actually sent.
Incremented for every message send issued and decremented when the message is actually sent.
EXP_CNT for VGPR exports, this is a graphics feature used for passing things like vertex positions and color values between shaders. This exists in CDNA but isn’t used so we’re going to conveniently ignore this counter
EXP_CNT for VGPR exports, this is a graphics feature used for passing things like vertex positions and color values between shaders. This exists in CDNA but isn’t used so we’re going to conveniently ignore this counter
The big issue here with the GCN/CDNA setup of these counters is that instructions of different types are returned out-of-order. And in the case of scalar loads they are also returned out-of-order even within the same instruction type.
This regretfully means that you are likely to need to insert more waits more than would otherwise be required, especially for scalar reads where waiting for the counter to reach zero is the only useful value.
So AMD redesigned this part of the instruction set for RDNA4 and made it much more fine-grained. We now have a lot more of these counters.
LOADcnt counts the number of vector loads issued but not yet completed, similar to the first part of VM_CNT.
LOADcnt counts the number of vector loads issued but not yet completed, similar to the first part of VM_CNT.
STOREcnt counts the number of vector stores issued but not yet completed, similar to the second part of VM_CNT.
STOREcnt counts the number of vector stores issued but not yet completed, similar to the second part of VM_CNT.
SAMPLEcnt counts vector sample/gather/msaa/lod instructions issued but not yet completed, similar to the third part of VM_CNT.
SAMPLEcnt counts vector sample/gather/msaa/lod instructions issued but not yet completed, similar to the third part of VM_CNT.
DScnt counts the number of LDS instructions issued but not yet completed, equivalent to the first half of LGKM_CNT.
DScnt counts the number of LDS instructions issued but not yet completed, equivalent to the first half of LGKM_CNT.
KMcnt counts the number of scalar loads and message sends issued but not yet completed, similar to the second half of LGKM_CNT.
KMcnt counts the number of scalar loads and message sends issued but not yet completed, similar to the second half of LGKM_CNT.
EXPcnt counts the number of export or parameter loads outstanding, similar to EXP_CNT.
EXPcnt counts the number of export or parameter loads outstanding, similar to EXP_CNT.
BVHcnt counts the number of vector memory BVH (raytracing) instructions issued but haven’t returned data yet. There’s obviously no GCN or CDNA equivalent since there’s raytracing hardware.
BVHcnt counts the number of vector memory BVH (raytracing) instructions issued but haven’t returned data yet. There’s obviously no GCN or CDNA equivalent since there’s raytracing hardware.
This means that RDNA4 has a lot more situations where it can avoid having to explicitly stall the core even though there is useful work to perform compared to GCN or CDNA prior to GFX1250.
And while this isn’t as big of a deal in compute tasks, in modern graphics tasks it does make a difference since for example you can issue texture reads far ahead and not have that block a regular memory load.
On GFX125x we don’t have the following graphics wait counts,
SAMPLEcnt doesn’t exist because there are no texture samplers.
SAMPLEcnt doesn’t exist because there are no texture samplers.
EXPcnt doesn’t exist because there is no rasterizer.
EXPcnt doesn’t exist because there is no rasterizer.
BVHcnt doesn’t exist because there are no raytracing instructions.
BVHcnt doesn’t exist because there are no raytracing instructions.
But to compensate we get new ones that do not exist in RDNA4 at all,
ASYNCcnt, counts the number of outstanding asynchronous load and store instructions to LDS.
ASYNCcnt, counts the number of outstanding asynchronous load and store instructions to LDS.
TENSORcnt, counts the number of outstanding tensor load and store instructions.
TENSORcnt, counts the number of outstanding tensor load and store instructions.
Xcnt, counts the number of memory operations that have not completed address translation yet and a wait here is required to be inserted before instructions that must not be retried in case of a page fault or similar.
Xcnt, counts the number of memory operations that have not completed address translation yet and a wait here is required to be inserted before instructions that must not be retried in case of a page fault or similar.
The final one is either the most or the least interesting one to you, since it allows a simpler way to guarantee atomicity during XNACK replay.
Summary
It seems like a good improvement and AMD is catching up on features with both Hopper and Blackwell, hopefully that means that they can not only beat Nvidia in theoretical performance but also challenge the market leader in practical performance with MI455X and Helios.