ROCm Composable Kernel #10574 — Why __restrict__ Was an Invalid Contract for Block-Shared LDS

ROCm CK LDS Aliasing, Part 1 of 2 — A workgroup-shared scratch pointer, LLVM noalias, and numerical values that became stale across a barrier

A barrier was already present.

The kernel still produced wrong results.

That combination can make the problem look like a hardware synchronization failure.

Threads write to LDS
        ↓
block_sync_lds()
        ↓
other threads read the LDS values

If the result is stale after that sequence, the first suspicion is often:

The barrier did not work.

ROCm Composable Kernel PR #10574 found a different boundary.

The shared-memory pointer itself carried a contract that the program did not satisfy.

Several gridwise GEMM functions declared their block-shared scratch pointer as:

void* __restrict__ p_shared

But p_shared referred to AMD’s workgroup-shared Local Data Share, or LDS.

The same physical memory was intentionally accessed by multiple threads.

Thread A
→ writes LDS

Thread B
→ reads the same LDS after synchronization

In the affected compiler path, Clang lowered __restrict__ into LLVM’s noalias contract.

That contract allowed alias analysis to assume that conflicting accesses from other threads could not modify the object during the function invocation.

The source was therefore saying two different things at once.

Program structure:

Other threads write this shared object.

Pointer contract:

Other accesses cannot conflict with this object
during the invocation.

The barrier described the required execution order.

__restrict__ described an access relationship in which that cross-thread modification was not supposed to exist.

When LLVM became less conservative about synchronization effects on non-escaping local memory, the contradiction became observable.

On MI300X, the PR author reported:

Current staging compiler
+
__restrict__ present
→ FAILED
→ max error 130
→ 137,150 incorrect values
→ 22.78% of values wrong

Current staging compiler
+
__restrict__ removed
→ PASSED

Baseline compiler with the relevant alias-analysis change reverted
+
__restrict__ present
→ PASSED

The patch did not add another barrier.

It did not strengthen the existing barrier.

It removed the false aliasing promise from four p_shared parameters.

The synchronization operation was not missing.
The compiler had been given an invalid description of the memory being synchronized.

What `p_shared` represented

The affected functions were gridwise GEMM implementations inside Composable Kernel.

Their signatures included ordinary global-memory operands such as:

A tensor pointer

B tensor pointer

output tensor pointer

and one scratch pointer:

void* p_shared

That scratch pointer referred to LDS allocated for the workgroup.

On AMD GPUs, LDS is memory shared by threads in the same workgroup.

It is commonly used to stage:

input tiles,
transformed values,
partial results,
layout-converted fragments,
and intermediate data reused by neighboring threads.

The broad pattern is:

Global-memory tile
        ↓
threads cooperatively write LDS
        ↓
workgroup barrier
        ↓
threads cooperatively read LDS
        ↓
matrix instructions or further transformations

The memory is useful precisely because its ownership is collective.

It is not private storage associated with one thread.

The pointer value may be identical in every thread

A simplified kernel can be imagined like this:

__device__ void run(void* p_shared) {
  auto* scratch = static_cast<float*>(p_shared);

  if (thread_id == 0) {
    scratch[0] = 42.0f;
  }

  block_sync_lds();

  if (thread_id == 1) {
    consume(scratch[0]);
  }
}

The intended result is:

Thread 0 writes 42

barrier completes

Thread 1 reads 42

Every participating thread can receive the same LDS base address.

Thread 0 p_shared
→ address X

Thread 1 p_shared
→ address X

Thread 2 p_shared
→ address X

That is not an accidental alias.

It is the design.

The pointer names the workgroup’s shared scratch region.

`restrict` is not merely a performance suggestion

It is tempting to read __restrict__ as:

Please optimize this pointer aggressively.

Its meaning is stronger.

It gives the compiler information about how the pointed-to memory may be accessed.

In an ordinary single-threaded example:

void update(
    float* __restrict__ dst,
    const float* __restrict__ src) {
  ...
}

the qualifier tells the compiler that accesses based on dst do not conflict with accesses through another unrelated pointer such as src.

That promise can enable:

load reuse,
store motion,
vectorization,
dead-load removal,
and stronger alias-analysis conclusions.

The compiler is allowed to transform the program according to that promise.

If the promise is false, the transformed program is not required to preserve the behavior the programmer expected.

The contract also matters across threads

The important detail in #10574 was not only that several pointer expressions could reach the same address inside one thread.

LLVM’s noalias contract also applies to conflicting accesses from other threads.

