Working the roofline for DeepSeek-V3 on Hopper

In part two, we looked at how our choices of parallelism, activation checkpointing and low precision influence how much memory we need to train DeepSeek-V3, and found a configuration that let us fit it on a 2048 H800 cluster."Found" is perhaps too strong a word: it is not a coincidence that if you keep the architecture fixed and look for a configuration that works well for it on the hardware DeepSeek trained on, you end up with the configuration DeepSeek picked. An interesting and important operating mode is codesign, where we (infra) can have input into aspects of the architecture (e.g., expert sparsity) so that it runs better. This is a discussion unto itself, which we will delay for later. However, you often have some choice in how exactly you get the model to fit, and we'd like to choose the configuration that maximizes our useful FLOPs/sec. For example, our first question was whether to use FSDP (ZeRO-3) or pipeline parallelism. It is quite painful to implement pipeline parallelism, so we would really like to decide on a strategy without having to build both systems and benchmark them on a full-size cluster. And even if we did build and run the benchmark, how would we know that we hadn't made a mistake that invalidated the comparison?

The classic tool for answering questions like this is a roofline analysis: at the end of the day, the number of FLOPs we need to run is fixed, and so the most important factor of our system's performance will be its distributed communication (much in the same way the most important factor of kernel performance is if it is compute or bandwidth bound). A roofline analysis will tell us that FSDP is a bad idea for DSv3, because it will be communications bound on InfiniBand bandwidth. If we consider speed of light (SOL), where we assume our compute and communications go as fast as physically possible (either according to the hardware spec sheets or controlled microbenchmarks), we can estimate the overall wall time of a training step with a spreadsheet.

The goals for this part are two-fold. First, we will walk through the speed of light rooflines for DeepSeek-V3 on Hopper. But second, and more importantly, I want to explain the underlying reasoning behind the approximations we will use in this analysis. For example, the 6ND formula is really useful, but it can also feel a bit like magic the first time you encounter it. I want to explain it in enough detail that you could derive it from first principles and also tell when you shouldn't use it.We are going to continue to omit MTP from this analysis for simplicity.

What's speed of light?

In physics, nothing can go faster than speed of light: it is a strict upper bound on the speed of objects in the universe. When doing roofline analysis, using speed of light quantities for your GPU's FLOP/s and your network's bandwidth tells you the best performance your hardware can potentially deliver. It may not be possible to achieve this performance, but it can roughly tell you how much you are leaving on the table.

However, there is a problem: what exactly is the speed of light for our hardware? An obvious choice is to use the TFLOP/s advertised in NVIDIA marketing materials, which are computed by taking the GPU's advertised clock frequency, number of arithmetic units and FLOP/cycle the units can take. But these numbers are essentially never achievable in practice for matrix multiplies, even under optimal microbenchmarking conditions. If you issue wgmma instructions in a loop with zeroed tensors, you can indeed get over 95% of the theoretical peak performance. But this doesn't translate into matrix multiply performance:

• There's more to matrix multiplies than just issuing wgmma in a loop. You have to load data from memory and the cache hierarchy, issue non-tensor core bookkeeping instructions and come up with a good schedule for all of this. Now, it is absolutely possible for a GEMM to be poorly optimized for sizes—you should roofline your GEMMs too—but a well optimized GEMM will still have a gap.
• Even when you are issuing wgmma in a loop, in some situations, if you are running it on non-zero data, you may run into power limits of the GPU which will reduce your clock speed. This leads to the surprising phenomenon where the performance of your GEMM depends on what kind of data is in the input tensors (better not tune on zero tensors!) It's also common to power cap GPUs to improve performance per watt.

Because the divergence between marketing and achievable matmul FLOPs can be quite large, we will prefer to use achievable FLOPs in our analysis below, e.g., by benchmarking cublas's GEMM speed for large enough inputs on the actual GPUs you plan to run your training on (e.g., in case they're power capped). This microbenchmark is not too hard to reproduce by hand, but in this analysis we're going to use the benchmarks from HuggingFace's Smol Training Playbook, assuming that the tensor core FLOPs of H100 match those of H800:

For network bandwidth, you can get pretty close to the spec sheet for InfiniBand, so we just use the full spec sheet 50 GB/s. However, DeepSeek reported themselves in their paper that they measured only 160 GB/s unidirectional bandwidth on H800's NVLink, so we will adopt that for our calculations (this won't matter much, except for mostly arguing that NVLink is still not the bottleneck, even with the nerfed H800 NVLink).

