Shape is a lie your framework tells you

Lesson 3 · strides and layout · the arithmetic underneath every tl.load

Lesson 2 proved your kernels are usually memory-bound. So the next question is the only one that matters: which bytes, in what order? That is decided entirely by strides — and once you see them, a whole class of mysterious 5× slowdowns stops being mysterious.

The definition that replaces the mystery

A tensor is four things, and only one of them is data:

Part	What it is	Your analogy
data pointer	One flat, 1-D block of memory.	A `malloc`'d buffer. That's it.
shape	How you'd like to index it, e.g. `(4, 6)`.	Pure metadata. Costs nothing, moves nothing.
strides	How far to step, in elements, per index.	The multipliers in your address arithmetic.
dtype	How many bytes one element takes.	The scale factor from elements to bytes.

There is no 2-D memory. There never was. There is a line of bytes and a formula:

The address formula: a logical index becomes a linear offset via the strides.

Listen first: the offset of logical position i, j equals i times stride zero plus j times stride one. For a normal row-major tensor of shape four by six, stride zero is six and stride one is one — stepping down a row jumps six elements, stepping across a column jumps one. So element two, three lives at two times six plus three, which is fifteen. That is the whole idea: shape tells you what is legal to ask for, strides tell you where it actually is.

offset(i, j) = i * stride[0] + j * stride[1]      # in ELEMENTS, not bytes

shape   = (4, 6)          # row-major, contiguous
strides = (6, 1)          # down a row = +6,  across a column = +1

offset(2, 3) = 2*6 + 3*1 = 15

Now the part that pays your rent. Transposing does not move data. It swaps the shape and swaps the strides, and returns a view onto the same buffer:

Transpose as a metadata edit: shape and strides swap, the buffer is untouched.

Listen first: after a transpose the shape becomes six by four and the strides become one and six. Not a single byte was copied. But notice what happened to the inner dimension: the stride along a row used to be one, and now it is six. Walking across a row of the transposed view means jumping six elements at a time through memory, which is exactly the access pattern that wrecks bandwidth.

x            shape (4, 6)   strides (6, 1)   contiguous
x.T          shape (6, 4)   strides (1, 6)   NOT contiguous — zero bytes copied

# the inner stride went 1 -> 6.  Reading a row is now a stride-6 walk.

Why Triton kernels take strides as arguments

In Lesson 1 the vector-add kernel computed offsets = block_start + tl.arange(0, BLOCK_SIZE) — a 1-D case where the stride is 1 and invisible. The moment you write a 2-D kernel, the stride stops being invisible and becomes a parameter you are handed:

A 2-D Triton kernel: the caller passes strides, and the kernel does the address arithmetic.

Listen first: the kernel takes a pointer plus two strides, stride_m and stride_n. It builds a row range and a column range with tl.arange, then forms a two-dimensional block of addresses: row offsets times stride_m, added to column offsets times stride_n. The colon-None and None-colon indexing is broadcasting — turning two 1-D ranges into a 2-D tile of addresses. The important point is that the kernel never assumes the data is contiguous. It is told the strides, so the same kernel works on a normal tensor and on a transposed view — at very different speeds.

@triton.jit
def kernel(x_ptr, stride_m, stride_n, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr):
    rm = tl.arange(0, BLOCK_M)                    # row offsets
    rn = tl.arange(0, BLOCK_N)                    # column offsets
    ptrs = x_ptr + rm[:, None] * stride_m + rn[None, :] * stride_n
    tile = tl.load(ptrs)                          # a 2-D tile of addresses

# The SAME kernel runs on x and on x.T — you just pass different strides.
# Same correctness. Wildly different speed.

This is the single most useful thing to internalise about Triton: you own the addresses. The compiler will happily coalesce, vectorise and swizzle — but only around the access pattern you described. It cannot fix a stride you chose.

Micro-world: watch the address arithmetic happen

Logical index → linear memory

The top grid is the shape you think in. The strip below it is the one flat buffer that actually exists. Click any cell to see the arithmetic.

Rows: 4

Columns: 6

View

Logical view — what your code indexes

Physical memory — the one buffer that exists (element offsets)

Click around the contiguous view. Offsets increase by 1 across a row and by "columns" down a column. Nothing surprising — which is exactly why the next step lands.
Switch to x.T and press "Walk one row". The logical row is still four neighbouring cells. In memory, they are six apart. Same data, same bytes, completely different access pattern.
Now press "Walk one column" on x.T. Suddenly it is contiguous again. Transposing did not make the tensor slow — it swapped which traversal is fast.
Try x[:, ::2]. A slice with no copy either: the inner stride becomes 2 and the buffer stays exactly the same size. Half the elements are simply never addressed.
Shrink columns to 2 and re-walk the transposed row. The gaps shrink with the stride. Stride is not "bad" — large stride is bad, and how large depends on your shape.

Check yourself

Cold-recall defense (one breath):
A tensor is a flat buffer + shape + strides + dtype. Memory is 1-D; shape is metadata; offset = Σ index[k] * stride[k] in elements. Transpose and slicing are free — they swap or scale strides and copy nothing — but they change the inner stride, and the inner stride is what decides whether a read is one packed burst or many scattered ones. Triton kernels take strides as arguments precisely because you own the addresses; the compiler optimises around the pattern you chose, it cannot choose a better one for you.

Primary source

Read torch.Tensor.stride and then is_contiguous (~5 min together). Then, in any Python REPL with torch: make a tensor, print .shape, .stride() and .is_contiguous(), transpose it, and print them again. Thirty seconds of that beats any amount of reading.

💬 I'm your teacher for this — ask me followups any time. Good next question: "when is .contiguous() worth the copy?" The answer is a bandwidth calculation, and you already have the tools for it from Lesson 2.