I wrote this while moving a JSON-extraction service (a quantized 8B model on a single T4, around ten concurrent users) to a bigger GPU on Azure. The question “which GPU, and how fast will it be?” turned out to be answerable with arithmetic you can do on paper. The numbers in the widgets all come from the model’s config, so you can check every one of them.

TL;DR: A GPU does not run “a model”. It runs thousands of tiny threads that each compute a few multiplications and additions. In this post we follow one token through Qwen3-8B, starting at its real config.json, count every weight, FLOP and byte, find out why generating text is memory-bound, compare an NVIDIA T4 with an A100, and finally see how CUDA indexing (threadIdx, blockIdx, blockDim, grids) maps the matrix-vector products onto the hardware. Everything has a widget you can play with.

I wrote this while moving a JSON-extraction service (a quantized 8B model on a single T4, around ten concurrent users) to a bigger GPU on Azure. The question “which GPU, and how fast will it be?” turned out to be answerable with arithmetic you can do on paper. The numbers in the widgets all come from the model’s config, so you can check every one of them.

All interactive widgets work in the light and the dark theme: they inherit your text color and use transparent backgrounds.

Contents

1. Where the numbers come from: config.json
2. The three ideas: FLOPs, bytes, arithmetic intensity
3. The architecture, drawn with the config values
4. One token’s journey: an animated simulation (interactive)
5. The KV cache
6. The hardware: T4, A100, H100 and what Azure offers
7. Memory-bound or compute-bound? The full calculation (interactive)
8. Prefill: why reading the prompt is a different problem
9. Quantization: what INT4, FP8 and FP16 really change
10. CUDA indexing: threads, blocks and grids (interactive)
11. Reading a pipeline timeline
12. A vLLM checklist
13. Limits of this post and sources

1. Where the numbers come from: config.json

Every model on Hugging Face ships a small file called config.json. It does not contain any weights, only the shape of the network. Everything in this post is derived from it. This is the real file for Qwen3-8B (source, fetched 2026-10-20):

{
  "architectures": ["Qwen3ForCausalLM"],
  "attention_bias": false,
  "attention_dropout": 0.0,
  "bos_token_id": 151643,
  "eos_token_id": 151645,
  "head_dim": 128,
  "hidden_act": "silu",
  "hidden_size": 4096,
  "initializer_range": 0.02,
  "intermediate_size": 12288,
  "max_position_embeddings": 40960,
  "max_window_layers": 36,
  "model_type": "qwen3",
  "num_attention_heads": 32,
  "num_hidden_layers": 36,
  "num_key_value_heads": 8,
  "rms_norm_eps": 1e-06,
  "rope_scaling": null,
  "rope_theta": 1000000,
  "sliding_window": null,
  "tie_word_embeddings": false,
  "torch_dtype": "bfloat16",
  "transformers_version": "4.51.0",
  "use_cache": true,
  "use_sliding_window": false,
  "vocab_size": 151936
}

Only ten of those keys matter for performance:

Three more keys matter for deployment:

• torch_dtype: bfloat16. The T4 has no bf16 support, so on a T4 you must run FP16 (--dtype half in vLLM) or a quantized checkpoint.
• rope_theta: 1000000 is the base for rotary position encoding. A large value helps with long contexts.
• use_cache: true is the KV cache from section 5.

Try it: compute everything from the config

Change any value (or paste the numbers from another model’s config.json) and see parameters, FLOPs per token and KV-cache size update. The defaults are Qwen3-8B.

2. The three ideas: FLOPs, bytes, arithmetic intensity

A GPU is limited by two different things, and which one hurts depends on the job.

• FLOPs (floating-point operations): one multiply or one add is one FLOP. A multiply-then-add counts as 2 FLOPs. The GPU’s calculating speed is measured in TFLOPS (10¹² FLOPs per second).
• Bytes moved: every number the GPU computes with first has to be fetched from its memory (VRAM). The fetching speed is the memory bandwidth, measured in GB/s.
• Arithmetic intensity = FLOPs ÷ bytes fetched, in FLOP/byte. It says how much math you do per trip to memory.

