SCALE rejects an inline-PTX `bra` to a label defined in a separate `asm()` statement ("unknown label")

[diskdog]

Problem

SCALE resolves inline-PTX labels only within a single asm() statement. A bra targeting a label defined in a different asm() statement fails with unknown label. nvcc concatenates all inline asm in a function into one PTX unit, so the cross-statement label resolves. The label name is irrelevant; the same label inside one asm() compiles.

Reproducer

__global__ void k() {
    asm volatile("bra RET;");
    asm volatile("RET:");
}

nvcc -arch=sm_75 -std=c++17 -c compiles. SCALE to NVIDIA (llvm/bin/nvcc -arch=sm_75 -std=c++17 -c) fails:

error: unknown label: RET

Likely cause

SCALE scopes inline-PTX labels per asm() statement rather than per function, so a branch cannot target a label emitted by a separate asm() statement. nvcc resolves such cross-statement labels, so this is an nvcc-compatibility gap.

Toolchain: SCALE unstable-2026.07.08 (clang 20.1.8), NVIDIA target sm_75.

[ChrisKitching]

This one has been known for a while, we just didn't prioritise it since we haven't yet found a real-world project that acutally does something this crazy! :D