What's a roofline?

There are other good introductions to roofline analysis; you might also consider reading Making Deep Learning Go Brrrr From First Principles or All About Rooflines.

The term roofline originally comes from the Williams, Waterman and Patterson paper in the context of understanding program performance when memory bandwidth was the constraining factor. Its name comes from the classic shape of the graph you get when you plot arithmetic intensity (the number of FLOPs you perform per byte you read from or write to memory) against attainable FLOP/s.

Roofline analysis is a natural thing to do when authoring GPU kernels. The roofline tells you if a kernel is compute bound or memory bound. When a kernel is memory bound, making the compute run faster is useless, since after you finish your compute you will simply idle waiting for bytes to show up from HBM. Instead, fusing multiple kernels together helps you amortize the cost of HBM reads over more compute, helping push you into the compute bound regime. Conversely, we don't usually worry about fusing GEMMs together (modulo attention or low batch inference), because matrix multiplication naturally has a high arithmetic intensity and is compute-bound in most regimes. It's also fairly difficult to "fuse" two GEMMs together (or even usefully fuse into a GEMM prologue), because input tiles are used multiple times when computing the GEMM: to avoid repeatedly recomputing an input tile, you need to store it somewhere that all of your SMs can access, and has enough space to store everything you need. Usually HBM is the most convenient thing that actually fits. The fact that Flash Attention can keep the score matrix on-chip without going to HBM is something of a minor miracle.

Rooflines in distributed training

We can also do a roofline analysis on an entire distributed train step. However, there are two big differences. First, the HBM is no longer the bottleneck: instead, the bottleneck will usually be your cross-node bandwidthBlackwell's NVL72 spans across nodes, but it's usually useful to think of your NVL domain as a "node" in this analysis., e.g., your InfiniBand bandwidth. Second, the amount of communication a distributed training setup performs depends on your choice of parallelization—in contrast, the amount of FLOPs you need to do, on a fixed architecture, is fixed, no matter how you parallelize it! A good analogy is this: imagine you have a pizza pie representing how much work (FLOPs) you need to do. No matter how you slice the pie, reducing the individual amount of work a given node needs to do, all of the slices still add up to the original pie. The analogy works less well when you think about recompute conjuring more pizza to eat—a good thing in real life, and a bad thing in distributed training!

In both cases, the rule of thumb is that you want to be compute bound, since that usually means you are efficiently utilizing the hardware resources you have available. In distributed training, this means you want your compute to take longer than your communications, since a well-implemented trainer will manage to overlap the comms and make them effectively free.Actually, it's not free. If you don't have enough compute to cover up your comms, you probably want to change something (your architecture, or your parallelism if your architecture is fixed) to put yourself back in the good compute bound regime.There are also architectural changes you might make which make it harder to be compute-bound, especially optimizations for compute efficiency. For example, switching your GEMMs to low precision can exacerbate comms-boundedness, since it makes your GEMMs run a lot faster without necessarily reducing the amount of communications you need to do.

We've promised that we would show that ZeRO-3 makes you comms bound for the original DeepSeek-V3 training setup. Here's how we will figure this out:

1. We'll count the FLOPs a single training step takes. Dividing by the speed of light FLOP/s of our hardware gives the time we spend in compute.
2. For each parallelism, we'll count the bytes it moves per link and what bandwidth it uses (NVLink inside a node, InfiniBand between nodes). Dividing the per-link bytes by the speed of light byte/s on that network link gives us the time we spend in comms.
3. By comparing these times, we can tell if we are compute bound (compute time is larger) or comms bound (the links are accounted separately and either could be the bottleneck, although typically the bottleneck is InfiniBand).

How many FLOPs do we need to do?

The classic training FLOPs approximation is 6ND, where N is the number of activated parameters in the model and D is the number of tokens (either the total number of trainable tokens, giving the total amount of FLOPs for the training run, or simply the number of tokens in a global batch, giving the total amount of FLOPs for a single train stepIt's fairly typical to treat the global batch size as a fixed quantity that can't be changed, as it is typically set by science: beyond a critical batch size, increasing batch size yields diminishing returns on optimization efficiency.). This formula is indifferent to our choice of parallelism: we always have to run these FLOPs somewhere, one way or another, no matter how we have split up our work.

