fix(codegen): match redux.sync's f32 form on the instruction, not on nearby text

[vyncint]

Fixes #1303 for the redux.sync case it was filed for.

contains_redux_f32_features asked whether redux.sync and .f32 appeared in one ;-delimited piece of the LLVM IR. ; starts a comment in IR, so that piece is usually the whole module, and a module reducing integers with any unrelated .f32 was pushed from sm_80 to sm_100a. The integer warp_reduce fast path in #811 is such a module.

How

The predicate decides on tokens.

• IR: lexed with quoted text consumed whole and comments dropped. A global symbol counts when its decoded name is llvm.nvvm.redux.sync. followed by a float operation — the f-prefixed ones, since the float callees carry no type suffix.
• Inline asm: the template is the string after LLParser's asm [sideeffect] [alignstack] [inteldialect] [unwind], decoded as llvm::UnEscapeLexed does, then lexed as PTX. A redux opcode counts when its modifier run — the .name tokens after it — starts with .sync and includes .f32. Whitespace and comments do not end the run; string literals are skipped as data, with a backslash literal inside them; a number consumes its whole run.

Text that only mentions the instruction or callee — a comment, a data or metadata string, a .file path, another symbol's quoted name — no longer counts. The quoted, escaped, unwind, spaced and comment-separated spellings of a real one do.

Tests

31 controls in four tests, each a line from a module llvm-as from the pinned toolchain accepts; 30 also lower through llc (the exception is a global whose quoted name contains a space and @, which NVPTX cannot emit). The PTX controls are checked with ptxas: the float cases assemble at sm_100a and are refused at sm_80 as redux.f32. Twelve mutations of the lexers each fail a test. cargo test -p cuda-oxide-codegen 1264 pass with the existing detection tests untouched; clippy -D warnings, fmt and all 17 scripts/check-*.sh clean.

The float intrinsic spelling was never matched by the old predicate, but that is not a live mis-target: those are catalog intrinsics, so select_target_with_generated takes sm_100a from their recorded requirements regardless.

#1303 stays open for the other conjunctive predicates, which still read raw module text.

[nihalpasham]

The reported integer-redux/f32 failure is real, but splitting by lines leaves the same root cause in place. LLVM accepts two declarations on one line; an integer redux declaration followed by an f32 shuffle declaration still incorrectly selects ReduxF32. An unrelated .f32 SSA name can trigger it too.

Please match the actual intrinsic callee or PTX opcode/modifiers, with negative tests for unrelated names and same-line declarations. The submitted regression passes, but the independent counterexample also passes LLVM verification and still fails feature detection.

This partially addresses #1303; it does not close it yet.

[vyncint]

You were right, and the line split was worse than not fixing it. Force-pushed a different change: the granularity approach is gone, and the predicate now matches the instruction.

Both of your counterexamples, plus the one they led me to

Same six inputs through the real detect_features_in_llvm_text, old predicate against new:

input	want	old	new
integer redux + f32 shfl, separate lines	none	ReduxF32	none
integer redux + f32 shfl, same line	none	ReduxF32	none
integer redux + %acc.f32 SSA name	none	ReduxF32	none
f32 redux, intrinsic @llvm.nvvm.redux.sync.fmin	ReduxF32	none	ReduxF32
f32 redux, @...fmax.abs.NaN	ReduxF32	none	ReduxF32
f32 redux, inline PTX redux.sync.min.abs.NaN.f32	ReduxF32	ReduxF32	ReduxF32

Five of six wrong today, and the last two are the ones I did not expect: the intrinsic spelling of a float reduction contains no .f32 at all. redux_sync_min_f32 lowers to llvm.nvvm.redux.sync.fmin — the operation names the type, and the symbol is not overloaded, so LLVM appends no suffix. The float forms are the f-prefixed ones (fmin/fmax, each with optional .abs/.NaN) against integer add/and/min/max/or/umin/umax/xor, per cuda-intrinsics-gen/src/resolve/families/redux_dotprod.rs.