Think of a cook and a fridge in another building. Bandwidth is how fast you can carry ingredients, TFLOPS is how fast you can chop. If you chop only a little per trip, you wait for the walk (memory-bound). If you carry a lot and chop for ages, the knife is the limit (compute-bound).

The golden rule for matrix × vector

A matrix with m rows and k columns times a vector of length k gives m outputs. Each output is a dot product: k multiplies and k adds.

FLOPs = 2 × m × k          (2 per weight)
bytes = m × k × bytes_per_weight    (each weight is fetched once)
intensity = 2 / bytes_per_weight   → 1 FLOP/byte for FP16

One FLOP per byte is tiny. If you now multiply the same matrix by n vectors at once (n users), the weights are still fetched once but used n times, so intensity becomes n FLOP/byte. This is why batching users is so effective.

The ridge point

Every GPU has a balance point called the ridge point:

ridge = peak FLOPS ÷ peak bandwidth

For an A100 that is 312 TFLOPS ÷ 1.935 TB/s ≈ 161 FLOP/byte. Below it you are memory-bound, above it compute-bound. Hold on to the number 161. In section 7 we will see how far below it LLM decoding sits.

3. The architecture, drawn with the config values

An LLM generates text by repeating one loop: take the text so far, output a score for every possible next token, pick one, append it. Inside the model every token is a vector of 4,096 numbers that travels through a stack of identical layers.

Inside each layer there are two blocks, each wrapped in a residual connection (the block’s output is added to its input, never replaces it):

• Attention lets the current token look at earlier tokens. It needs four matrices: Wq, Wk, Wv, Wo.
• MLP (feed-forward network) processes each token on its own. It widens the vector to 12,288 numbers and narrows it again. Qwen uses the SwiGLU design with three matrices: Wgate, Wup, Wdown.
• RMSNorm layers before each block keep the numbers in a stable range. Their cost is negligible.

Counting the weights

The “8B” in the name is just the sum of all matrix entries. Using rows × columns from the config:

The few thousand norm weights per layer are negligible. In FP16 that is 16.4 GB, in INT4 roughly 4 to 5 GB (depending on what the quantizer leaves in FP16).

FLOPs per token

The golden rule says 2 FLOPs per weight that takes part in a multiplication. The embedding table is only a lookup, so it does not count; the LM head does:

weight FLOPs per token = 2 × (6,945,767,424 + 622,329,856) = 15.14 GFLOPs

Attention adds math that depends on the context length L. For each of the 32 query heads in each of the 36 layers the new token does a dot product with L cached keys (256 FLOPs each: 128 multiplies + 128 adds) and then mixes L cached values (another 256 per position):

attention FLOPs per token = 36 layers × 32 heads × 512 × L
                          = 589,824 × L

At L = 8,000 that is 4.72 GFLOPs, so one token costs about 19.9 GFLOPs in total.

4. One token’s journey: an animated simulation

Below is a miniature Qwen3. It follows the same recipe as the real model, but with 8 numbers instead of 4,096 (4 query heads, 2 key/value heads, 16 numbers in the MLP). The weights are random, so the word it predicts is meaningless, but every number on screen is really computed: the dot products, the RMSNorm, the rotation, the softmax.

How to use it:

• Press Play to watch all 16 steps in a row, or click a step number to jump to it.
• Drag the slider to scrub through the animation of one step, forwards and backwards.
• Under the animation you get a plain-words explanation, the formula and the real Qwen3-8B sizes for that step (weights, FLOPs, KV-cache reads).
• Head changes which attention head you look at in steps 7 to 9. Context changes how long the real context is for the real-size numbers.
• Colors: every box is one number. Orange is positive, blue is negative, a stronger color means a larger magnitude.

What to look for