Most matrix multiplies in DSv3 are done in FP8 precision. So we can get a quick and dirty estimate of step time by simply dividing the FLOPs by our speed of light FP8 TFLOP/s:

Where does 6ND come from?

The core of the 6ND approximation is based off the idea that your non-attention matmuls take the bulk of the FLOPs in your model (in part 1, we used this idea to justify treating the matmuls specially in our architecture diagram). The number of FLOPs per token that a learned linear transformation computes scales linearly with the parameters of that matmul. We can see this by assigning each FLOP we need to do to a parameter in this diagram:

So you derive 6ND by multiplying N (the parameter countStrictly, we shouldn't include non-matmul params like the RMSNorm parameters or the router bias in the parameter count, but these parameters are proportionally very small and don't make a big difference either way.) by D (number of tokens) and then scaling it by 6 (two FLOPs per matmul times three matmuls total: one in forwards, two in backwards computing the gradient of the input and the gradient of the weight).

Why is this an appropriate approximation to make? We are making an argument about asymptotic behavior when we scale compute. However, upon close reflection, you might think it a bit strange to reason about the asymptotic behavior of DeepSeek-V3: after all, there's only one size for this model—671B! What exactly are you scaling?

In pretraining research, it is natural to think of an architecture as actually a family of architectures of different sizes, designed to be as similar to each other as possible, so that we can do science on the smaller sizes, fit scaling laws, and then extrapolate them to the expensive hero run. What happens to the FLOPs of a model as we increase compute, trying to keep the architecture otherwise the same? Given some compute budget C, here is a standard recipe for defining this family (for example, this matches this scaling ladder design in the MAI-Thinking tech report):

• Per Chinchilla, we will choose to keep tokens-per-parameter (TPP) fixed, so that our parameter count and number of tokens grow proportionally.
• We will choose to keep the aspect ratio (AR) of our model fixed, so that the depth and width of our model are proportionally related as we increase parameter count.
• We will choose to keep our compute allocation between attention and MLP fixed, so the hidden size and intermediate size (the MLP hidden size) are proportionally related as we increase the width of our network.
• We will choose to keep the MoE sparsity structure fixed: so the number of experts and how many experts we route to stay fixed even as compute grows.

When we increase depth, all of the transformer block FLOPs increase uniformly linearly with depth. However, when we increase width, while non-matmul compute increases linearly, the learned linear transformations (matmuls that are activation × weight) increase quadratically, since both the input and output spaces get bigger. This means that, in the limit, learned linear transformation FLOPs dominate everything else.

Is this model too simple?

Yes. There are three main refinements, ordered in how much they impact your overall wall time roofline.

1. The 6ND approximation undercounts attention FLOPs, since the attention core (the QKᵀ and PV matmuls) doesn't actually have any trainable parameters (it scales on sequence length). The attention is non-negligible and we need to account for it.
2. We are doing some recompute to help us fit the activations in memory. We need to add these recompute FLOPs to our FLOP count.
3. Not all matrix multiplies in DSv3 are FP8; notably attention, router and LM head are in BF16. We have a different speed of light number of BF16 FLOPs, so we should use a different divisor for these FLOPs. In DSv3, this isn't a huge difference (and thus last on our list), because most GEMM compute is in FP8, but this could be a substantial correction factor if you were analyzing a model that didn't have one dominant compute dtype.

Let's take these in order. Inside attention, there are attention activation matmuls that don't (directly) take a weight matrix. These scale more slowly than the parameter count (linearly, not quadratically, in width), but they also scale with sequence length. It's common to pretrain with some fixed sequence length, and then only scale the sequence length in some midtraining for context length. This commonly results in a situation where the 6ND approximation is good for pretraining (sequence length not too long) but bad for midtraining (long sequence length means attention is now a bottleneck). The diagram below shows us how the proportion of FLOPs changes as we scale up our compute and sequence length:The LM head (yellow) segment of this graph is worth a comment. The compute needed by the LM head scales with your vocabulary size, which is typically fixed throughout your entire scaling ladder. So on small models, the LM head is an extremely expensive and a very noticeable part of your compute, but because its compute is only linear in hidden size, its share asymptotes to nothing in the limit (like the non-matmul operations).

