[ck_tile] FastGeluAsm::operator()<float, float> returns wrong results from -O1 up: inline-asm scratch operand lacks earlyclobber

[doplxyz]

Summary

ck_tile::element_wise::FastGeluAsm::operator()<float, float> declares its inline-asm scratch
operand as [v_tmp] "+v"(tmp) with tmp uninitialized, and marks nothing earlyclobber. The
asm writes tmp before it has finished reading its inputs, so the register allocator is free to
give tmp the same register as one of those inputs — which it does from -O1 upward, and the
function then returns wrong values.

This is a legal allocation for the constraints as written, not a compiler defect. The one-line
fix is to describe the scratch operand accurately.

The packed fp32x2_t specialization directly below did not misbehave in any configuration I
measured.

Affected code

include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp, struct FastGeluAsm, the
CK_TILE_DEVICE operator()<float, float> specialization (line ~1264 at
5a74dec07a894484b9489d0c0e00cd3b52652d18, unchanged on develop as of 2026-08-19):

const uint32_t c1     = 0xbd92220c;
const float c2        = -2.0 * 0.797885f;
const uint32_t log2e_ = 0x3fb8aa3b;
float tmp;                                   // never initialized

asm volatile("v_mul_f32 %[v_tmp], %[v_x], %[v_x]        ; x*x\n"
             ...
             "v_mul_f32 %[v_y], %[v_tmp], %[v_x]        ; x * 1/(emu+1f)\n"
             : [v_y] "=v"(y), [v_tmp] "+v"(tmp)         // no earlyclobber
             : [v_x] "v"(x), [s_c1] "s"(c1), [v_c2] "v"(c2), [s_log2e] "s"(log2e_)
             :);

There are two separate constraint defects here, and it is worth keeping them apart:

• What allows the overlap is that the scratch output is not earlyclobber. Without &, an
  output may share a register with an input, and this asm writes tmp before it has finished
  reading x and c2. This is the one that produces the wrong results.
• Separately, "+v"(tmp) declares tmp read-write, i.e. that its incoming value matters. It
  does not, and tmp is uninitialized, so an undefined value is being passed in as an input.

The matrix below separates them: making the scratch write-only without earlyclobber
("=v"(tmp)) removes the second defect and leaves the first, and it fails at every level.

Reproduction

ec_check.hip: 256 inputs evenly spaced over [-8, 8], one thread each, compared against a
double-precision reference x / (1 + exp(x * (c1*x*x + c2))) with c1 = -2*0.035677,
c2 = -2*0.797885. max_rel is taken only where |ref| > 1e-3, because relative error is
unstable where the reference is close to zero (at x = -8 it is about -3e-21), not because
anything underflows.

x0 below is the sample nearest to 1 — index 143 of 256. The inputs are generated in float
(-8.f + 16.f * i / (n - 1)), so the exact value is x0 = 0.9725494384765625, not 248/255;
exact 1.0 is not one of the sample points. At that coordinate the reference is 0.811557765, and
the passing builds give f(x0) = 0.81155771, i.e. |f(x0) - ref(x0)| = 5.495e-08 — under the
reported max_abs, which occurs elsewhere in the sweep.

clang++ -I<ck>/include -I<rocm>/include --offload-arch=gfx1201 -std=gnu++20 -O3 -DNDEBUG \
        ec_check.hip -o ec_check && ./ec_check

Same hardware (RX 9070 XT, gfx1201), same inputs, same CK pin, both compilers, stock header
versus the one-line change below:

		-O0	-O1	-O2	-O3	-Os	-Oz
AMD clang 23.0.0git (ROCm 7.14)	stock	ok	bad	bad	bad	bad	bad
	patched	ok	ok	ok	ok	ok	ok
AMD clang 22.0.0git (ROCm 7.2.4)	stock	ok	bad	bad	bad	bad	bad
	patched	ok	ok	ok	ok	ok	ok

• "ok" = max_abs 3.696e-07, max_rel 7.710e-07, f(x0) = 0.81155771 (reference 0.811557765)
• "bad" = max_abs 8.000e+00, max_rel 3.152e+03, f(x0) = 0.290346831

At -O3 the function returns ≈x for large negative input and ≈0 for large positive — the sign
of the exponent argument is effectively inverted.

The aliasing, in the generated ISA

Both compilers, -O3, gfx1201, stock header — %[v_tmp] and %[v_c2] are both v3:

v_mul_f32 v3, v2, v2        ; x*x
v_fma_f32 v3, v3, s0, v3    ; c1*x*x+c2      <- v_tmp and v_c2 are both v3
v_mul_f32 v3, v3, v2        ; x*(c1*x*x+c2)
v_mul_f32 v3, v3, s1        ; log2e*x*(c1*x*x+c2)

With the change, tmp gets v4 and c2 keeps v3, and the results match the reference.

Which input tmp collides with depends on the surrounding code: in a different TU where the
specialization is inlined into a kernel that also uses the packed form, I saw it take x's
register instead (v_mul_f32 v1, v1, v1 ; x*x). The constant is that the scratch has no
earlyclobber, not which operand it lands on.