1. Steps 3 and 4. Watch the row-by-row dot products. The whole network is built from this one operation, repeated billions of times.
2. Step 5. The arrow length never changes, only its direction. That is why RoPE can encode position without destroying the content.
3. Steps 7 to 9 with different heads. Change the Head selector and see that each head attends to the earlier tokens differently, even in this tiny random model.
4. Step 10 and 13. After every block the vector is old + correction. Compare the strips: most of the numbers move only a little.
5. The FLOP bar. At 8,000 tokens of context the MLP accounts for roughly 55% of all FLOPs and attention math for about 24%. Switch Context to 1,000 and attention nearly vanishes, switch it to 40,000 and it rivals the MLP.
6. Where the KV cache is read. Only in steps 7 and 9 (scores and mixing). Every other weight-heavy step reads weights that are shared by all users. That asymmetry is the story of section 7.

Glossary: the words of this section

5. The KV cache

The keys and values of earlier tokens never change. Instead of recomputing them for all L tokens at every step (which would multiply the work by L), the model stores them once and reads them back. That store is the KV cache.

bytes per token = 2 (K and V) × 36 layers × 8 kv_heads × 128 head_dim × 2 bytes (FP16)
                = 147,456 bytes ≈ 147 KB

Grouped-query attention (GQA) is why this is affordable. With 32 KV heads instead of 8, every number above would be 4× larger (4.7 GB per user at 8,000 tokens). Notice the asymmetry: the cache stores 8 heads (bytes) but the attention math runs over 32 query heads (FLOPs). Every cached byte is therefore used by 4 query heads, which gives the KV part a fixed intensity of about 4 FLOP/byte, no matter how many users you have. Remember this when we reach section 7.

6. The hardware: T4, A100, H100 and what Azure offers

Only two spec-sheet numbers matter for this analysis:

The A100 bandwidth figure matches the numbers published for Azure’s NC A100 v4 series (A100 PCIe, 80 GB, 1,935 GB/s). The T4 and H100 figures are the vendor’s published values as I know them, so check them against the current NVIDIA datasheets before you rely on them.

The ratios are what count: the A100 has 6× the bandwidth and 4.8× the compute of a T4.

What you can ask for in Azure

Practical notes:

• GPU quota is 0 by default. Request it under Subscriptions → Usage + quotas for the right VM family (1× A100 needs 24 vCPUs, 1× H100 NVL needs 40).
• Microsoft’s NC A100 v4 documentation says that only the NCads H100 v5 series receives net-new capacity, so A100 availability varies by region. Check Azure Regions by Product.
• One third-party comparison listed on-demand prices of roughly $3.20/h (A10), $4.41/h (A100), and $8.30/h (H100 NVL). These move, so use the Azure pricing calculator for your region. 1-year and 3-year reservations are discounted.
• Spot VMs can be evicted at any time, which is a poor fit for a production API.

7. Memory-bound or compute-bound? The full calculation

This is the core of the post. We take decode (one new token per user per step), n = 6 users, L = 8,000 tokens of context each, INT4 weights (≈ 5 GB) on both GPUs.

Step 1: bytes read per step

Everything the GPU must fetch to produce one token for every user:

weights (shared by all users)      = 5.0 GB          (INT4, one read)
KV cache (private per user)        = 6 × 8,000 × 147,456 B = 7.08 GB
total per step                     = 12.08 GB

Step 2: FLOPs per step

weights:    15.14 GFLOPs per user
attention:  589,824 × 8,000 = 4.72 GFLOPs per user
per user:   19.86 GFLOPs
6 users:    119.1 GFLOPs

INT4 does not reduce this: the weights are expanded back to FP16 inside the kernel, and the multiplication still happens at FP16.

Step 3: arithmetic intensity

119.1 GFLOPs ÷ 12.08 GB = 9.9 FLOP/byte