To actually compute attention FLOPs, the conventional formula is 4S²d (where S is the sequence length and d is the total attention width across all heads), but if you have causal masking, this halves the number of FLOPs.A causal query at position i attends to i keys, itself included, so the average is (S + 1)/2 keys rather than S/2. But that's only 0.02% more so we keep the simpler formula. We're also going to assume a FlashAttention style kernel implementation, which implies some built-in recompute in the backwards.

Next, we need to account for recompute. Our chosen recompute policy from part 2 recomputes many norms, which aren't modeled in our 6ND FLOPs approximation; the only GEMMs we recompute are the outputs of the MLA up-projections. We can count specifically the parameters involved in recompute:

This recompute is very cheap; it adds very little to our overall wall time!

Finally, let's properly account the precision each matmul occurs in. Here is the architecture diagram from part 2 again, with the matmuls color coded by their compute precision:

The majority of FLOPs are in FP8, so we'll compute this factor by subtracting out the BF16 FLOPs from our already computed FP8 FLOPs and then figuring their wall time with separate speed of light numbers (attention is skipped here since we already computed its wall time above). You can see that the BF16 compute is negligible; we only needed to adjust our net wall time by 121 ms.

Putting it all together:

Expert parallelism

How much time do the dispatch and combine communications take in the expert parallel region? This is a bit more annoying to do in the Hopper era, because our EP spans across nodes, rather than being in a single NVL72 domain as would be the case in GB200/GB300. But it is not too complicated. Let's assemble our assumptions for our roofline:

• We'll assume uniform routing: the nodes a token is routed to are a uniform pick among the nodes of the EP group (including our own).
• We will assume use of deduplication. The idea is that if multiple GPUs on the same node require the same token, we don't have to broadcast it multiple times on IB; instead, we can dispatch it once and then use NVLink to locally distribute it to the correct GPU. Additionally, DSv3 has a routing constraint that a maximum of four nodes can be routed to, so our worst case InfiniBand cost is four copies of the token's hidden vector.
• InfiniBand is full-duplex, meaning that we can simultaneously transmit data at the link rate in both directions simultaneously. Thus we will analyze the overall communication needed simply by looking at how much data any given node gets on the receive side.
• We will assume that the local NVLink is fast enough relative to InfiniBand that we can neglect its bandwidth for the purpose of this calculation. It's only about 3.2× faster on H800 (160 GB/s vs 50 GB/s), but the original DeepSeek-V3 technical report notes that "communications via IB and NVLink are fully overlapped, and each token can efficiently select an average of 3.2 experts per node without incurring additional overhead from NVLink." With four nodes for eight experts, on average there will only be two experts per node the token needs to be routed to.

In practice, we will be able to do a little better than having to send four copies of each token's hidden vector over InfiniBand, since if the expert is colocated on the same node no IB communication is needed. A nice and simple ratio is to say we need to send each token's hidden vector to 3.5 remote nodes, which is the expected number of communications you'd get if you picked exactly four of the eight nodes uniformly at random per token.This is an overestimate, since DeepSeek-V3 limits routing to at most four nodes, not exactly four. Specifically, the router picks the four nodes whose two best experts score highest, then takes the top eight experts among those nodes. This means that potentially one of the four selected nodes could end up hosting none of the active experts. The graph has an option to switch to DeepSeek-V3 routing: about one token in 22 touches only three nodes, pushing the mean down to 3.46, assuming independent, uniform affinity scores.

With this, we can work out the speed of light wall time that we will spend on EP communications:

We shouldn't really round FP8 to one byte: the DeepSeek-V3 FP8 format quantizes activations in 1×128 tiles, each with its own scaling factor, so every copy of a token's hidden vector carries 7,168 / 128 = 56 scales, 224 B at four bytes each. Accounting for this:DeepEP also has an 8 B header, the token's top-k indices and weights (32 B each), and padding to a multiple of 16 B. These I omit for brevity.

As a reminder, this was our compute wall time for a step:

We can see that at speed of light, EP communication alone uses up 84.8 % of our compute time.

ZeRO-1

In our memory analysis, we determined that we would need to shard our optimizer states (including the FP32 master parameters). When we do this, we need to do a reduce-scatter to get each GPU the summed gradients for its shard, and then after the (sharded) optimizer step, an all-gather that sends the updated BF16 parameters back to every GPU. Calculating the cost of these collectives is a little nontrivial due to the need to reason about hierarchical all-gather in the presence of EP, so we'll take a detour to explain this. We'll then reuse these numbers later when we analyze FSDP/ZeRO-3.

