MIT 6.5940 L11 TinyEngine and Parallel Computing: From Loop Tiling to In-Place Depthwise

🌏 中文版

Version note: This post is based on Lecture 11 (2024-10-10) of MIT 6.5940 Fall 2024. The main materials are Lec11-TinyEngine.pdf (79 pages) and the lecture recording. Page numbers refer to PDF pages. Facts were checked against the official materials on 2026-09-30. Access level A3: slides, video, and the example code repos the slides cite are all public. What you can't get is Canvas submission and grading feedback.

Fall 2026 comparison: The F26 schedule puts the same lecture on October 20. As of 2026-09-30 its slide and video links are still empty.

Series: previous L10 MCUNet and tinyML | next L12 Transformer and LLM (bridge) | Series overview

Lecture 10 described MCUNet as a co-design of TinyNAS and TinyEngine, but only covered TinyNAS. Lecture 11 fills in the other half. Once the model is fixed, how does the system layer make it run fast and use less memory?

This lecture has more code than the earlier ones, but the structure is simple. The first half takes one matrix multiply and applies loop optimizations, SIMD, multithreading, and CUDA in turn, reporting a speedup at each step. The second half switches to convolution and covers four inference tricks TinyEngine uses. The Lecture Plan on page 2 has exactly these three parts: edge AI and MCU characteristics, parallel computing techniques, and inference optimizations.

Where an MCU is small

Page 5 lines up four platforms in a table. The two ends make the point:

Page 6 compares the STM32F746 with an Apple M1 Ultra MacBook Pro. The clock is 216MHz vs 3200MHz. The MCU has only an 8KB L1 cache, with no L2, no L3, and no operating system. Memory differs by 210,000x and storage by 8,400,000x.

Page 7 shows the memory hierarchy: lower levels are larger, slower, and cheaper. Citing "Latency Numbers Every Programmer Should Know," the slide lists about 0.5ns for L1, about 100ns for DRAM, and about 1ms for storage. Every loop optimization that follows aims to keep data in the upper levels.

Parallel computing techniques: six versions of one matmul

Page 9 lists four families: loop optimization (reordering, tiling, unrolling), SIMD, multithreading, and CUDA. Example code lives in mit-han-lab/parallel-computing-tutorial. Apart from CUDA, all speedups were measured on an Intel Xeon 4114.

Loop reordering: 12x

Pages 10–12. In the triple loop i, j, k, as the innermost k changes, B[k][j] jumps down a column. Matrices are stored row-major, so those jumps keep missing the cache.

# Before: i, j, k — B[k][j] is read down a column, poor locality
for i in range(N):
    for j in range(N):
        for k in range(N):
            C[i][j] += A[i][k] * B[k][j]

# After: i, k, j — the innermost j reads B[k][j] along a row
for i in range(N):
    for k in range(N):
        for j in range(N):
            C[i][j] += A[i][k] * B[k][j]

Just changing the loop order gives a 12x speedup, per page 12.

Loop tiling: 19x

Pages 14–20 handle the next problem. When B is much larger than the cache, data gets evicted before it's reused. The fix is to split the iteration space into TILE_SIZE blocks so each block's data fits in cache. After tiling j, k, and i in turn, each access to A, B, and C touches TILE_SIZE² elements instead of N². Page 19 goes one level further with multilevel tiling for L1 and L2. Page 20's C implementation with BLK_SIZE 32 is 19x faster.

Loop unrolling: 2.85x

Pages 22–24. Loops have overhead: pointer arithmetic, the end-of-loop test on every iteration, and branch prediction. Copy the loop body 4 times and step by 4 instead of 1. Pointer arithmetic and loop tests drop to a quarter, but the innermost code gets 4x larger. The slide calls this a trade-off between binary size and overhead, and on an MCU with 1MB of Flash that trade-off is real. Page 24 unrolls j by 8 and k by 4 for a 2.85x speedup.

SIMD: one instruction, four numbers

Pages 26–27 cover instruction set background first. CISC (x86) has many specialized instructions. RISC (Arm, RISC-V) implements only the common ones. For C = A + B, CISC might need one instruction, while RISC might need four: load, load, add, store.

Pages 28–30 use 128-bit vector registers to process four 32-bit floats at once, cutting arithmetic instructions from N³ to N³/4. The slides show two intrinsic families side by side: x86 SSE (_mm_load_ps, _mm_mul_ps, _mm_add_ps) and Arm NEON (vld1q_f32, vmulq_f32, vaddq_f32). They also decode the names: ps means packed single-precision, and q means quadword. Page 30's implementation transposes B first so both operands are read contiguously. That page reports no speedup figure.

Multithreading: 4.1x

Pages 32–37. Threads in one process share memory but each has its own stack and program counter, so they can run on different cores. Page 35 uses Pthreads to split the matrix rows evenly across 4 threads, for a 4.1x speedup. Pages 36–37 switch to OpenMP, where one #pragma omp parallel for line before the loop is enough; the slide notes this is much cleaner than Pthreads.

CUDA and Tensor Cores: 94x