Compare with the ridge points: 203 (T4) and 161 (A100). Your workload sits at about 5% and 6% of the ridge. Both GPUs are deep in memory-bound territory: the tensor cores are idle more than 90% of the time.

Step 4: time per step

I assume real kernels reach 60% of peak bandwidth, 30% of peak FLOPS on the T4 and 50% on the A100. These percentages are assumptions, not measurements.

The A100 is 6.0× faster, exactly the bandwidth ratio (1,935 ÷ 320 = 6.05). Same bytes, same FLOPs, both memory-bound, so only bandwidth matters. The A100’s much larger compute capacity does nothing for decode here.

Step 5: will more users make it compute-bound?

Weights are shared, but every user brings a private KV cache. The intensity as a function of users is:

I(n) = n × 19.86 GFLOPs ÷ (5 GB + n × 1.18 GB)

The curve flattens at 16.8, an order of magnitude below the ridge points. At 8,000 tokens of context, decode never becomes compute-bound, however many users you add. More users still raise the aggregate throughput, but with diminishing returns, because each user also brings their own cache to read. With short contexts (a few hundred tokens) the weights dominate and the ceiling is much higher.

Try it: the roofline calculator

Move the sliders. The bar shows how many bytes per step are weights (blue) and how many are KV cache (orange). Each GPU card shows intensity, step times, tokens per second, the bottleneck, whether everything fits in VRAM, and the time to prefill one document of L tokens.

Things to try:

• 48 users: intensity creeps toward 16.8 and never reaches either ridge point.
• 500 tokens of context: weights dominate the bytes, and intensity rises.
• FP8 KV cache: the orange part of the bar halves.
• FP16 weights on the T4: it reports that nothing fits. The weights alone are 16.4 GB, which is why INT4 is mandatory on that card.

8. Prefill: why reading the prompt is a different problem

Before the first output token, the model must read the whole prompt. During this prefill phase all T prompt tokens are processed together, so each weight matrix multiplies a [T × 4,096] matrix instead of a single vector. Every weight is fetched once but used T times, so intensity is in the hundreds to thousands of FLOP/byte: prefill is compute-bound.

prefill FLOPs ≈ T × 15.14 GFLOPs + 589,824 × T² / 2
8,000 tokens  ≈ 121 TFLOPs + 19 TFLOPs ≈ 140 TFLOPs

Two consequences:

1. Quantizing the weights does not speed up prefill. The memory term is tiny and the compute term is unchanged.
2. Bursts hurt. If ten users submit documents at the same instant, their prefills queue up: ten × 7.2 s means the last user waits over a minute on a T4 before seeing the first token.

9. Quantization: what INT4, FP8 and FP16 really change

Quantization only changes bytes (and with them decode speed). It does not change FLOPs.

On the T4 (Turing): no bf16, no FP8, and the fast Marlin INT4 kernels need Ampere (compute capability 8.0) or newer. The realistic choices are FP16 or INT4 (AWQ or GPTQ). AWQ usually preserves quality a little better, GPTQ is sometimes slightly faster. Benchmark both. bitsandbytes NF4 is for fine-tuning, not for serving.

Decode speed per user on an A100 (n = 6, L = 8,000, 60% of peak bandwidth):

Notice that with INT4 weights the 7.08 GB KV cache is already more than half of the bytes you read each step. At that point the cache, not the weights, is what you should shrink next.

My recommendation for an extraction workload: on the T4 you have no choice but INT4. On the A100 start with FP16 or FP8 weights for the best accuracy (wrong digits or broken schemas are expensive in structured extraction), and move to INT4 only if you measure that you need the speed and your accuracy test set still passes. Support for a quantized KV cache differs between GPUs and vLLM versions, so check it for your setup.

10. CUDA indexing: threads, blocks and grids

We have established that decode is a long sequence of matrix × vector products, each limited by how fast the weights stream out of VRAM. Now let us look at how a GPU actually executes one of them.

The hierarchy

