Naive GEMM Is Not the First CUDA Kernel You Optimize

It’s the First Kernel That Shows Why GPUs Need Memory Reuse

Almost everyone who starts learning CUDA writes the same matrix multiplication kernel first.

One thread computes one output element.

Inside the thread:

for k:
    sum += A[row][k] * B[k][col]

It feels correct. Clean. Parallel.

At first glance it even feels efficient.

There are thousands of threads running. Every thread is doing useful work. Memory accesses look structured. Nothing seems obviously wrong.

So naturally the next assumption is:

this must be the baseline version of matrix multiplication.

Later we will optimize it.

But something strange happens when you actually look at how the GPU executes this kernel.

The kernel is not slightly inefficient.

It is fundamentally structured in a way GPUs dislike.

And once that becomes visible, naive GEMM stops looking like a beginner version of a fast kernel.

It starts looking like a demonstration of what not to do with memory on a GPU.

The First Surprise: The Kernel Already Looks Parallel Enough

When people first see naive GEMM, the mapping feels perfect:

one thread → one output element

For an (N \times N) matrix, that means launching (N²) threads.

That sounds ideal.

GPUs like parallelism. We created parallelism. So performance should be good.

This logic feels completely reasonable.

But GPUs are not limited by how many threads you launch.

They are limited by how efficiently those threads share data.

And naive GEMM does not share data at all.

The Second Surprise: Memory Access Is Already Coalesced — And Still Slow

A common explanation says naive GEMM is slow because memory accesses are not coalesced.

But something interesting happens if you actually trace what a warp is doing.

Take 32 threads computing:

C[row][0]
C[row][1]
C[row][2]
...
C[row][31]

All those threads read:

B[k][0]
B[k][1]
B[k][2]
...
B[k][31]

These are consecutive memory locations.

That is already coalesced.

At the same time, those threads all read:

A[row][k]

the exact same value.

Modern GPUs broadcast this efficiently across the warp.

So both matrices are accessed in a way the hardware understands well.

And still the kernel performs badly.

So clearly, the real problem is not coalescing.

Something deeper is happening.

The Real Problem Appears When You Look One Step Further

Instead of looking inside one warp, look across nearby threads.

Imagine two threads computing:

C[0][0]
C[1][0]

Both of them repeatedly read:

B[k][0]

for every iteration of the loop.

They are loading the same number.

Again and again.

From global memory.

Nothing is shared between them.

Now extend that thought across thousands of threads.

Across multiple warps.

Across multiple thread blocks.

Across the entire grid.

The same column of matrix B keeps getting fetched repeatedly.

The same rows of matrix A keep getting fetched repeatedly.

This is where the real issue lives.

Not in how memory is accessed.

But in how often memory is accessed.

GPUs Are Fast Because Threads Cooperate — Naive GEMM Pretends They Don’t

When people first learn CUDA, it is natural to imagine threads behaving like tiny independent CPU cores.

Each thread:

loads its data does its work stores its result

That mental model works for correctness.

It fails for performance.

GPUs are designed around cooperation.

Threads are supposed to load data together.

Share that data.

Reuse it many times.

Then move on.

Naive GEMM ignores this completely.

Every thread behaves as if it is alone on the GPU.

The Kernel Looks Efficient Until You Count the Loads

Inside the inner loop of naive GEMM, every iteration performs:

load A[row][k]
load B[k][col]
multiply
accumulate

Then repeat.

Again.

And again.

And again.

Now imagine this happening thousands of times per thread.

And thousands of threads are doing the same thing.

The GPU is not spending most of its time multiplying numbers.

It is spending most of its time fetching the same numbers repeatedly.

That is the real bottleneck.

The Next Thing the Kernel Quietly Reveals: What the GPU Is Actually Executing

At the CUDA level, the inner loop looks simple:

sum += A[row][k] * B[k][col]

But underneath that line, the GPU performs something closer to:

ld.global
ld.global
fma

for every iteration of k.

Two memory loads.

One arithmetic instruction.

Repeated thousands of times.

This ratio matters much more than it appears at first.

Because global memory loads are expensive compared to fused multiply-add instructions, computation keeps waiting for data instead of continuing smoothly.

Even though matrix multiplication is mathematically compute-heavy, this implementation behaves like a memory-driven workload.

This Is Where Arithmetic Intensity Explains Everything

A useful way to describe what is happening here is arithmetic intensity.

Instead of asking:

How many operations does the kernel perform?

arithmetic intensity asks:

How much computation happens per byte loaded from memory?

In naive GEMM, every multiply-accumulate operation depends on two values arriving from global memory.