The clarified LLVM rule states that conflicting accesses not based on the noalias pointer must either:

happen before the function invocation begins

or:

happen after that invocation has finished

An access performed concurrently by another thread during the invocation does not satisfy that boundary merely because the two threads meet at a barrier inside the function.

For a workgroup-shared object:

Thread A function invocation
→ writes p_shared

Thread B function invocation
→ reads or writes the same object

both invocations overlap

The cross-thread access is part of the kernel’s normal execution.

That is incompatible with treating each invocation’s pointer as an exclusive noalias access root.

The source contained a false ownership claim

The affected declaration looked approximately like this:

__device__ static void Run(
    const ADataType* __restrict__ p_a_grid,
    const BDataType* __restrict__ p_b_grid,
    EDataType* __restrict__ p_e_grid,
    void* __restrict__ p_shared,
    ...
);

The global tensor pointers and the LDS pointer do not necessarily have the same ownership contract.

Global input and output pointers

A kernel can legitimately require that:

A

B

output

do not overlap in prohibited ways.

The caller can provide separate allocations.

The __restrict__ declaration may describe a valid kernel precondition for those arguments.

LDS scratch pointer

p_shared is intentionally shared among threads.

Many thread invocations

→ same scratch object

→ conflicting writes and reads

→ synchronization inside the kernel

The same qualifier therefore does not automatically make sense for this parameter.

The problem was not that __restrict__ is always unsafe in GPU code.

The problem was that it was attached to a pointer whose purpose contradicted its exclusivity contract.

A barrier orders accesses that the compiler still believes can exist

A workgroup barrier normally establishes an execution boundary such as:

writes before barrier

happen before

reads after barrier

But that description assumes the compiler recognizes that another thread may modify the relevant memory.

Now add the false noalias fact.

From one thread’s optimized view:

I have an exclusive access contract for this object.

Other threads cannot conflict with it during my invocation.

The barrier may still synchronize control flow.

But alias analysis can conclude that the shared object itself was not modified by any access the current invocation must consider.

A previously loaded LDS value may therefore appear reusable.

load LDS value

barrier

load same address again

can be reasoned about as though it were:

load LDS value once

barrier

reuse register value

if the compiler believes no legal external write could have changed the object.

The hardware barrier cannot force a reload that the compiler removed before machine code was emitted.

A conceptual stale-value example

The following is a minimal conceptual model, not the exact CK reproducer:

__device__ void example(void* __restrict__ p_shared) {
  float* shared = static_cast<float*>(p_shared);

  float before = shared[0];

  if (thread_id == 0) {
    shared[0] = 42.0f;
  }

  block_sync_lds();

  float after = shared[0];

  consume(after);
}

The programmer expects:

before
→ old value

after
→ value written by Thread 0

But if the compiler accepts the noalias contract as true:

No conflicting other-thread modification is legal
during this invocation.

then it may reason:

shared[0] cannot have changed
through a relevant external access

therefore

after == before

The exact instruction sequence removed or retained in the failing CK kernel was not published in the PR.

The important source-level contradiction is visible without that final disassembly:

cross-thread write is required by the algorithm

while

cross-thread modification is excluded by the pointer contract

Why the failure appeared after an alias-analysis change

The source qualifier existed before the reported failure.

That does not mean it was correct before.

An invalid optimization contract can remain latent while the compiler is conservative.

False source assumption

+

optimizer does not exploit it

→ output happens to remain correct

A later compiler update can begin using the information more aggressively.

Same false source assumption

+

stronger alias analysis

→ transformation becomes legal under the declared contract

→ wrong result appears

The PR specifically linked the failure to an LLVM alias-analysis change titled:

No synchronization effects
for never-escaping identified local

LLVM traditionally models fences and stronger atomic operations as reading and writing broad memory regions so that optimizations respect synchronization.

The referenced change stopped applying those broad effects to certain identified local objects that do not escape.

That can be a valid optimization.

A genuinely private local object cannot be modified by another thread merely because a synchronization instruction exists.

The affected CK pointer, however, was not genuinely private in the way its noalias contract implied.

Optimizer improvement

+
incorrect source-level alias fact

→ latent defect becomes observable

The compiler update did not invent the sharing

The shared access already existed in the algorithm.

Thread group cooperatively uses LDS

Removing an overly conservative memory effect did not create that behavior.

It exposed a mismatch that had already been present.

This distinction matters.

Incorrect explanation

The newer LLVM compiler broke LDS synchronization.