A CUDA kernel launch creates a grid of blocks, and each block contains threads. Inside the kernel every thread can read built-in variables that tell it who it is:

The classic recipe for a unique 1-D thread number is:

global_tid = blockIdx.x * blockDim.x + threadIdx.x

The hardware groups the threads of a block into warps of 32 threads that execute in lockstep. A block runs on one streaming multiprocessor (SM). A T4 has 40 SMs, an A100 has 108. The number of blocks that can be resident on an SM at the same time is limited by thread count (T4: 1,024 threads per SM, A100: 2,048), so a large grid is processed in “waves”.

A matrix-vector kernel

y = W · x for a matrix W with M rows and K columns. The natural mapping is: one warp per row, so the 32 lanes of a warp read 32 neighbouring elements of the same row (one coalesced load) and then add up their partial sums:

__global__ void gemv_warp_per_row(const half* W, const half* x, float* y, int M, int K) {
    int global_tid = blockIdx.x * blockDim.x + threadIdx.x;
    int row  = global_tid / 32;          // which row this warp owns
    int lane = threadIdx.x % 32;         // position inside the warp
    if (row >= M) return;                // the whole warp leaves together

    float acc = 0.f;
    for (int k = lane; k < K; k += 32)   // lanes stride across the row
        acc += __half2float(W[row * K + k]) * __half2float(x[k]);

    for (int off = 16; off > 0; off >>= 1)          // sum the 32 partial results
        acc += __shfl_down_sync(0xffffffff, acc, off);

    if (lane == 0) y[row] = acc;
}

// launch: threads = 256 means 8 rows per block
// int threads = 256;
// int blocks  = (M * 32 + threads - 1) / threads;
// gemv_warp_per_row<<<blocks, threads>>>(W, x, y, M, K);

Every weight is read exactly once and used for exactly one multiply and one add. That is intensity 1 FLOP/byte for a single vector, which is memory-bound by construction. Production kernels (Marlin, FlashInfer, cuBLAS) use vectorised loads and more tricks, but the indexing logic is the same.

Play with the indexing

Pick a strategy and a block size. Every chip is one row of a small toy matrix and shows which block and thread (or warp) owns it. Click a chip for the exact index arithmetic. The lower part scales the same launch up to the real Qwen3-8B matrices and shows block counts, waves per GPU, how many bytes a warp pulls in per load, and the lower bound on time from memory bandwidth.

What to notice:

• One thread per row with a toy matrix of 16 rows and a block of 128 threads launches 112 idle threads, and one warp’s loads are scattered over 32 separate rows.
• One warp per row needs 32× more threads, but each load is one contiguous 64-byte read. That is why this layout is the textbook choice for matrix-vector products.
• Small matrices underfill big GPUs. Wk (1,024 rows) with a 256-thread block makes only 128 blocks, which does not even give each of the A100’s 108 SMs two blocks. The Wq and Wgate matrices produce enough blocks to keep everything busy. This is one reason why kernels for small matrices split the work further across the K dimension.
• The “time at 60% of bandwidth” line is the lower bound for streaming that matrix once. Add the 36 layers and seven matrices and you get back to the 62.9 ms and 10.4 ms steps from section 7.

11. Reading a pipeline timeline

I plotted every document of a batch run as a horizontal bar and coloured the phases: beige = waiting for the LLM (queued), orange = prefill until the first token, dark green = decode. With the numbers above you can read such a chart quantitatively.

What it showed on the T4 and what it means:

1. Decode (green) dominated, around 80 to 90% of each bar. Prefill was small. At ≈ 15 t/s per user, a 100 to 450 s green bar is roughly 1,500 to 6,800 generated tokens, much more than I had assumed. Generation length, not prompt length, was the cost.
2. About six parallel slots. The first six documents started together and every later one started exactly when an earlier one finished. A client-side worker pool or a server limit (max_num_seqs) caps concurrency. On a T4 with 8,000-token contexts, the KV cache allows about 6 users anyway (the calculator in section 7 reports 6).
3. The first orange bars were long (45 to 90 s). Six prefills at the same moment compete for compute, the burst effect from section 8.
4. One document ran for 2,100 s, which at ≈ 15 t/s is about 30,000 tokens. That is a runaway generation (a repetition loop or a missing stop condition). It blocked one of six slots for the whole run and decided the length of the whole batch.