Which constraint is responsible

Seven constraint variants x six optimization levels, gfx1201, clang 23 (42 configurations; the
stock-vs-patched A/B above is a separate set of 24 runs covering both compilers):

[v_tmp] / [v_y]	-O0	-O1	-O2	-O3	-Os	-Oz
"+v"(tmp) / "=v"(y) — as shipped	ok	bad	bad	bad	bad	bad
"=v"(tmp) / "=v"(y) — write-only, no earlyclobber	bad	bad	bad	bad	bad	bad
"+v"(tmp) / "=&v"(y) — earlyclobber on the output only	ok	bad	bad	bad	bad	bad
"+&v"(tmp) / "=v"(y)	ok	ok	ok	ok	ok	ok
"=&v"(tmp) / "=v"(y)	ok	ok	ok	ok	ok	ok
"+&v"(tmp) / "=&v"(y)	ok	ok	ok	ok	ok	ok
"=&v"(tmp) / "=&v"(y)	ok	ok	ok	ok	ok	ok

The cleanest comparison is the pair that differs only in the &:

"=v"(tmp)  / "=v"(y)   ->  bad at every level
"=&v"(tmp) / "=v"(y)   ->  ok at every level

Both pass tmp as a write-only operand, so the uninitialized-input defect is absent from both;
the only difference is earlyclobber, and it is the difference between wrong and right.

Two other rows are worth noting. Earlyclobber on the output alone does not fix it — it produces
a different wrong answer (f(x0) = 0.00350051583 rather than 0.290346831), which tells us the
output is not where the problem is, not that allocation is deterministic. And write-only without
earlyclobber is worse than what ships: it fails at -O0 too, which is where the current "+"
happens to save it.

Within these 42 configurations — this asm, this compiler, gfx1201, these optimization levels —
every variant carrying earlyclobber on the scratch passed all six levels and every variant without
it failed from -O1 up. I have not tested other compilers, targets or inputs against this matrix.

Suggested fix

--- a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp
+++ b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp
@@
                      "v_mul_f32 %[v_y], %[v_tmp], %[v_x]        ; x * 1/(emu+1f)\n"
-                     : [v_y] "=v"(y), [v_tmp] "+v"(tmp)
+                     : [v_y] "=v"(y), [v_tmp] "=&v"(tmp)
                      : [v_x] "v"(x), [s_c1] "s"(c1), [v_c2] "v"(c2), [s_log2e] "s"(log2e_)
                      :);

One operand. =&v states what is actually true: the asm's first instruction
(v_mul_f32 %[v_tmp], %[v_x], %[v_x]) writes tmp before anything reads it, so it is write-only
and early-clobbering. Dropping the "+" also stops an uninitialized tmp from being passed in as
an input. y is written only by the final instruction, after every input has been read, which is
consistent with it not needing earlyclobber — and the table above shows earlyclobbering it does
not help.

Happy to open a PR if that is easier than applying it directly.

Scope and severity

At pin 5a74dec, FastGeluAsm appears three times in the whole tree: the struct, its name
member, and one use at
example/ck_tile/15_fused_moe/instances/fused_moegemm_api_internal.hpp:28 (Activation == 0).
That call site reaches only the fp32x2_t specialization
(fused_moegemm_pipeline_flatmm_uk.hpp:351,385), which was correct at every optimization level I
measured. Its operands are "+v" on initialized values, which removes the uninitialized-input
defect, but by the same argument I apply to SiluAsm below that does not by itself settle the
earlyclobber question — I have not analysed its full asm data flow, only observed that it produced
correct results everywhere I ran it. So no in-tree caller at this pin reaches the specialization
this issue is about; I have not assessed out-of-tree users, and this is a public header.

SiluAsm, the other struct in that file with asm volatile, is inside an #if 0 block (the
comment above it says it is kept "purposely if in the future ppl want to try"), so it is not
compiled at this pin and is out of scope here. I have not assessed whether it would have the same
problem if it were enabled — its operands are "+v" on values that are initialized, but that on
its own does not settle the earlyclobber question.

Environment

GPU	AMD Radeon RX 9070 XT (gfx1201, RDNA4). gfx942 and gfx950 are compile targets only here
Compilers	AMD clang 23.0.0git (ROCm/llvm-project 46fcb339 +PATCHED:440716f8) in rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0; AMD clang 22.0.0git (roc-7.2.4 f58b06dc) in rocm/pytorch:rocm7.2.4_ubuntu24.04_py3.12_pytorch_release_2.9.1
Host	x86_64, Ubuntu 24.04, kernel 6.14.0-37
CK	5a74dec07a894484b9489d0c0e00cd3b52652d18 (PyTorch's current submodule pin; identical on develop as of 2026-08-19)

The same v_tmp aliasing appears in the generated assembly when targeting gfx942 and gfx950.
That is a static observation of the emitted ISA only — I have no such hardware and have not run
anything there.

Reproducer, the constraint matrix and full logs: https://gist.github.com/doplxyz/8b157ad3a3a33b6478f857567466054f/63fcd75782259db429e28c94eabdb28f50c83501