So a module reducing floats through that path is detected only when an unrelated .f32 happens to sit nearby — the direction that selects a target which cannot run the module. My line split would have removed exactly that accidental coverage while keeping the root cause, which is what your review is pointing at.

Stated as latent, not live: I have not produced a mis-targeted build from it. redux_f32 is pinned to exact sm_100a in the smoketest, so the scan's opinion is not load-bearing there, and I have not checked whether its module carries some other .f32. I will settle that on hardware when I next have a box up for #1305 and report either way.

What it does now

Two matchers, one per spelling that reaches the IR:

• inline PTX through the existing contains_instruction_family_modifier(contents, &["redux.sync"], "f32") — it already bounds the token and splits the modifier list, so redux.sync.min.abs.NaN.f32 matches and redux.sync.add.u32 does not. Reusing it rather than writing a second token scanner also means @llvm.nvvm.redux.sync.add cannot match, since redux.sync there is preceded by . and so is not at an instruction boundary.
• NVVM callee anchored at @, taking the float family by the operation's leading f rather than by listing today's eight spellings — missing one selects a target too old, so the open form is the safer default here, matching the "broad family matching is intentional" note further down the file.

cargo test -p cuda-oxide-codegen 1215 pass (58 existing detection tests untouched), clippy -D warnings clean, fmt clean, all 17 scripts/check-*.sh pass. The diff is +77/−2 in two files; the 21-site granularity change is withdrawn entirely.

#1303 stays open as you said — the other conjunctive predicates still ask their question over module-scope text, and I still have no reproducer for any of them, so I have not touched them.

[vyncint]

Retracting half of my last comment. I said the intrinsic spelling being invisible to this predicate was "a false negative in the direction that matters", and promised to settle on hardware whether it was live. It is not, and the claim was wrong.

I built redux_f32 with nothing pinned, on an A10G, against both predicates:

this branch (token-scoped)       selected: .target sm_100a
upstream (redux.sync && .f32)    selected: .target sm_100a

Its IR is what I expected — eight float callees, no type suffix, and .f32 on zero lines:

@llvm.nvvm.redux.sync.fmin{,.NaN,.abs,.abs.NaN}
@llvm.nvvm.redux.sync.fmax{,.NaN,.abs,.abs.NaN}
lines containing .f32 : 0

So the predicate really is blind to that spelling. What I missed is that it is not the only thing selecting the target: redux_sync_min_f32 is a catalog intrinsic, and select_target_with_generated takes sm_100a from its recorded hardware requirements regardless of what the text scan says. Selection is right today by a different route, and I should not have described it as a live mis-target on the strength of reading the predicate alone.

What stands, unchanged:

• The false positive is real and is what blocked warp_reduce runs a 10-instruction shuffle tree where Ampere's redux.sync does it in one: 4.74x measured, but no way to gate it on the target #811 — integer redux plus any unrelated .f32 reads as the f32 extension. That is the bug this PR fixes, and the counterexamples you gave are in the test.
• The predicate now matching the intrinsic spelling is worth keeping as the thing its name claims, and it covers an f32 redux that arrives outside the catalog path. But it is defence in depth, not a repair of something currently broken, and the PR body says so now.

Nothing in the diff changes; only the claim does.

[nihalpasham]

The narrower redux change fixes the earlier same-line and .f32 SSA cases, and it now recognizes unquoted float intrinsics.

There are still two lexical gaps at b244421d:

• An ordinary data string, comment or metadata value containing @llvm.nvvm.redux.sync.fmin now triggers ReduxF32 and selects sm_100a. The @ prefix does not make the text a callee. This false positive is new.
• In external LLVM without generated catalog requirements, a real call to @"llvm.nvvm.redux.sync.fmin", including its escaped-identifier spelling, misses ReduxF32 and selects sm_80.

The catalog requirement protects the typed Rust intrinsic path, as you clarified. LLVM 23 verifies these external-LLVM probes, and the public detector/target selector reproduce both results despite all 1,260 ordinary tests passing. Please distinguish actual decoded intrinsic symbols and inline-PTX instructions from comments and data, with regressions for these cases.