Accounting for hierarchical all-gather over IB and NVLink

To get accurate numbers for how long our all-gathers take, we need to reason about whether the all-gather for a parameter happens entirely over InfiniBand, or if some communications can be handled in NVLink. Suppose we are doing an all-gather over the entire cluster (DP2048). We don't actually need to send our parameters to 2048 GPUs over IB: we can use a hierarchical collective. First, each GPU all-gathers over IB with only the GPUs that share its local rank on the other 255 nodes, which gets it one eighth of the parameters. Then the eight GPUs on a node all-gather their eighths over NVLink. This nets out to an eight-fold reduction on IB comms compared to the full parameters. In the diagram below, mouse over a GPU to see the route its shard takes to arrive at our GPU.

However, the routed expert parameters are handled differently. With EP, nodes on the same NVLink domain already have different experts and don't need to all gather across NVLink (we use the NVLink to do the EP all-to-all comms). These parameters only all-gather over IB and don't use NVLink in their all-gather at all, thus we have to pay their full cost on IB.

The upshot is that we have to split the parameters in our model into two categories: the non-expert parameters (sharded over DP2048) and the expert parameters (sharded over EDP32) and charge them differently for InfiniBand usage. This split is pretty intuitive, but here it is on the architecture diagram (green boxes are the routed experts and blue boxes are everything else that holds parameters):

Let's write the spreadsheet. We will carefully denominate our values as parameters (not bytes), because the comms dtype differs depending on what we are doing: we reduce-scatter FP32 gradients, but we all-gather BF16 updated weights. First, the non-experts, which communicate with IB over all ranks modulo NVLink:

Note there's one small corrective factor above: we don't communicate all of Ψ! We don't need to do IB communications for the slice that lives on our node, so we only need to communicate (1 − 1/IBDP). It's worth noting that besides this factor, the amount of data we need to transmit in an all-gather isn't affected by how many nodes we are sharded to.

Continuing with experts, where our whole EDP slice crosses IB:An interesting observation is that the experts (after EP sharding) are actually less parameters than the non-expert parameters, but they still end up taking up more IB communications because none of them can be communicated over NVLink.

This gives us the following costs for ZeRO-1's all-gather and reduce-scatter:

If we only applied ZeRO-1 (without PP or FSDP), this would be our overall IB cost, pushing us squarely into the comms bound regime. However, we'll see that in PP (but not FSDP), the model sharding reduces the number of parameters the ZeRO-1 groups need to operate on.

ZeRO-3/FSDP

ZeRO-3, aka FSDP, requires more all-gathers (as we now need to all-gather our parameters for both forwards and backwards). The fact that EP + ZeRO-1 is already comms bound is bad news for doing FSDP, whose extra parameter all-gathers and reduce-scatters would also compete for IB bandwidth. Let's see exactly how bad.

Non-PP microbatching with gradient accumulation

Straight off the bat, there is a complication. Usually, when we say FSDP, we mean all-gathering parameters before forwards, all-gathering parameters before backwardsYou can skip resharding the parameters for the very last transformer block and then skip its backwards all-gather, but as you increase the number of layers this optimization matters proportionally less, so we don't show this optimization in the two-layer toy diagrams below. and reduce-scattering gradients to their shards, as you see in this diagram:How to read the memory graph? Each color specifies what (parameters, gradients, activations) the memory is used for. We assume that staging buffers, e.g., what the FSDP all-gather gathers into, are allocated persistently, rather than deallocated when not needed. While technically we don't need this staging buffer until the point when the all-gather is initiated, it doesn't really help with the high watermark to reduce the span except for the very last all-gather, and it's common to prefetch the next all-gather at the beginning of a transformer block, so you need it allocated at that point anyway.

However, recall from part 2 that we not only need to cut down on parameter memory, but also saved activation memory. When we used PP, this reduced the activation memory on GPUs because each GPU held fewer layers, and activation memory is stored per layer.We covered this in the previous part, but it bears repeating: although PP implies microbatching, since you must have more than one microbatch to cover up pipeline bubbles, the additional memory savings you would get from running microbatches are used up by the fact that our PP schedule will have multiple microbatches in flight at once—this is the 8.5 microbatches factor. However, FSDP doesn't reduce our activation memory (compared with ZeRO-1, it only shards parameters and gradients). We cannot fit on a Hopper without doing something about the activation memory!You can get the activations to fit without microbatching if you adopt a full recompute policy (recompute every transformer block immediately before backwards). However, this costs us a full forward's pass worth of wall time (a third more of the overall compute), as well as needing to run the EP comms again. This usually is a pretty bad trade!

