Grid barrier state is initialized without inter-block synchronization

[tyler-utah]

Summary

decode_kernel initializes its global grid-barrier state from block 0 and then uses __syncthreads() before other blocks access that state. Because __syncthreads() only synchronizes threads within one block, other blocks can reach atomicAdd(barrier_counter, 1) concurrently with block 0's initialization.

This affects the current main branch at commit bab084443f14e6ab3ba76f8e8e3f721ad2e56144.

Problem

The kernel currently performs the equivalent of:

if (blockIdx.x == 0 && threadIdx.x == 0) {
    *barrier_counter = 0;
    *barrier_generation = 0;
}
__syncthreads();

if (threadIdx.x == 0) {
    atomicAdd(barrier_counter, 1);
    // ...
}

There is no grid-wide ordering between the initialization and the atomic arrival from another block. For example, another block can increment barrier_counter before block 0 stores zero, allowing the initialization store to erase that arrival.

The barrier can consequently enter an inconsistent counter/generation state. Depending on scheduling, blocks may wait for a generation that is never reached or otherwise proceed with incorrect barrier state, potentially hanging or mis-synchronizing the persistent decode kernel.

The barrier tensors being initialized to zero when the Python model is constructed does not address repeated launches because the generation changes during every decode step.

Suggested fix

Reset barrier_counter and barrier_generation on the CUDA stream immediately before launching decode_kernel, remove the in-kernel initialization/bootstrap barrier, and initialize the kernel-local generation to zero.

This makes initialization stream-ordered and complete before any block begins executing the kernel.

I have a minimal patch prepared for this change.

Research context

This issue and fix were identified, analyzed, and documented during a fine-grained, interactive debugging session with GitHub Copilot, powered by GPT-5.6. The work was conducted as part of a research project on GPU testing and verification that is currently under submission.