More precise explanation

The kernel declared a block-shared LDS pointer as restrict.

A newer alias-analysis rule relied more strongly on that declaration.

The resulting optimization no longer preserved the behavior
the source author expected from cross-thread sharing.

The patch repaired the source contract rather than reverting the compiler optimization.

The before-and-after matrix isolated both sides

The PR author tested three combinations on MI300X.

Compiler	Source qualifier	Result
Current staging	`__restrict__` present	Failed
Current staging	`__restrict__` removed	Passed
Baseline with AA change reverted	`__restrict__` present	Passed

This matrix is useful because each comparison answers a different question.

Same compiler, different source

Current staging + restrict
→ fail

Current staging - restrict
→ pass

This isolates the source qualifier.

Same source, different alias-analysis behavior

Current staging + restrict
→ fail

baseline/reverted AA + restrict
→ pass

This identifies the compiler change that made the invalid assumption operationally visible.

Together:

Invalid source promise

×
optimizer willing to exploit it

→ numerical failure

The failure was numerical, not merely theoretical

The published reproducer was:

TestGroupedConvndBwdData2d/9.Test2D

The affected implementation was a grouped-convolution backward-data kernel built through gridwise GEMM code.

Before the patch, the reported output included:

max error:
130

number of errors:
137,150

incorrect values:
22.78%

After removing __restrict__:

1 test passed

The test did not merely check compilation.

It performed a numerical comparison and observed incorrect results before the patch.

That moves the finding beyond static analysis.

Invalid alias contract
→ identified in source

Wrong numerical result
→ reproduced on MI300X

Minimal source repair
→ test passes

The patch changed four declarations

The landed diff removed __restrict__ from p_shared in four places.

Three were in:

gridwise_gemm_multiple_d_xdl_cshuffle.hpp

One was in:

gridwise_gemm_xdlops_v2r3.hpp

The change was structurally small:

- void* __restrict__ p_shared
+ void* p_shared

The patch did not change:

workgroup dimensions,
tile dimensions,
synchronization placement,
LDS allocation size,
matrix-instruction selection,
input layouts,
output layouts,
or numerical tolerances.

It removed one compiler promise that was not true for the object.

A one-word qualifier can change generated code

At the C++ surface, this looks like a small declaration edit.

remove __restrict__

At the compiler boundary, it changes the alias relation available to optimization passes.

Before

p_shared
→ exclusive noalias access root

After

p_shared
→ may be affected by accesses not derived
from this thread invocation’s pointer

This can force the compiler to preserve loads or memory dependencies that were previously treated as redundant.

No new runtime instruction needs to be written manually in the source.

Changing the semantic contract changes which machine-code transformations are legal.

More optimization information is useful only when it is true

__restrict__ can improve performance when it describes the program accurately.

For example:

A allocation
→ read only

B allocation
→ read only

C allocation
→ distinct output

Telling the compiler these regions do not overlap can be valuable.

The same reasoning does not justify placing the qualifier on every pointer in a function signature.

A pointer should be marked __restrict__ only when the program can satisfy the corresponding access contract.

The correct question is not:

Would alias analysis become stronger?

It is:

Is the stronger alias statement true for every legal execution?

For block-shared LDS with cross-thread writes and reads, the answer was no.

The physical address alone does not define the violation

It would be too simplistic to say:

Two threads have the same pointer value,
therefore restrict is always invalid.

Read-only access from several threads does not create the same conflicting-write problem.

The relevant combination is:

overlapping thread invocations

+

same logical memory object

+

at least one conflicting write

+

accesses expected to communicate through that object

The CK scratch buffer met that condition.

Threads produce values in LDS

other threads consume those values

barrier orders the communication

The object was a communication channel between threads, not merely a shared immutable lookup table.

A workgroup barrier is evidence of shared ownership

The presence of:

block_sync_lds()

near the scratch-memory path is itself a strong architectural clue.

Private memory does not require a workgroup-wide barrier.

The barrier exists because:

one subset of threads produces state

another subset may consume it

That is exactly the condition a reviewer should compare against a __restrict__ declaration.

Pointer marked exclusive

+

workgroup barrier protecting the same object

→ ownership contract deserves inspection

The barrier does not automatically prove the qualifier is wrong.

But it tells the reviewer that memory ownership is not purely local to one thread.

Shared-memory APIs need to express collective ownership

A raw pointer type does not reveal whether the address refers to:

thread-private memory

workgroup-shared memory

device-global memory

host-visible memory