The CUDA introduction on pages 39–44 borrows from Stanford CS149. Threads form blocks and blocks form a grid. Each thread uses blockIdx and threadIdx to find the element it owns. Host and device have separate address spaces, and cudaMemcpy moves data between them. Kernels see three kinds of memory (per-thread private, per-block shared, and global), each with different locality and access cost. Page 44's implementation loads tiles of A and B into shared memory first. On a 2080Ti it runs 94x faster end to end than the naive CPU version.

Pages 45–52 go one level lower to Tensor Cores. A CUDA core does 1 FP32 or 2 FP16 MACs per cycle. A Tensor Core finishes a whole small matrix multiply per cycle (4×4×4 on Turing, 8×4×8 on Ampere). Page 46 measures on an A6000 that Tensor Cores are 3.8x faster than CUDA cores once N is large. Pages 47–52 show how to build a 16×16×32 multiply from 16×8×16 MMA intrinsics.

Inference optimizations: TinyEngine's four tricks

Page 54 lists the second half's four techniques together. The first two are about memory, the last two about speed.

Im2col: turn convolution into matrix multiply

Pages 55–57. Rearrange the input activation into a matrix and the convolution can call a general matrix multiply (GEMM) directly. The upside is that every matmul optimization from the first half applies. The downside is extra memory. The slides mention implicit GEMM as the fix: a variant of direct convolution that works directly on the original weight and activation tensors.

In-place depthwise: peak from 2×C×H×W to (1+C)×H×W

Page 59 gives the motivation. MobileNetV2's inverted residual block uses depthwise convolution to save model size and FLOPs. But its middle layer expands channels 6x, so peak memory grows 3–6x. That's why Lecture 10 said MobileNetV2 "shrinks parameters but not activations."

Pages 60–64 give the fix. A normal depthwise convolution holds input and output at once, for a peak of 2×C×H×W. Depthwise channels are independent, so you can compute one channel at a time with a single channel-sized (H×W) temporary buffer, then write the result back where the input was. The peak becomes (1+C)×H×W.

NHWC for pointwise, NCHW for depthwise

Pages 66–69 look at memory layout. A pointwise (1×1) convolution takes a weighted sum across channels at each pixel. NHWC, with channels innermost, makes that read contiguous, so TinyEngine uses NHWC for pointwise. A depthwise convolution slides spatially within each channel. The in-place scheme above already walks channel by channel, so NCHW is more contiguous there.

Within one network the two kinds of convolution alternate, and the layout follows the operation. General-purpose frameworks often leave this on the table; a specialized inference engine can pick it up.

Winograd: 2.25x fewer multiplications

Pages 71–76. Computing 4 outputs (2×2) of a 3×3 convolution directly takes 9×C×4 MACs. Winograd transforms the input tile and the filter into another space, multiplies element-wise, and transforms back. That takes only 16×C MACs, 2.25x fewer. The filter transform can be precomputed offline; input and output transforms happen at inference time. The full formula Y = Aᵀ[(GgGᵀ) ⊙ (BᵀdB)] comes from Lavin & Gray 2015.

The link to Labs 4 and 5

Page 77 recommends two repos, TinyEngine and TinyChatEngine. Page 78 previews Labs 4 and 5. You quantize an LLM to 4-bit with AWQ and deploy it as a local chatbot. You implement loop unrolling/reordering, SIMD, and multithreading yourself, then measure the latency gain from each. The slide calls the engine "TinyLLMEngine"; the F24 Lab 5 handout and starter repo call it TinyChatEngine. Details are in the Lab 4 + Lab 5 guide.

In other words, the CPU techniques from this lecture come back unchanged to speed up LLaMA2-7B's linear kernels.

What to do after this lecture

• Tonight: clone parallel-computing-tutorial and run the naive and reordering versions on your own machine. Compare your speedup with the slide's 12x. The gap itself tells you about cache size and memory bandwidth.
• For a fuller treatment of the GPU side: CS149 L2: Multicore, SIMD, Hardware Multithreading and CS149 Lecture 7: GPU Architecture and CUDA.
• From convolution-as-matmul and tiling all the way to FlashAttention: CS149 L9: Running DNNs Efficiently on GPUs.

Further reading

• Locality and arithmetic intensity: CS149 L6 Locality and Communication
• GPUs and Triton kernels: CS336 Lecture 5: GPUs, CS336 Lecture 6: Triton kernels
• CUDA in an assignment: CMU 11-868 HW1: CUDA Programming

References

• MIT 6.5940 Fall 2024 course page: L11 date, slide and video links
• MIT 6.5940 Fall 2026 course page: L11 scheduled for October 20, materials not yet released
• Lec11-TinyEngine.pdf (Fall 2024): source of every page number, speedup, and formula in this post
• EfficientML.ai Lecture 11 - TinyEngine (YouTube)
• mit-han-lab/parallel-computing-tutorial (GitHub): the loop optimization, SIMD, multithreading, and CUDA examples the slides cite
• mit-han-lab/tinyengine (GitHub)
• mit-han-lab/TinyChatEngine (GitHub)
• Lavin & Gray, Fast Algorithms for Convolutional Neural Networks (2015): Winograd convolution
• Lin et al., MCUNet: Tiny Deep Learning on IoT Devices (NeurIPS 2020): where TinyEngine was introduced
• Latency Numbers Every Programmer Should Know: the source page 7 cites for memory hierarchy latencies