What would change on an A100 (estimates from the model above):

Fix these before buying hardware:

• Set max_tokens to what your JSON can plausibly need and add a request timeout.
• Use JSON-schema guided (structured) decoding so the model cannot loop outside the schema.
• Write compact JSON: no indentation, short keys, only the fields you need. Every token you do not generate is time saved, because decode is the expensive phase.
• For Qwen3, check that thinking mode is off (enable_thinking=False in the chat template). Reasoning tokens would easily explain multi-thousand-token outputs.
• Raise concurrency only after moving to the A100.

12. A vLLM checklist

Flag names change between versions, so confirm them with vllm serve --help.

# T4: FP16 compute, INT4 weights, bounded context
vllm serve <your-awq-checkpoint> \
  --quantization awq \
  --dtype half \
  --max-model-len 8192 \
  --max-num-seqs 6 \
  --gpu-memory-utilization 0.90

# A100: native bf16, higher concurrency, longer headroom
vllm serve Qwen/Qwen3-8B \
  --dtype bfloat16 \
  --max-model-len 8192 \
  --max-num-seqs 32 \
  --gpu-memory-utilization 0.90

• --max-model-len should match your real maximum prompt plus output. A smaller number frees VRAM for more concurrent users.
• --max-num-seqs is the server-side concurrency. Compare 6, 16 and 32 and look for the knee in aggregate throughput.
• To measure with your own prompts, use vLLM’s benchmark_serving.py at 10, 20 and 50 parallel requests and record time to first token, tokens per second per user and extraction accuracy (not only speed).
• For a quantized KV cache (--kv-cache-dtype), check what your GPU and vLLM version support before relying on it.

13. Limits of this post and sources

What is exact: every shape, weight count, FLOP count and KV-cache size is derived from the real config.json. They add up to the advertised 8B parameters.

What is a toy: the animated simulation in section 4 runs a miniature network (8 numbers per token, random weights). Every number it shows is really computed with the true formulas, but because the weights are random, the word it predicts means nothing. The “real Qwen3-8B” panels next to it use the true sizes from the config.

What is an estimate:

• The efficiency factors (60% of bandwidth, 30% and 50% of peak FLOPS) and the 5 GB size of an INT4 checkpoint are assumptions. Real kernels vary.
• The model ignores softmax, normalization, activation functions, rotary encoding, kernel-launch overhead and the cost of dequantizing INT4. Matrix-vector attention also often cannot use tensor cores efficiently. None of this changes the memory-bound verdict, because compute is only a few percent of the step time, but it does shift the absolute numbers.
• The overhead of 1.5 GB and 90% usable VRAM in the calculators are rough defaults. Your real limit depends on --gpu-memory-utilization.
• Hardware figures other than the A100 bandwidth come from vendor data as I know it, and prices from a third-party comparison. Verify them before you buy.

What to do with the numbers: use them to understand why things are slow and what to measure, then benchmark on the real hardware with your own prompts.

Sources

• Qwen3-8B config.json: huggingface.co/Qwen/Qwen3-8B
• Azure NC-family GPU sizes: learn.microsoft.com, NC family
• Azure NC A100 v4 series: learn.microsoft.com, NC A100 v4
• Hugging Face on Azure, supported hardware list: huggingface.co/docs/microsoft-azure
• Third-party Azure GPU price comparison: thundercompute.com
• CUDA programming guide (threads, blocks, grids, warps): docs.nvidia.com/cuda

Next time: what happens when we replace the single-GPU setup with tensor parallelism and why the interconnect suddenly matters.