Skip to main content

Welcome to a Grasp lesson

Create a free account to keep studying and create your own study path.

Lesson 58

Shared Memory Tiling: Performance Benefits

Recommended resources

Introduction

In our last lesson, you successfully implemented the "hero" kernel for matrix multiplication using a tiled algorithm and shared memory. We moved from a naive approach where every thread fended for itself to a cooperative strategy where threads within a block work as a team. The central idea was to reduce the punishingly high number of accesses to slow global memory.

Today, we will prove that our hero kernel is worthy of its name. This lesson is all about analysis and quantification. We will answer two critical questions:

  1. How much faster is it? We'll measure the performance improvement in concrete terms.
  2. Why is it faster? We will go beyond intuition and perform a rigorous analysis of memory traffic and data reuse to understand the source of our speedup at a fundamental level.

This process of measuring, analyzing, and understanding performance bottlenecks is the core daily work of a kernel engineer. By mastering this, you are taking a significant step toward your goal of achieving true high-performance computing.


The "What": Measuring the Performance Gain

Before we analyze why our tiled kernel is faster, let's see the results. A common and crucial metric for measuring the performance of scientific computations is GFLOPS, or Giga-Floating-Point Operations Per Second. It tells us how many billion calculations the processor is completing each second.

For a matrix multiplication of size , each of the output elements requires multiplications and additions. This gives a total of floating-point operations. The GFLOPS can then be calculated as:

A higher GFLOPS value signifies better performance. Now, let's see a direct comparison between our naive kernel and the tiled kernel.

CUDA Crash Course: Cache Tiled Matrix Multiplication

The video "CUDA Crash Course: Cache Tiled Matrix Multiplication" by Nick provides a direct comparison of the execution times using the NVIDIA Nsight profiler. This gives a clear, practical demonstration of the speedup.

Watch the section from this timestamp. The presenter profiles a naive kernel, noting an execution time of about 18,000 microseconds (18ms). He then profiles the tiled kernel, which finishes in just 8,000 microseconds (8ms). This is more than a 2x speedup on his hardware for a 1024x1024 matrix.

This is a significant improvement, but academic and industrial research shows we can expect even more. A formal study provides a more systematic analysis.

[PDF] Performance Analysis of CUDA-based General Matrix Multiplication ...

This research paper provides rigorous benchmarking results, comparing a baseline kernel (similar to our naive one) with a shared memory tiled version.

First, look at Table 4.2 in the Results section. It shows the GFLOPS for the tiled kernel, reaching over 970 GFLOPS. Next, examine Figure 4.1 in the Comparative Analysis. This graph visually contrasts the flat, low performance of the baseline kernel (~230 GFLOPS) with the dramatically higher performance of the tiled kernel. Finally, read the first two paragraphs of the Discussion section. The authors attribute the more than 4x speedup directly to "enhanced data locality," which is precisely what we aimed for with our tiling strategy.

The data is clear: tiling provides a massive performance boost. Now, let's dissect the fundamental reason for this improvement.


The "Why": Arithmetic Intensity and Data Reuse

The performance of any kernel is ultimately governed by the balance between computation and memory access. We can formalize this relationship with the concept of Arithmetic Intensity.

Arithmetic Intensity is the ratio of floating-point operations performed to the bytes of data transferred from main memory.

A kernel with low arithmetic intensity is memory-bound; its speed is limited by how fast it can fetch data. A kernel with high arithmetic intensity is compute-bound; its speed is limited by how fast it can perform calculations. Our goal is always to move kernels from being memory-bound to compute-bound, as this means we are fully utilizing the GPU's powerful processing cores.

Let's analyze the arithmetic intensity for our two kernels.

Tiled GEMM |

This article from vvnasantosh.net provides an excellent, quantitative breakdown of the data reuse and arithmetic intensity for both naive and tiled GEMM.

Read the section titled "Data Reuse Analysis" carefully. It compares a naive implementation with a tiled one for a 256x256 matrix multiplication using 32x32 tiles. Focus on: Naive GEMM Limitations: The text explains how calculating adjacent output elements like C(0,0) and C(0,1) requires redundant loads of the same data, because the hardware cache cannot hold it long enough. Memory Access Calculation (Naive): The author calculates that for a single thread block, the naive approach has an arithmetic intensity of just 0.25 FLOPs/Byte. Tiled GEMM Solution: The author then calculates the same for the tiled kernel. Notice how each byte loaded from global memory is reused many times within the tile. This leads to an arithmetic intensity of 8.0 FLOPs/Byte.

The analysis shows a 32x increase in arithmetic intensity! This is the core reason for our performance gain. By loading data into shared memory once and reusing it many times, we have drastically reduced the demand on the slow global memory bus and allowed the SMs to spend their time on computation rather than waiting for data.

This tiered data movement is a central pattern in high-performance computing.

This diagram illustrates the full data flow in our tiled algorithm. Large chunks of the matrices are in slow Global Memory. A thread block cooperatively loads smaller tiles into fast Shared Memory. From there, warps load data from shared memory into the even faster Register File for computation by the SM CUDA Cores. Our optimization focuses on the first step: minimizing traffic from global to shared memory.

The load -> synchronize -> compute -> synchronize loop you implemented in the previous lesson directly maps to this flow. The following diagram shows how this plays out over time within the main loop of your kernel.

This diagram shows the pipeline of work within one iteration of the tile loop. First, threads cooperate to load a tile from global to shared memory. A `__syncthreads()` barrier ensures this is complete. Then, the block performs many computations, loading data from shared memory into registers for the math operations. The loop-carried dependency back to the start is protected by the second `__syncthreads()` call.

Conclusion

In this lesson, you have moved from being a kernel implementer to a kernel analyst. You not only saw the significant performance improvement from tiling but also delved into the fundamental reasons behind it.

Key Takeaways:

  • Performance is measurable: GFLOPS is a standard metric to quantify the computational throughput of a kernel.
  • Tiling yields massive speedups: We observed a greater than 2-4x performance increase, transforming our kernel from a toy example into a genuinely high-performance implementation.
  • Arithmetic Intensity is key: This ratio of FLOPs to Bytes accessed from global memory determines whether a kernel is memory-bound or compute-bound.
  • Tiling increases Arithmetic Intensity: By facilitating data reuse in fast shared memory, tiling drastically increases the number of computations performed for each byte loaded from slow global memory. This is the fundamental source of the performance gain.

You have successfully optimized the data path from global memory to shared memory. However, the data path from shared memory to the registers and CUDA cores in the diagram above still presents an opportunity.

In our next lesson, we will introduce register-level tiling. We will look inside a single thread's computation and see how we can restructure the inner loop to maximize data reuse from the register file—the fastest memory available on the GPU—to push our kernel's performance even closer to the hardware's theoretical peak.