The semantic meaning exists in:

allocation origin,
address space,
calling convention,
surrounding synchronization,
and documentation.

When an API exposes all of those through:

void*

qualifiers and comments become especially important.

A misleading qualifier can be more dangerous than an absent one because it supplies the compiler with a false proof.

The patch did not add a stronger synchronization primitive

This is important because the visible symptom involved stale shared data.

A common repair might have been:

add another barrier

add a fence

make access volatile

use stronger atomic operations

PR #10574 did none of those.

That is consistent with the root cause.

Actual algorithmic ordering
→ already present

Compiler model of the memory
→ incorrect

Adding synchronization without correcting the alias contract could leave the compiler free to continue treating the object as unaffected by legal external writes.

The source first needed to stop promising exclusive access.

`volatile` would have answered a different question

Marking accesses volatile can force individual loads and stores to remain observable in particular ways.

It does not repair the meaning of an invalid noalias promise.

volatile
→ concerns whether accesses may be removed or reordered
under volatile semantics

restrict / noalias
→ concerns which accesses may refer to the same object

Using volatility to compensate for false alias information would mix two contracts.

The landed fix changed the one that was wrong.

An extra barrier would also have answered a different question

Another barrier could establish another execution rendezvous.

But the kernel already had a barrier separating producers and consumers.

producer writes
        ↓
existing block_sync_lds
        ↓
consumer reads

The mismatch was not:

consumer can physically run before producer

It was:

compiler can conclude that the producer’s cross-thread write
is not a legal modification of this noalias object

More rendezvous points do not make the false exclusivity statement true.

The patch scope was narrow

The change should not be generalized into:

Composable Kernel cannot use __restrict__.

or:

LLVM noalias is unsafe for GPU code.

The confirmed scope is:

four p_shared parameters

in two gridwise GEMM headers

used by a failing grouped-convolution backward-data path

under the reported staging compiler

on MI300X / gfx942

Other pointers may satisfy the qualifier.

Other kernels may use LDS through different signatures.

Other compiler versions may not expose the same wrong result.

The patch repairs the declarations that the reproduced path demonstrated were invalid.

The failure was not established as an MI300X hardware defect

The hardware executed the generated program.

The public evidence points to:

source alias contract

×

compiler optimization

not an LDS coherency defect in the device.

The same hardware passed when:

the qualifier was removed

and the same source passed when:

the relevant compiler alias-analysis behavior was reverted

That comparison is inconsistent with a simple hardware barrier failure.

The exact stale instruction was not published

The PR describes the compiler as being able to retain an LDS value in a register across the barrier.

Its before-and-after matrix strongly supports the alias-analysis mechanism.

But the public PR does not include:

a complete pre-patch LLVM IR dump,
a complete post-patch LLVM IR dump,
the exact load eliminated by alias analysis,
or a machine-code diff identifying one reused register.

The article should therefore distinguish:

Directly reproduced:

wrong numerical output with restrict
correct output without restrict
correct output with older/reverted AA behavior

from:

Mechanistic explanation reported by the PR:

invalid noalias allowed LDS state to remain stale
across synchronization

The source-level contract and validation matrix are strong.

The precise final instruction transformation was not separately published.

The patch reused an existing numerical test

The PR did not add a new dedicated regression-test file.

It used the existing:

TestGroupedConvndBwdData2d/9.Test2D

to show:

before
→ FAIL

after
→ PASS

That is meaningful runtime evidence.

It leaves a maintenance question.

A future change could reintroduce __restrict__ or an equivalent alias contract without a newly added test explicitly documenting:

This test protects cross-thread LDS visibility
under current alias analysis.

The existing numerical test may continue to protect the path indirectly.

The patch does not add a minimal compiler-specific regression asserting the alias or memory-dependency boundary itself.

A focused regression could pin the contract directly

A small GPU test could construct a workgroup-shared buffer where:

one lane writes

all lanes synchronize

another lane reads

the read value participates in the output

The test would then compile under the optimization configuration that exposed the CK failure.

Its purpose would not be to test the hardware barrier in isolation.

It would test:

shared pointer contract

+
cross-thread modification

+
compiler alias analysis

+
barrier visibility

A compiler-level test could separately inspect whether a second LDS load is retained when external thread modification is possible.

These are proposed test directions, not tests added by #10574.

What the patch directly confirms

The merged PR confirms that:

p_shared was declared with __restrict__ in four gridwise GEMM function parameters,
those parameters referred to block-shared LDS used across threads,
the qualifier was removed from all four sites,
the reported grouped-convolution backward-data test failed numerically with the qualifier,
the same test passed after the qualifier was removed,
the qualifier-bearing source passed when the relevant LLVM alias-analysis change was reverted,
the reproduced environment was MI300X / gfx942,
and the patch was merged into rocm-libraries develop.

What the patch does not confirm

The public evidence does not establish that:

every CK kernel using LDS had the same defect,
every use of __restrict__ in CK is invalid,
all AMD GPU architectures reproduce the same numerical failure,
every LLVM version containing the referenced AA change exposes it,
the exact eliminated LDS load is documented in public IR,
the failure affects every grouped-convolution shape,
the patch changes kernel performance,
or a new minimal regression test was added specifically for the alias contract.

The supported conclusion is narrower:

Four block-shared LDS pointer parameters carried an exclusivity contract that the reproduced kernel did not satisfy, and removing that contract restored numerical correctness under the reported compiler.

The source bug and the compiler trigger were both necessary

The failure can be expressed as two conditions.

Condition 1 — invalid source fact

p_shared is workgroup-shared

but

p_shared is declared restrict

Condition 2 — optimizer uses the fact

synchronization no longer forces broad memory effects
for an object considered local and non-escaping

Together:

cross-thread LDS update

→ excluded from the compiler’s legal memory model

→ old value can survive optimization

→ numerical output becomes wrong

Removing either condition hides the failure.

Remove restrict
→ compiler must account for aliasing
→ PASS

Use more conservative AA
→ false promise not fully exploited
→ PASS

Only the first repair corrects the source contract.

The deeper boundary is ownership, not syntax

The code change was one word.

The bug was not grammatical.

void* __restrict__ p_shared

compiled correctly.

The compiler accepted the declaration exactly as written.

The mismatch was between:

declared ownership

and:

actual ownership

The source described the object as exclusive.

The algorithm used it collectively.

Shared state must remain visible to the compiler as shared state

A GPU barrier is meaningful only when the compiler preserves the memory communication it orders.

Producer thread writes state

        ↓

barrier establishes ordering

        ↓

consumer thread reloads state

If a pointer contract tells the compiler that the producer’s write cannot legally affect the consumer’s object, the dependency graph is broken before the barrier reaches hardware.

Synchronization orders real dependencies.
It cannot recover a dependency that an invalid alias contract told the compiler did not exist.

For block-shared LDS, the pointer’s type and qualifiers must leave room for the cross-thread writes that make the scratch buffer useful.

That was the contract repaired by ROCm Composable Kernel #10574.

Part 2: why the barrier could not repair `noalias`

The second and final article examines the compiler side in more detail.

C++ __restrict__

→ LLVM noalias

→ other-thread access semantics

→ synchronization effects in alias analysis

→ non-escaping local-memory reasoning

→ stale value retained across block_sync_lds()

It will distinguish:

hardware execution ordering

from:

compiler memory-dependence modeling

and explain why LLVM’s optimization was valid under the declared noalias contract even though the program behavior expected by CK was not.

Related material

Patch status: Merged into ROCm rocm-libraries develop
Affected component: Composable Kernel gridwise GEMM
Affected pointer: Workgroup-shared LDS scratch pointer p_shared
Source change: Removed four __restrict__ qualifiers across two headers
Observed failure: Numerical mismatch in TestGroupedConvndBwdData2d/9.Test2D
Reported hardware: MI300X, gfx942
Pre-fix result: Max error 130, 137,150 errors, 22.78% incorrect values
Post-fix result: Test passed
Compiler control: Qualifier-bearing source passed with the relevant LLVM AA change reverted
Direct patch scope: Alias contract only; no new barrier, tile, layout, or kernel algorithm
Regression evidence: Existing numerical test used; no dedicated new test file added
Evidence boundary: Exact final load/register transformation was not published as a complete IR or ISA diff

This is Part 1 of a two-part series on ROCm Composable Kernel’s block-shared LDS aliasing boundary.

This article examined why p_shared could not truthfully carry an exclusive __restrict__ contract when the GEMM algorithm intentionally communicated through that LDS object across threads.

Part 2 examines how that source qualifier became LLVM noalias, why noalias constrains accesses from other threads, and why an in-function workgroup barrier could not restore a dependency the compiler had already been told was impossible.

#ROCm #ComposableKernel #AMD #MI300X #LDS #GPUProgramming #Restrict #NoAlias #AliasAnalysis #CompilerOptimization #CodeAnalysis