Once those values are used, they are not reused by nearby threads.

So computation keeps requesting new data instead of continuing with data already available.

This produces extremely low arithmetic intensity.

And low arithmetic intensity means the kernel becomes memory-bound.

The Bandwidth Problem Is Not What It First Looks Like

At this stage, it is natural to assume the GPU simply does not have enough memory bandwidth.

But something surprising happens when the numbers are considered more carefully.

If each element of matrices A and B were loaded only once, the total memory traffic required would be relatively small compared to what modern GPUs can support.

Instead, naive GEMM reloads the same rows and columns repeatedly across different threads.

So the limitation is not that bandwidth is unavailable.

The limitation is that bandwidth is being used inefficiently.

The kernel is not waiting because the memory is slow.

It is waiting because memory is being asked to do the same work again and again.

This Is Usually the Moment Experienced Kernel Engineers Recognize the Pattern

When someone sees naive GEMM for the first time, indexing logic stands out immediately:

row = blockIdx.y * blockDim.y + threadIdx.y
col = blockIdx.x * blockDim.x + threadIdx.x

The kernel looks structured and correct.

But experienced kernel engineers usually ask a different question:

How many global memory loads happen before the next multiply-accumulate instruction?

In naive GEMM, the answer is always the same:

two loads one computation

Repeated thousands of times.

That ratio alone already predicts the performance behavior of the kernel.

What New Kernel Writers Usually Notice First

When encountering naive GEMM early in CUDA learning, attention naturally goes to things that are easy to verify.

The kernel launches the correct number of threads.

Each thread computes exactly one output element.

Accesses across the warp appear contiguous when reading matrix B.

Accesses across the warp share the same element when reading matrix A.

Nothing looks obviously inefficient.

So the conclusion often becomes:

This kernel is correct and just needs optimization later.

This interpretation makes sense from a correctness perspective, but it hides the real performance limitation.

What Experienced Kernel Engineers Usually Notice First Instead

An experienced kernel engineer often reads the same inner loop and mentally translates it into something closer to this:

load from global memory
load from global memory
compute one fused multiply-add

and then immediately asks:

How many times will this pattern repeat?

The answer is:

thousands of times per thread

and across thousands of threads

for the same matrix elements.

That observation alone is enough to predict that the kernel will be limited by memory traffic long before profiling tools are even used.

The difference is subtle but important. The focus shifts away from thread count and indexing correctness toward the relationship between memory movement and computation.

A Small Example That Shows the Difference in Perspective

Consider two nearby threads computing:

C[0][0]
C[1][0

A new kernel writer usually sees this as two independent computations happening in parallel.

An experienced kernel engineer sees something else happening at the same time.

Both threads repeatedly load:

B[k][0]

for every iteration of the loop over k.

So even before thinking about optimization strategies, the kernel already reveals that multiple threads are fetching identical data independently instead of sharing it.

That is not just a small inefficiency.

It is a signal that the memory strategy itself needs to change.

Why This Difference in Perspective Matters So Much

At this stage, the kernel still looks mathematically correct and structurally clean. Nothing appears broken. But experienced engineers recognize that performance problems on GPUs rarely come from incorrect indexing.

They usually come from incorrect assumptions about reuse.

Instead of asking:

Are accesses coalesced?

The question becomes:

How many multiply–accumulate steps happen after each global memory load?

In naive GEMM, the answer is always the same.

One.

That single observation explains why the kernel cannot scale toward peak performance.

Why Tiling Stops Looking Like an Optimization

At first, shared-memory tiling is often introduced as an optimization step applied after naive GEMM.

But once the memory-to-computation ratio becomes visible, something else becomes clear.

Tiling does not make naive GEMM slightly better.

It changes what the kernel is doing with memory entirely.

Instead of each thread loading values independently:

threads load tiles together

reuse them across multiple multiply-accumulate steps

and keep the computation running continuously before requesting new data

So tiling is better understood as a reconstruction of the kernel’s memory strategy rather than a small optimization applied afterward.

The Real Lesson Hidden Inside Naive GEMM

Naive GEMM is often introduced as the first step toward high-performance CUDA matrix multiplication.

In practice, it serves a different role.

It demonstrates that parallel execution alone is not enough to achieve performance on a GPU.

The kernel is mathematically correct.

Its indexing is clean.

Its accesses are aligned with hardware expectations.

Yet it performs poorly because it ignores the most important rule of GPU kernel design:

Data should be reused after it is loaded.

Once this becomes visible, matrix multiplication stops looking like a nested loop mapped across threads.

It starts looking like a data-movement problem.

And understanding that shift is exactly where high-performance CUDA programming begins.