How did we find a GPU instruction that rewrites the block ID?

What does a GPU do with an instruction its compiler will never write? NVIDIA's tools list one called SETCTAID.X, "Set CTA ID", from Volta through Blackwell. CTA stands for Cooperative Thread Array, which is PTX and SASS's name for what CUDA code calls a thread block, so the CTA ID is the same value you read as blockIdx.x. Since there is no PTX spelling for SETCTAID.X, the only way to find out what it does is to patch it into a cubin by hand, and on a GB10 (sm_121a, CUDA 13) the answer turned out to be more interesting than a crash.

Short version#

PTX swears %ctaid is read-only. Slip SETCTAID.X into a running kernel anyway, and the block ID a warp reads back can change under it. Block 1 read 1 before the instruction, 0 after SETCTAID.X RZ, and 7 after SETCTAID.X R11 with 7 in R11, in 3 of 3 fresh-process repeats. The new value only showed up after the warp had issued an LDS, and why that load matters is still open. Shared memory stayed unchanged and both blocks stayed on SM 47.

The first test fooled us#

The first time SETCTAID.X ran in a live kernel, both block ID reads came back the same. It looked like the instruction did nothing, and we almost left it as a listed SASS op with no effect.

The mistake was in our observer. It read the block ID before the patch and wrote that same saved value into both result slots, so the two numbers had to match. In September a short thread posted the fixed result: block 1's S2R came back 0 after SETCTAID.X RZ, even though PTX still calls %ctaid read-only. That thread never explained the LDS link, and it never showed the bad first test which is covered further in this article.

How to read this#

A GPU kernel runs as many blocks, and each block learns its own position from blockIdx.x, which the hardware keeps in a special register called SR_CTAID.X. A warp reads it with an instruction called S2R, which copies a special register into an ordinary one.

PTX is the portable text form CUDA compiles to, and SASS is the machine code the chip actually runs. So an instruction can exist in SASS and have no spelling in PTX at all.

The setup#

SASSquatch, an open-source GPU instruction auditor, compiled a small CUDA kernel into a cubin and found four NANOSLEEP slots in it. Those slots got overwritten with candidate instructions, and every launch ran in a fresh child process under timeout -k 2 8, with GPU visibility checked after each one.

To put two test blocks on the same SM, 47 large blocks went out first and each parked on their own SM. So the two test blocks were forced onto the one SM left, SM 47, which also let us ask whether one block could see the other's shared memory.

The instruction under study is SETCTAID.X, which NVIDIA's CUDA Binary Utilities list as "Set CTA ID" from Volta through Blackwell. It has no PTX opcode, takes one source register and writes it into SR_CTAID.X. The two encodings used here were SETCTAID.X RZ (0x00000000ff05731f), where RZ always reads as 0, and SETCTAID.X R11 (0x000000000b05731f), where R11 held 7.

The false lead#

The first proof of concept said the instruction did nothing. The observer in that kernel read the block ID before the patched instruction and then stored that saved value twice, so both outputs matched and the instruction looked like a no-op.

The rebuild made the observer perform a real S2R of SR_CTAID.X after the instruction. The repaired sequence is short enough to show as an image.

What the repaired probe showed#

With the repaired observer, block 1 read 1 before the instruction and 0 after SETCTAID.X RZ, identical in 3 of 3 fresh-process repeats. With SETCTAID.X R11 and 7 in R11, the later read returned 7, again 3 of 3.

So the instruction writes the value from its source register into the block ID latch, and a later S2R can return it. That overturned the earlier "always a no-op" conclusion, at least for this encoding and kernel shape.

The LDS condition#

A simpler repaired kernel still read 1 after the instruction, so something else mattered. Comparing variants showed that the new value only appears once the warp has issued an LDS, a load from shared memory.

Allocating __shared__ memory or waiting on the timer was not enough. Whether the LDS commits the written value or flushes a cached launch-time ID is unknown, and special registers and shared memory are separate files, so this link between them is unexplained.

What shared memory did#

The same runs also tested whether a block that changes its own ID can then reach another block's shared memory, and it could not. The attacker block's word held 0x11111111 before and after the instruction, the victim block's word kept its 0xDEADBEEF, both blocks stayed on SM 47, and their DRAM pointers did not move. The instruction changed the number S2R hands back, and nothing about where the block runs or what memory it can touch.

---

Put together, SASS has a way to write the block ID that CUDA never generates, and on this GB10 that write shows up in the next S2R of SR_CTAID.X once the warp has issued an LDS. It showed up in every repeat with that load in place, and neither a shared-memory store nor a globaltimer wait could take the load's place.

That does not make it a security hole. Shared memory stayed the same, both blocks stayed on the same SM, and the only thing that changed was the number a block gets when it asks for its own ID, which gives it no way into another block's data.

It is not a world first either, because the instruction name already sits in NVIDIA's documentation. What seems to be missing from NVIDIA's spec, as far as we could find, is the behavior measured here, where the new value lands in the latch and S2R keeps returning the old one until an LDS. Looking through GitHub, the web, X and forums turned up nothing, though an empty result proves little, and NVIDIA's closed libraries were never checked. Production kernels only use the block ID and never set it, PTX documents %ctaid as read-only, and every dump we found with the name in it came from one of the cubins we patched ourselves.

The whole result comes from one chip, one CUDA version and a few small kernels, which leaves three questions without answers: why the LDS matters, what the instruction was built for, and whether other encodings and kernel shapes behave the same way. Some patches were kept out on purpose, namely the known hang opcode 0x3d0 and anything with BPT, TRAP, EXIT, BAR, BSSY or WARPSYNC, and no full opcode scan was run. After every child process exited, the GPU was still visible to the system.