Fix/pliron codegen regressions - #1508
Merged
Merged
Conversation
The SROA pass can scalarize a local array into a plain variable, which mem2reg then promotes with a cube.poison initial value threaded through loops as an iter-arg. cube.poison had no SPIR-V lowering, so emission failed on any kernel where this happened (e.g. the plane vec-mat matmul), and the failure was an .expect() panic on the device thread: the launch silently produced uninitialized output instead of a compile error. Lower poison to spirv.Undef and propagate to_spirv failures as a CompilationError. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01N6VLz7xTBVNmCwqhWyZVkY
A kernel with two early returns lowers the second one via
predicate_on_flag with the insertion cursor inside the then-block of the
first predication — a block that already ends with the yield placed when
that predication was built. Unconditionally appending another yield left
the block with two terminators, which fails module verification on every
backend ("Basic block ... has a terminator that is not the last
operation in the block"), and the launch error was silently swallowed:
the kernel simply never ran. Terminate the block only when it isn't yet,
mirroring terminate_yield.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01XbtZ8f9KveQLpd3kvL79tg
The unroll pass splits vectors wider than the target's maximum (e.g. 8-wide q4 lanes against SPIR-V's max of 4) op by op, and cast is element-wise — but CastOp never opted into TriviallyUnrollable, so any kernel casting a wide vector (f32 -> f16 on an 8-lane accumulator, say) panicked the device thread inside UnrollPass. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01XbtZ8f9KveQLpd3kvL79tg
The cpp printer emitted fma() verbatim for any operand type. CUDA/HIP headers have no vector overloads, so a vector fma failed hiprtc/nvrtc compilation outright, and no __half/bf16 overload either, making the scalar half call ambiguous between the double/float/_Float16 candidates. Register FmaOp with the cpp unroll rewrite like the other math functions, and pick __hfma/__hfma2 for half types outside Metal (whose fma is generic). Also mark FmaOp trivially unrollable in the IR so the core unroll pass can split wide vector fmas instead of panicking — the same latent gap CastOp had. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01XbtZ8f9KveQLpd3kvL79tg
wingertge
requested changes
Aug 12, 2026
wingertge
reviewed
Aug 12, 2026
The pliron migration dropped the pre-pliron dialect's conditional includes. __half survives on hiprtc's builtin header, but __hip_bfloat16 does not — the builtin only declares the legacy hip_bfloat16 struct — so every bf16 kernel failed hiprtc compilation with 'unknown type name'. Register the include on the type, the same way the CUDA target ties cuda_bf16.h to its bf16 types. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01XbtZ8f9KveQLpd3kvL79tg
The HIP bf16 max polyfill (bf16 has no __hmax on HIP, so min/max promote through f32) was a copy-paste of min_bf16 that still called .min(). Every bf16 max on ROCm returned the minimum — an online-softmax running max walks toward -inf, its exp() overflows, and attention outputs go NaN across the board. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01XbtZ8f9KveQLpd3kvL79tg
Drivers only accept the exact cooperative matrix (element type, rows, columns, use) combinations they advertise; any other shape is undefined behavior at runtime — on RADV the loads and stores silently produce garbage. Walk the assembled module for subgroup-scope OpTypeCooperativeMatrixKHR declarations in validate_ir and reject fragments the device doesn't support with a proper CompilationError. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01N6VLz7xTBVNmCwqhWyZVkY
Two changes to keep dynamic pool memory bounded on long-running processes, motivated by an OOM kill on a unified-memory APU (ROCm): - The SubSlices tail pool used max_page-sized sliced pages, so the first allocation bigger than the ladder materialized a page a quarter of device memory in size. Use exact-size exclusive pages instead, with a scaled dealloc period. - The pools' dealloc_period machinery was dormant: nothing invoked cleanup periodically, so freed pages were never returned to the driver outside explicit cleanups. Drive a non-explicit cleanup from reserve — each pool self-gates (sliced pools ignore it, exclusive pools check their period), and the direct pool now explicitly opts out to preserve its documented hold-below-the-ceiling behavior. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01N6VLz7xTBVNmCwqhWyZVkY
A kernel that fails at launch (e.g. a compilation error) pushes its error onto the stream, and sync/profiling flush strictly — but read_resources flushed with ignore+no-flush, so a plain readback returned stale bytes with Ok instead of the recorded error. That made every compile failure on wgpu silent unless profiling happened to be on; the HIP runtime already fails such reads. Flush strictly in read_resources and resolve the read future to the error. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01XbtZ8f9KveQLpd3kvL79tg
wingertge
reviewed
Aug 12, 2026
Review feedback: the comments explained why marking the ops TriviallyUnrollable fixed a regression, which doesn't belong in the code permanently. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01C982k1iD6aZRL9j73NuKNb
The pass runner around ToSpirvDialectPass unwrapped, so a conversion error panicked the device thread and poisoned the server: every subsequent kernel on the stream read garbage instead of surfacing one clean error. Propagate like the to_spirv step already does. The cube.poison lowering to OpUndef stays for now: SROA + mem2reg promote local arrays with poison as the initial loop iter-arg, which is always overwritten before use, and real kernels (register matmul with promotion, bf16/quantized heads) rely on it compiling. It can be removed once uninitialized arrays no longer reach the backend. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01C982k1iD6aZRL9j73NuKNb
fma wasn't called directly by any test, so the cpp scalarization and half intrinsic paths never got validated. Cover vectorization 1, 2 and 4 to exercise the scalar, __hfma/__hfma2 and unroll paths. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01C982k1iD6aZRL9j73NuKNb
The CPU server unwrapped the compile result on the device thread, so a compilation error panicked instead of surfacing at flush/sync like the other servers. Push it on the stream's error queue instead. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01C982k1iD6aZRL9j73NuKNb
Review feedback: fragment shapes and types come from the high-level cube instructions, so scanning the assembled SPIR-V module was both backend-specific and redundant. Check the declared fragment against the device's advertised configs in Matrix::uninitialized, through the existing scope error channel every backend drains at compile time, and drop the SPIR-V module scan and its wgpu validate_ir hooks. Add a runtime test asserting an unsupported fragment fails compilation. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01C982k1iD6aZRL9j73NuKNb
nathanielsimard
force-pushed
the
fix/pliron-codegen-regressions
branch
from
August 13, 2026 13:31
5f00c16 to
d0c5dc4
Compare
wingertge
approved these changes
Aug 13, 2026
…gressions # Conflicts: # crates/cubecl-cpp/src/shared/operation.rs
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.