This is progress on #1303, but it still needs changes before landing.

[vyncint]

Both real, and the first one I introduced. Fixed by giving the predicate a lexical pass over the IR so it is asked about code rather than about text.

The same nine inputs, previous revision against this one

input	want	b244421d	now
callee named in a trailing comment	none	ReduxF32	none
callee named in a c"..." constant	none	ReduxF32	none
callee named in a metadata string	none	ReduxF32	none
PTX text in a comment	none	ReduxF32	none
PTX text in a data constant	none	ReduxF32	none
real call, @"llvm.nvvm.redux.sync.fmin"	ReduxF32	none	ReduxF32
real call, @"llvm.nvvm.redux.sync.\66min"	ReduxF32	none	ReduxF32
real call, bare identifier	ReduxF32	ReduxF32	ReduxF32
real inline-asm template	ReduxF32	ReduxF32	ReduxF32

Seven of nine. Your two are in there; the two PTX rows are the same point generalized — I took "distinguish instructions from comments and data" to cover the inline-PTX half as well, since a redux.sync.min.f32 written in a comment is no more an instruction than a quoted callee is a call.

What it does

normalize_ir_for_detection rewrites the module once, then both matchers run on the result:

• a comment (; to end of line) is blanked — strings are consumed first, so a ; inside one is not a comment;
• ordinary string data is blanked: c"..." constants, metadata, attributes, and the inline-asm constraint string;
• an inline-asm template is kept, since that is how PTX reaches the IR. It is the string whose preceding token is asm or one of its modifiers, which is also what separates it from the constraint string beside it;
• a quoted global identifier is unquoted and its \XX escapes decoded, so all three spellings of the callee read alike.

Blanking rather than deleting keeps newlines and token boundaries where they were, so nothing can be made to match across a hole. LLVM writes a literal quote inside a string as \22, so the next " always closes one — which is what makes the scan simple enough to be worth trusting.

The normalization is used by this predicate only. The other conjunctive predicates in that file still ask their question over raw module text, I still have no reproducer for any of them, and #1303 stays open for that.

Tests

The nine rows above, as two tests — ..._is_scoped_to_the_instruction for the token-scope cases and ..._reads_code_and_not_prose for these. Seven of the nine fail on b244421d, which I checked by restoring that predicate and rerunning rather than by assuming.

cargo test -p cuda-oxide-codegen 1216 pass, the 58 existing detection tests untouched; clippy -D warnings, fmt, and all 17 scripts/check-*.sh clean.

[nihalpasham]

The module-wide false positive is fixed, but the new normalization still loses lexical context. At 135327f4, four LLVM-verified inputs are misclassified:

Input	Expected	Selected
asm sideeffect unwind "redux.sync.min.f32 ..."	float redux	sm_80
Inline PTX with \72edux.sync.min.f32 (\72 is r)	float redux	sm_80
// redux.sync.min.f32 ... inside an asm template	no redux	sm_100a
Ordinary symbol @"unused @llvm.nvvm.redux.sync.fmin"	no redux	sm_100a

The unwind and ordinary-symbol cases are regressions from the original predicate. The other two show why preserving the raw template is insufficient.

Please retain token identity instead of flattening decoded symbols and asm into one search string. Decode LLVM template escapes, recognize the complete asm modifier sequence, and inspect PTX instructions after excluding PTX comments. Add these controls beside the existing tests.

The unmodified codegen suite passes serially; the independent controls fail. One earlier parallel run hit an unrelated fake-opt executable-busy error. This needs changes before #1304 or its dependent #1305 can land. #1303 still covers the broader detector work.

[vyncint]

All four reproduce exactly as you described on 135327f4, and you were right that the flattening was the root cause — keeping the raw template and merging decoded names back into one string lost exactly the boundaries these cases depend on. Replaced it with a lexer, at 279ffe0f.

Your four, plus the neighbours they implied

Same nine inputs, 135327f4 against this revision:

input	want	135327f4	now
asm sideeffect unwind "redux.sync.min.f32 …"	ReduxF32	none	ReduxF32
template with \72edux.sync.min.f32	ReduxF32	none	ReduxF32
// redux.sync.min.f32 … inside a template	none	ReduxF32	none
@"unused @llvm.nvvm.redux.sync.fmin"	none	ReduxF32	none
/* redux.sync.min.f32 … */ inside a template	none	ReduxF32	none
// note\0Aredux.sync.min.f32 … (comment, escaped newline, real instruction)	ReduxF32	ReduxF32	ReduxF32
module asm "redux.sync.min.f32 …"	ReduxF32	ReduxF32	ReduxF32
"\\20redux.sync.min.f32 …"	none	none	none
block label redux.sync.min.f32:	none	none	none

The first five fail on 135327f4. The last four pass under both — they are not regressions fixed, they pin claims the new code makes, and each is load-bearing for a specific mechanism (below).

What it does now

The IR is lexed once — comments dropped, quoted text consumed whole before anything inside it is read — and the predicate asks two questions of the tokens:

• a global symbol, sigil removed and quoting decoded, whose name is llvm.nvvm.redux.sync. then an f operation. A quoted name stays one token, so @"unused @llvm.nvvm.redux.sync.fmin" is one symbol named unused @llvm… and is not a call.
• an inline-asm template containing the instruction. The template is the string at the end of LLParser's own sequence — asm [sideeffect] [alignstack] [inteldialect] [unwind], each optional, in that order. I confirmed that against llvm-as from the pinned toolchain rather than reading it off the grammar: every in-order subset is accepted, and unwind sideeffect, alignstack sideeffect and a repeated sideeffect are rejected. The template is then decoded as llvm::UnEscapeLexed does (\\ is one backslash, \XX a byte), PTX comments are removed after decoding, and only then is it split into instructions.

Decode-then-strip is what the escaped-newline row checks: the raw template is one line, so stripping // before decoding would appear to run to its end and lose the real instruction after \0A.

Every control is valid IR — and two of mine were not

I checked all 23 controls across the three tests with llvm-as, including the earlier rounds', and found two of my own were invalid: the data constants in …_reads_code_and_not_prose had array lengths that did not match their bytes ([31 x i8] for 27 bytes, [26 x i8] for 25), which llvm-as rejects as a type mismatch. Both are corrected; all 23 now parse. The escape cases were also round-tripped through llvm-dis, and llc confirms \\20redux… is emitted as a literal backslash followed by 20redux, not an instruction.

Mutation-checked, which caught a hollow control

Seven mutations, each asserted to have applied: dropping unwind, skipping the template decode, keeping // comments, keeping /* */ comments, matching a symbol anywhere in its name, reading a label as a word, and dropping the \\ rule. The first run caught six. The seventh passed, because my control for the \\ rule — \\72edux — leaves a leading backslash under either decoding and so could not tell them apart. It is replaced by \\20redux, which a decoder without the rule reads as a space followed by a real redux.sync.min.f32; that mutation is now caught too.

cargo test -p cuda-oxide-codegen 1266 pass across all suites, with the 58 existing detection tests and the upstream bare-PTX redux test untouched; clippy -D warnings, fmt, and all 17 scripts/check-*.sh clean. #1305 is rebased onto this and updated.

One scope note: a bare word that is itself the instruction also counts, because an existing upstream test feeds a raw PTX fragment (redux.sync.min.abs.NaN.f32 $0, $1, $2;) and expects ReduxF32. Both production callers pass LLVM IR, where PTX never appears outside a string, so that rule cannot match valid IR. It does mean several bare PTX statements on one line would be read with IR comment rules after the first ; — a shape no IR input can have, and I have not changed how the other feature families handle it.

[nihalpasham]

The four earlier cases now pass at 279ffe0f. Keeping LLVM token identity fixes those boundaries. Two issues remain in the PTX template scan:

Valid PTX	Expected	Selected
redux.sync.min .f32 $0, $1, $2;	float redux	sm_80
redux.sync.min/* qualifier */.f32 $0, $1, $2;	float redux	sm_80
redux .sync .min .f32 $0, $1, $2;	float redux	sm_80
.file 1 "name redux.sync.min.f32 label.cu"	no redux	sm_100a

These are complete LLVM modules, not malformed snippets. All four pass llvm-as, llc and CUDA 13.4 ptxas at sm_100a. The quoted-file control also assembles at sm_80; the spaced reduction fails there with Instruction 'redux.f32' not supported on .target 'sm_80'. The actual crate reproduces all four wrong selections.

Please recognize the PTX opcode and modifiers across whitespace/comments, and keep quoted directive data out of instruction detection. Add these controls to the tests. Keep the helper comments focused on their current contract.

The unmodified codegen suites pass (1,263 tests/doctests), as do strict Clippy and formatting. The independent PTX controls fail. #1304 and its dependent #1305 still need changes; the original four LLVM cases are resolved.

[vyncint]

All four reproduce on 279ffe0f, and I checked each ground truth with the same chain you used — llvm-as, llc, then ptxas (13.3 here): the three spaced forms assemble at sm_100a and are refused at sm_80 as redux.f32; the .file control, as module asm, assembles at both. Fixed at 6b4ac7d0.

What changed

Templates are now lexed as PTX rather than split on whitespace. Whitespace and // / /* */ comments separate tokens without ending anything, so an opcode's modifiers are just the run of .name tokens after it — redux.sync.min.f32, redux .sync .min .f32, min/* q */.f32 and a newline between modifiers all read as one instruction. String literals are skipped as data, which takes the .file path out.

Two lexer rules came from checking ptxas rather than from the PTX spec, and both were bugs in my first pass:

• A backslash is literal inside a PTX string. ptxas accepts .file 1 "C:\dir\", so the next quote always closes a string. I had written C-style escape handling, which would have swallowed the rest of the template after such a path. Control: .pragma "nounroll\"; followed by a real redux.sync.min.f32 in the same template, which ptxas refuses at sm_80 as redux.f32.
• A number consumes its whole run. Without that, \\20redux… — a literal backslash then 20redux — lexed as 2, 0, redux and conjured an opcode. The existing \\20redux control caught it.

Controls, 279ffe0f against this revision

input	want	279ffe0f	now
redux.sync.min .f32	ReduxF32	none	ReduxF32
redux.sync.min/* qualifier */.f32	ReduxF32	none	ReduxF32
redux .sync .min .f32	ReduxF32	none	ReduxF32
.file 1 "name redux.sync.min.f32 label.cu"	none	ReduxF32	none
redux.sync.min\n.f32	ReduxF32	none	ReduxF32
.pragma "nounroll\"; then a real float redux	ReduxF32	ReduxF32	ReduxF32
redux .sync .min .u32 (spaced integer)	none	none	none
{ .reg .f32 t; redux.sync.min.u32 …; mov.f32 t, …; }	none	none	none

The last three pass under both. They pin rules the new lexer has to keep: strings close at the next quote, spacing does not make an integer reduction float, and .f32 counts only as a modifier of the redux opcode, not anywhere in the template. That last one was missing until a mutation replacing the modifier check with ".f32 anywhere" survived; it is caught now.

Twelve mutations across the IR and PTX lexers, each asserted to have applied, each failing a test.

One claim I corrected before posting

I had written that every control lowers through llc. Checking all 31 (not just this round's): 30 do. The exception is your earlier @"unused @llvm.nvvm.redux.sync.fmin" global, which is valid IR that NVPTX refuses to emit ("Symbol name with unsupported characters"). It still does its job — it tests that a quoted name stays one token — and the commit message now says so rather than overclaiming.

Comments are trimmed to the current contract. cargo test -p cuda-oxide-codegen 1264 pass across all suites (your 1263 plus the new test); clippy -D warnings, fmt and all 17 guards clean. #1305 is rebased onto this.