Even if we're not using PP, we can still use microbatching (aka gradient accumulation) to reduce the memory usage of activations. We split our local batch into microbatches and run what is essentially a degenerate PP schedule where there is only one pipeline stage and you always run the backwards immediately after the forwards. The key that makes it work is that every microbatch's gradient can be computed individually and then summed together to compute the gradient of the full global batch (thanks to the linearity of derivatives). Here is an example of two microbatches:

Notice that the overall compute time stays the sameIn reality, it won't stay the same, because smaller batch size makes it harder for a matrix multiply to utilize all of the hardware. (computation is conserved!) but the number of collectives increases. This is the cost of microbatching: the more finely we split our local batch, the less compute there is to overlap the comms with. You might also wonder why we have to reduce-scatter over every microbatch: the problem (once again), is memory—without reduce-scattering, you are forced to hold onto the gradients of all parameters, not just your local FSDP shard! Eventually, there is not enough compute and our comms are exposed:

We can see intuitively from the diagrams above that eight microbatches is comms bound, but it is also the only microbatching setup that fits our activations in memory! In the above diagrams, I used our computed all-gather and reduce-scatter costs from the ZeRO-1 section to size the comms; let's now work it out on the spreadsheet.

Computing the cost of FSDP

In full FSDP, we do two BF16 all-gathers (one for forward, and one for backward) and one FP32 reduce-scatter on the parameters for every microbatch, replacing the single per-step collectives that the ZeRO-1 requires.

For reference, these were the other numbers we've already computed:

Our EP and FSDP traffic at one microbatch alone would make us comms bound; the fact that we have to do multiple microbatches to make activations fit in memory is added insult to injury. We can in fact plot the percentage of how much exposed time we have as we increase the number of microbatches.The graph interpolates between the actual physically possible microbatch counts. You could in fact achieve these effective microbatch sizes by tweaking your global batch size!

At the eight microbatches we need to fit, 69% of the step is exposed communication. Spoiler alert: PP will do much better (4%). To summarize: if you want to use FSDP, you are stuck in a "pincer" of two competing constraints:

1. FSDP doesn't save activation memory, so you want to decrease the microbatch size (increase the number of microbatches) to make things fit. But...
2. FSDP communications become more and more exposed the smaller your microbatch size is, so you end up with a lot of exposed comms with many microbatches.

Pipeline parallelism

Let's see now why pipeline parallelism is a better trade. Pipeline parallel only needs to move the residual stream across pipeline stages, which is relatively cheap since it scales only in hidden size. As in part 2, we analyze DualPipeV on 8 ranks, each holding two chunks of the model.DeepSeek's own run used DualPipe on 16 ranks, which also places two of the same 16 chunks on every rank, so the per-rank parameter share, the bubble and the residual-stream traffic below come out the same either way. We will easily not be IB bottlenecked:

However, we do pay a penalty: the PP bubble.

When we start a training step, later stages have to wait for microbatches to be processed by the earlier stages before they can start doing work. Similarly, at the end of the step, stages may run out of work to do while we are finishing running the backwards. It's beyond the scope of this article to do the full derivations for pipeline schedule bubbles, but a standard 1F1B schedule would have a bubble of (PP − 1)(F + B), where F and B are one stage's forward and backward time for one microbatch. Each stage does m microbatches of F + B work per step, so the bubble's share of compute is just (PP − 1)/m:Notice this depends only on the number of stages and microbatches, not on the model's size. In general, you can reduce your pipeline bubble by having more microbatches (at the cost of worse GEMM utilization at smaller microbatch size), but once you run out of local batch size you cannot increase the microbatch count any more.

This is already massively better than our FSDP exposed time. But it gets even better. If we use a DualPipeV scheduleIn the original DualPipeV article, Table 1, you can see the tradeoffs DualPipeV has compared to DeepSeek's original DualPipe schedule: at equal device count, it uses half the parameter memory, but twice the PP communication and a slightly larger bubble, (d − 1)/2 versus d/2 − 1 in common time units., we can reduce the bubble even further: (PP − 1)(F&B + B − 3W), where F, B and W are one chunk's times, W is the weight-gradient half of the backward (W ≈ F) and F&B is a forward and a backward run overlapped.

