GPU Lab EXTRA PRACTICE / AFTER CHAPTER 5
Reading a tensor
Follow a warp’s addresses from tensor elements to memory sectors.
Before you start: Complete chapter 5: How do threads read those values? first, then go straight to the memory access controls. The explanations below recap storage and strides if you need them.
START WITH A PYTORCH TENSOR
Where does a tensor keep its numbers?
PyTorch calls the memory holding a tensor’s numbers its storage. A tensor on the CPU normally uses system RAM; a CUDA tensor (on an NVIDIA GPU) normally uses GPU memory, also called VRAM. Here, storage means memory, not a saved file on an SSD or disk.
We often draw a tensor as a grid. In a contiguous, row-major tensor, its values occupy neighbouring positions in memory: all the values of the first row, then the second row, and so on. Each number in this example is FP32 (32-bit floating point), occupying four bytes.
One tensor: a grid and its memory layout
Drawn as 2 rows × 4 columns
The same 8 values in memory order
First 4 positions: row 0 · Next 4: row 1. Small labels are byte offsets from the start of storage, not RAM addresses.
The shape tells PyTorch how to interpret those values as a grid: here it is (2, 4). The strides tell it how many element positions to move when an index changes. Strides (4, 1) mean that moving down one row advances four elements (16 bytes), while moving right one column advances one element (4 bytes).
Now imagine a batch of 32 examples with eight features each. We can make a larger tensor whose values are easy to follow:
import torch
x = torch.arange(32 * 8, dtype=torch.float32).reshape(32, 8)
print(x.shape, x.stride()) # shape (32, 8), strides (8, 1)This tensor has shape (32, 8): 32 rows and eight columns. Its strides (8, 1) mean a move down one row advances 8 elements (32 bytes), and a move right one column advances 1 element (4 bytes).
With arange, the values happen to match their element offsets (positions counted from the beginning of storage). This is a convenient coincidence, not a general property of tensors. Byte offsets are those positions multiplied by four; actual memory addresses are different and are not shown.
Example: move right one column
Tensor view: x[row, column]
… 32 rows in total · 8 columns across →
Storage order (row-major)
0 → 1 · +1 element = +4 bytes
Example: move down one row
Tensor view: x[row, column]
… 32 rows in total · 8 columns across →
Storage order (row-major)
0 → 8 · +8 elements = +32 bytes
Both slices below are views: they share x’s storage. Reading the first row visits neighbouring elements. Reading the first column visits every eighth element. contiguous() copies this strided column into packed storage; that copy also takes work and memory.
Compare: x[:, 0] (view) versus x[:, 0].contiguous() (copy)
x: original storage (first 4 rows)
… 32 rows in total · 8 columns across →
Column view: x[:, 0] · stride 8 · shares x's storage
Values 0, 8, 16, 24 … point into x's existing storage at offsets 0, 8, 16, 24 … No copy is made; the slice is a view with stride 8.
Packed copy: x[:, 0].contiguous() · stride 1 · new storage
The same values 0, 8, 16, 24 … are copied into new storage at offsets 0, 1, 2, 3 … The copy has stride 1. .contiguous() on an already-contiguous tensor may return the same tensor without copying.
More settings: explore other starting positions
One index step, two views
Dashed border = start. Solid filled border = destination. The underlined storage cells lie between them.
Tensor view: 32 rows × 8 columns
Showing rows 0–3; each cell shows x[row, column].
← 8 columns across · 32 rows down ↓
The same values in storage order
A window starting at element 0. Neighbouring boxes are neighbouring stored elements.
0 → 8
+8 elements = +32 bytes
Top number: element offset and value in this arange example. Bottom: byte offset from x’s storage base.
x[0, 0] → x[1, 0]: increase the row index by one. Storage offset 0 → 8; 8 elements × 4 bytes = 32 bytes.
A view keeps the original offsets; a packed copy changes them
x: original storage
First four rows of x. Highlighted values belong to the selected slice.
x[:, 0]
View points into x’s same storage
First four reads shown; 32 values in the complete slice.
Storage is shared: no values are copied to make this view. The resulting slice has stride 8.
Full runnable example
import torch
x = torch.arange(32 * 8, dtype=torch.float32).reshape(32, 8)
row = x[0, :] # features of example 0: 0, 1, ..., 7
column = x[:, 0] # feature 0 of each example: 0, 8, ..., 248
print(x.shape, x.stride()) # shape [32, 8], strides (8, 1)
print(row.stride()) # (1,)
print(column.stride()) # (8,)
packed_column = column.contiguous()
print(packed_column.stride()) # (1,)The buttons below map one participating lane to each selected value. The row has eight values, so its experiment uses eight active lanes. The column has 32 values and uses the full warp. Both start at offset zero.
This CPU example needs no GPU. The same shape and stride reasoning applies to the corresponding CUDA tensor. PyTorch can choose its own kernels, thread mapping and copies; slicing alone does not launch the particular load modelled here.
Same values. Different addresses.
Imagine reading one feature from each row of a tensor. The GPU can give a different row to each participating thread.
New concepts
- Kernel
- A function launched to run on a GPU. Many threads execute it.
- Thread
- One execution of the kernel, with its own thread identity and work.
- Warp
- A group of 32 threads that an NVIDIA GPU schedules together.
- Lane
- A thread’s position within its warp, numbered 0 through 31.
Each square represents one lane (one thread). Eight per row is just the drawing layout, not a separate hardware group. Some lanes may be inactive for an instruction.
Each of the 32 threads in a warp can read a different row. When every lane participates, the warp requests 32 values together. Each value here is a 32-bit floating-point number (FP32) occupying four bytes.
The GPU can combine nearby memory requests. Our teaching model groups addresses into aligned 32-byte sectors: bytes 0–31, then 32–63, and so on. Combining nearby requests into fewer memory transactions is called coalescing.
Model comparison: what a warp reads
This model counts distinct aligned 32-byte sectors touched by one warp's FP32 load. It is not measured GPU timing. Each FP32 value is 4 bytes; each sector holds 8 values (32 bytes).
Row read: 32 lanes, 32 adjacent values
4 sectors · 128 bytes requested · 128 bytes covered · 100% utilisation
Column read: 32 lanes, every 8th value
32 sectors · 128 bytes requested · 1,024 bytes covered · 12.5% utilisation
The comparison above shows the key result: reading 32 adjacent values touches 4 sectors; reading every 8th value touches 32 sectors, even though both request the same 128 bytes of useful data. Use the experiment below to explore how stride, offset, and active lanes change the sector count.
Choose a pattern, then predict or show the result.
The warp
Choose a lane to follow its addressMemory sectors
Each small cell is one four-byte valueFilled cells are requested values; outlined cells share their sector. The selected lane’s value has a bright border. Rows show only touched sectors; address gaps are omitted.
What this model tells you
We count distinct aligned 32-byte sectors touched by one warp’s global-memory load, with four-byte aligned FP32 values and a sector-aligned base pointer. This matches the coalescing accounting described for devices of compute capability 6.0 or newer in NVIDIA’s CUDA Best Practices Guide.
These counts are not DRAM transactions or elapsed time. Caches, reuse between warps, stores, instruction issue, and memory-controller behaviour are outside this model. A sector diagram does not show execution order. More sectors can increase memory traffic; it does not imply a proportional slowdown.
Read NVIDIA’s coalescing guide ↗CONNECT THE MODEL TO CODE
A small CUDA kernel.
CUDA lets you write the program each GPU thread runs. This kernel copies selected input values into consecutive output positions. Change stride or offset above and the input indexing below changes with it.
Threads are also grouped into blocks. Here, threadIdx.x is a thread’s index in its block; blockIdx.x identifies the block; blockDim.x is its thread count. Together they calculate i, the output index.
The simulator shows the input load for the first warp of the first block. Its active-lane control corresponds to a bound n smaller than a warp. The output store is not counted.
TRY IT YOURSELF
Make a prediction you can explain.
- Choose Tensor column. Think of a row-major tensor with eight features per row. The warp reads the same feature from successive rows. Predict the sector count, then check it.
- Switch to Contiguous. Now it reads neighbouring features. Compare requested bytes and sector coverage. Which changed?
- Choose Shift by one. Explain what happens at the two ends of the address range.
- Set 8 active lanes, stride 1, and offset 0. Predict what changes if you move the offset to 1.
Explain the result in your own words, then compare
Each sector holds eight FP32 values. Neighbouring lanes can use different values in the same sector. With a stride of eight elements, each lane reaches a different sector. Shifting a consecutive range by one element can make it cross an extra sector boundary.
In an AI workload, changing how threads map onto tensor dimensions can change this access pattern. A transpose or a different tensor layout may help one operation and add costs elsewhere. Count the work across the whole operation before choosing a layout.
Run the example on your desktop
Download this chapter’s opening PyTorch example. See the complete example index and setup appendix for CPU PyTorch, the supported GPGPU-Sim programs and optional NVIDIA GPU execution.
The browser still requires no installation. Downloads are generated from the displayed source; CPU checks and simulator checks do not establish GPU performance.
References
Use these official references to check the concepts and code on this page.
- PyTorch: Tensor.stride (new tab)
Relate a view’s strides to the byte addresses in the diagram.
- PyTorch: tensor views (new tab)
Check which operations share storage.
- CUDA Best Practices: coalesced access (new tab)
Compare the model’s sectors with the documented access rules.
These browser exercises do not run an installed PyTorch or CUDA toolkit. For optional local execution, choose documentation matching your installed tools. Links open online in a new tab; the questions and answer comparisons below work offline.
Use the documentation
Use the references to explain each answer, then apply it to this page’s example. No installation is needed.