One final note! If you do PP, you still need to do your ZeRO-1 collectives, but you only need to do them for the parameters on your pipeline stage, reducing the comms by a factor of PP.

We can plot our idle time and see that our bubble decreases as we increase the number of microbatches.Why does the activation memory usage fall as we increase the number of microbatches? Increasing the number of microbatches indeed doesn't affect the number of microbatches in flight for a given schedule (always 8.5, in the DualPipeV schedule). But if we keep our local batch size fixed, it will decrease the size of our microbatches, thereby also decreasing memory.

Putting it together

We have seen that PP is clearly superior for training DSv3 on 2048 H800s. If we assume perfect overlap, we can now estimate what our train step time should be. We see that with PP, the InfiniBand terms sum to less than the compute, so at speed of light the step is compute bound, with the DualPipeV bubble as the only idle time. The margin is thin: comms are within 10% of compute, likely an artifact of infra-model codesign!

We can compare this against the self-reported pretraining time of DSv3 and see how different our SOL roofline is compared to the actual training numbers.

We are off by 3×! Where does the difference come from?

Comparing against the DeepSeek trace

DeepSeek published a profiling trace of one rank of a 64-GPU H800 profiling run of their EP64 setup. It captures a forward chunk of four MoE layers for a 4,096-token microbatch, overlapped with a backward chunk of four MoE layers for another. Routing is simulated as perfectly balanced and PP communication is left out. We can use this to localize some of the difference between the real run and SOL, by going through each kernel in the profile one by one and comparing it with what our SOL says it should be:Trace times are averaged over the chunk's four layers. The routed-expert GEMMs are additionally normalized to 32,768 routed rows (the capture delivered slightly fewer: 32,512 forward, 32,128 backward), and their weight-gradient entries are the elapsed span of four concurrent expert kernels rather than a per-kernel mean, so the totals are not raw wall times.

You can see that our speed of light uniformly overestimates our GEMM performance, which is not surprising because on real sizes it can be difficult to achieve microbenchmark performance in optimal conditions.

However, careful inspection of the profile reveals another issue which is reducing GEMM performance: our speed of light analysis assumes that overlapping compute with communications is free, but it's not; DeepSeek's DeepEP kernel requires 20 SMs to handle the communication, which implies our SOL time should grow by 18% (you can toggle this on the table). We can also see that in reality, ignoring all non-matmul compute cost us 2.7 ms of the forward's 9.3 ms, and 5.3 ms of the backward's 14.9 ms.

We should also take a look at the actual comms performance compared to SOL:

Same conventions as the kernel table, though averaging is somewhat misleading here, as the DeepEP notify handshake is quite variable, since it blocks on remote ranks. You can also see that SOL systematically overestimated our network performance, but that the proportions are correct: for example, dispatch to combine stays near 1:2 in both columns. In fact, the DeepSeek trace reveals that their compute and comms are balanced on a knife's edge: in this single-rank capture, compute and communication occupy the GPU almost equally (about 31 ms against 29 ms per layer, and one of the four layers tips the other way), so it is at best narrowly compute bound. This is unlikely to be an accident; they probably tuned their system based on real-world profiles to balance it as closely as possible!

This is not a complete explanation for the end-to-end 3× difference, but the trace accounts for most of it. The compute runs 2.2 to 2.5× its SOL time; comms are also slower by the same 2.2 to 2.5×. The profile is not a full train step and it measures a steady-state layer with balanced routing. We could imagine that real routing imbalance, the pipeline bubble, MTP, the batch-size ramp early in training, checkpointing and stragglers combine to cover the rest of the gap.

Conclusion

We have come a long way. By carefully reasoning out the mechanics of how our distributed training works, we can put together a simple set of formulas that let us model how our system will behave, even if we don't actually have a cluster of 2048 H800s. The obvious benefit of this is that we can now make predictions about our system performance (for example, we see that FSDP will be hopelessly comms bound). But to me, the more important takeaway is that the act of rooflining forces us to truly understand what it is our system is doing. For example, we took 6ND apart to understand how the characteristics of our job change as we scale compute. The roofline will never exactly predict real-world behavior, but the clarity of a simple model means that we can figure out what went wrong (either in the model, or in our implementation).

⤓ Download all the sheets as one .xlsx