CUDA Tile (cuTile) first working port in Java

Developer mikepapadim opened pull request #1083 against the TornadoVM project, adding a CUDA Tile (cuTile) execution path to the CUDA backend as a TileContext API across 30 commits. The change introduces PartitionView.atomicAdd for single-kernel split-K matmul and brings the tile test suites to 91 passing tests with checkstyle clean. The PR was submitted for review by kotselidis, mairooni, orionpapadakis, and stratika on September 11, 2026.

cuTile Java CUDA Tile cuTile execution path to the CUDA backend as TileContext API - 1083 mikepapadim /mikepapadim wants to merge 30 commits into \ cuTile\ Java CUDA Tile cuTile execution path to the CUDA backend as TileContext API top 1083 mikepapadim /mikepapadim wants to merge 30 commits into \ cuTile\ Java CUDA Tile cuTile execution path to the CUDA backend as TileContext API top 1083 mikepapadim /mikepapadim wants to merge 30 commits into Conversation mikepapadim /mikepapadim requested review from kotselidis /kotselidis , mairooni /mairooni , orionpapadakis /orionpapadakis and stratika /stratika September 11, 2026 09:24 event-30961982952 mikepapadim /mikepapadim changed the title ~~ CUDA Graph Capture a JIT + cuTile + cuBLAS pipeline in a single graph~~ cuTile Adds the CUDA Tile cuTile execution path to the CUDA backend: a TileContext Sep 11, 2026 event-30962062069 mikepapadim /mikepapadim changed the title ~~ cuTile Adds the CUDA Tile cuTile execution path to the CUDA backend: a TileContext~~ cuTile Adds the CUDA Tile cuTile execution path to the CUDA backend as TileContext API Sep 11, 2026 event-30962078312 | Atomic accumulation through a view — split-K now finishes in one kernel Next item from the parity list 7bbfad2 . PartitionView.atomicAdd tile, blockIndices... accumulates a tile so every concurrent block's contribution survives, which collapses the two-kernel split-K arrangement write per-split partials, reduce in a second pass into a single launch with no intermediate buffer. It does not lower like a store, and the header is why I checked the 13.3 header rather than assuming a symmetry: partition view has no atomic read-modify-write. It offers atomic load and atomic store only; ct::atomic add takes a tile of pointers : template<memory order Order, thread scope Scope = default thread scope , storeable tile Ptrs, non narrowing tile convertible to<tile load t<Ptrs Values tile builtin tile load t<Ptrs atomic add Ptrs ptrs, Values values, ... ; Two consequences that shaped the implementation: Order is memory order relaxed , scope is thread scope device . The scope is not cosmetic — measured SASS: The split-K kernel gets RED rather than ATOMG because the returned old value is unused and the compiler picks the fire-and-forget reduction form. Default system scope would emit the more expensive .SYS variants. TestTileAtomics — 4 tests, none of which can pass by accident | | Kernel | Instruction | |---|---| | row accumulator | ATOMG.E.ADD.F32.FTZ.RN.STRONG.GPU | | split-K matmul | RED.E.ADD.F32.FTZ.RN.STRONG.GPU | Every test targets an output region that several blocks write, so a plain store would lose contributions non-deterministically. A test that would also pass with a normal store would prove nothing. | Test | Contention | |---|---| | testAccumulateRowsRank2 | 96 blocks accumulate into the same 64-element row | | testAccumulateRank1 | 64 blocks, rank-1 view the simpler pointer arithmetic | | testSplitKTwoSplits / testSplitKFourSplits | single-kernel split-K matmul, K split along bidZ , 2 and 4 blocks contending per output tile | Tile suites now 91 tests, all passing ; checkstyle clean. Two failures on the way, both worth knowing for anyone extending the plugins - receiver.get true leaves a NullCheckNode behind when a plugin does not consume the receiver as a node input, and the CUDA backend has no emitter for one — it failed with a bare TornadoInternalError: unimplemented at LIR generation with no indication of the cause. This plugin uses get false . - Bypassing the view node means the view has no usages and is eliminated , so nothing marked the accumulator as written. The dataflow analysis called it read-only, no TRANSFER DEVICE TO HOST was emitted, and the accumulator read back as all zeros while the kernel was in fact doing the atomics correctly. CUDATileAtomicAddNode now implements MarkArrayParameterAccess itself and reports READ WRITE . This is the same class of bug as the half-float read marker fixed earlier in this branch, and the second time an access-marking omission has produced silently wrong data rather than an error — worth remembering that any new node touching a buffer must mark it. Parity list, updated Closed since the audit: comparisons + select , minimum / min / prod / fma , int8/double/fp8 views, atomicAdd . Still open: the rest of the atomic family atomic sub / min / max , bitwise, xchg , compare exchange , the masked forms, and a choice of order/scope , rank ≥ 3 views , permute / reshape / broadcast / extract as explicit operations, scans partial sum / partial prod , the remaining scalar math, per-operation rounding tags, and still no num ctas / occupancy / latency hints. The shape operations, each replacing a workaround that earlier kernels in this branch had to use: - broadcast: the mask kernels stretched an iota row to the score shape by adding a zero tile of that shape, using arithmetic to do a layout job. TestTileMasking now says broadcast, which is what it meant. - reshape: changing a rows, 1 reduction result to another shape of the same size previously meant transpose, a data movement rather than a reinterpretation. It is also the only way to change rank. - extract: splitting a packed row meant loading the same row twice through a narrower partition view; one load and two extracts is what the kernel actually wants. extract's indices are in sub-tile units, not elements - verified by running it rather than reading it, since the header does not say: extract row, shape<1,32 , 0, 1 of a 1, 64 tile returns elements 32..63. permute is deliberately not added: on a rank-2 tile the only non-identity permutation is the transpose that already exists. All three lower through one node and statement, emitting ct::fn tile, ct::shape<... {} with extract's indices appended. The plugin builds its signature as a single Type with the receiver first, because the InvocationPlugin constructor takes Type... and a nested Class is one argument of the wrong type rather than a spread signature. TestTileShapes 4 tests : a bias row broadcast across a tile, a row-sum column broadcast to normalise, a reshape of a 1,1 reduction to rank 1 for a rank-1 output, and SiLU-and-multiply splitting a packed row with two extracts. Tile suites 95/95; checkstyle clean. | Shape operations: broadcast , reshape , extract 6582500 | Each of these replaces a workaround that earlier kernels in this branch were using, so the tests are as much cleanup as coverage: | Operation | What it replaces | |---|---| | broadcast | the mask kernels stretched an iota row to the score shape by adding a zero tile of that shape — arithmetic doing a layout job. TestTileMasking now says broadcast , which is what it meant. | | reshape | changing a rows, 1 reduction result to another shape of the same size meant transpose — a data movement rather than a reinterpretation. It is also the only way to change rank . | | extract | splitting a packed row meant loading the same row twice through a narrower partition view; one load plus two extracts is what the kernel actually wants. | extract indexes in sub-tile units, not elements. The header does not say which, so I ran it rather than guessing: extract row, shape<1,32 , 0, 1 on a 1, 64 tile returns elements 32..63 , i.e. block-style addressing like a partition view. Exposing that backwards would have been a silent correctness trap. permute is deliberately not added: on a rank-2 tile the only non-identity permutation is the transpose that already exists, so it would be surface without capability. That is now what the docs say instead of listing it as missing. All three lower through one node and one statement emitting ct::fn tile, ct::shape<... {} , with extract 's indices appended. TestTileShapes — 4 tests - a bias row broadcast across a tile - a row-sum column broadcast to normalise each row the shape rule in the other direction - a 1, 1 reduction reshaped to rank 1 for a rank-1 output - SiLU-and-multiply splitting a packed row with two extracts from a single load Tile suites now 95 tests, all passing ; checkstyle clean. <sub One incidental: InvocationPlugin 's constructor takes Type... , so a plugin with a computed arity has to pass a single Type with the receiver first — a nested Class is one argument of the wrong type rather than a spread signature.</sub Parity status after four iterations Closed: comparisons + select + logicalAnd / Or , minimum / min / prod / fma , log / log2 / exp2 / tanh / abs / floor , int8 + double + fp8 views with the fp8 capability gate , atomicAdd , broadcast / reshape / extract . Still open, in what I read as value order: 1. num ctas / occupancy / latency hints — nothing emits them, and this is the one item where a JIT can do better than a human author rather than merely catching up: TornadoVM knows the shapes at compile time. The FP16 ladder put the tile GEMM 2.8x ahead of hand-written mma.sync but still 3.3x behind cuBLAS, and untuned launch hints are a plausible part of that gap. 2. rank ≥ 3 views — every attention kernel currently folds batch and head into the row index. 3. the rest of the atomic family sub / min / max /bitwise/ xchg / compare exchange , masked forms, order and scope choice . 4. scans partial sum / partial prod . 5. remaining scalar math tan sinh cosh atan2 isnan isinf mulhi remainder , element bitcast , logicalNot and per-operation rounding/NaN tags. | Rank-3 views 6d1e916 — the last substantive gap from the audit | view , partition , load / store and their masked forms and reshape now take three extents, so a batch or head dimension is a block index instead of arithmetic folded into the row index. The contrast is with TestTileLlmKernels bmm , the batched matmul already in this PR. Over a rank-2 view of batch m, k it has to derive the batch from the block index, multiply by the number of K tiles, and add that to every operand index. None of it is wrong; all of it is bookkeeping the view can do: // before: rank-2 view, batch folded into the row index int flatRowBlock = tc.bidX ; int batchIndex = flatRowBlock / rowBlocks; int bRowBase = batchIndex kBlocks; acc = tc.mma av.load flatRowBlock, step , bv.load bRowBase + step, columnBlock , acc ; // after: rank-3 view, batch is just an index PartitionView aView = tc.partition tc.view a, batch, m, k , 1, TILE, TILE ; Tile left = tc.reshape aView.load batchIndex, rowBlock, step , TILE, TILE ; acc = tc.mma left, right, acc ; Only the API surface and the plugin registrations were rank-limited — the LIR statements already loop over their extents and indices, so the emitter needed no change whatsoever. That was a pleasant surprise and worth recording: the rank limit was in the Java signatures, not in the lowering. Rank 3 is for addressing, not arithmetic. mma , the reductions and transpose are rank-2 operations, so a 1, M, N load is reshaped to M, N before use — the same shape NVIDIA's kernels take, loading a rank-4 tile from a B, H, S, D view and reshaping before the first multiply. TestTileRank3 — 4 tests - batched matmul at three sizes, including a batch of one where the rank-3 view is pure bookkeeping and still has to be right - per-batch scaling where every element of batch b is multiplied by b + 1 , so a kernel that ignored the batch index fails rather than merely being slower One thing the tests document, because I hit it while writing them: view extents may be runtime values, tile extents may not. Passing the row width in as a parameter was rejected at sketch time by the compile-time-shape rule — correctly, since a tile shape is part of the kernel's type — so the width is a constant while the batch and row counts stay parameters. Tile suites now 100 tests, all passing ; checkstyle clean. Where parity stands Six iterations of audit-then-close. Everything that blocked or awkwardly shaped a kernel is done: comparisons + select , the missing reductions and math, int8/double/fp8 views, atomicAdd , broadcast / reshape / extract , launch hints, rank-3 views. What is left is genuinely minor, and I would not add any of it speculatively: | Remaining | Why it can wait | |---|---| | rank ≥ 4 views | only needed to stop folding two leading dimensions; rank 3 covers batch or head | | rest of the atomic family sub / min / max /bitwise/ xchg / compare exchange , masked forms, order and scope choice | mechanical once atomicAdd exists; nothing has needed them | | scans partial sum / partial prod | no ported kernel uses a prefix operation | | tan sinh cosh atan2 isnan isinf mulhi remainder , element bitcast , logicalNot | one line of registration each, on demand | | per-operation rounding / NaN tags | defaults have matched every reference so far to the tolerances in the tests | | permute | stays out: rank-2's only non-identity permutation is the transpose that exists | The one item I would not treat as a gap to close is automatic hint derivation — the measurement two comments up shows the compiler's default beating every occupancy I forced, so that needs a model, not an implementation. The API reference had grown by appending a row per feature and had started to show it: maximum, sum and max each appeared twice, several operations added over the last few commits were missing, and nothing said which ranks an operation actually supports. Rewritten as a reference organised by category - block indices, views and memory, tile creation, elementwise arithmetic, comparisons and select, math, reductions, matrix multiply, shape and type, control flow, query methods - with the CUDA Tile spelling and a note for each. Checked mechanically rather than by eye: every public method of TileContext, PartitionView, TensorView and Tile, and every DType constant, now appears in the document. Added a "rank support at a glance" table, because the coverage is uneven on purpose and that was not written down anywhere: views, loads, stores and reshape go to rank 3; tile creation, atomicAdd, broadcast and extract stop at rank 2; mma, the reductions and transpose are rank 2 by nature. Also: - restored the note that a reduction result cannot be stored back into the view it came from same shape refuses a 32x1 tile against a 32x32 view , which was dropped when the table was reorganised; - said that every operation has a JVM implementation, and that this is why a tile kernel which fails to compile can still pass a test when bailout is enabled - check --printKernel; - noted that setBlockIndex/setBlockCount are JVM-path helpers with no device meaning verified: they are not intrinsified ; - corrected two cross-references that said "below" for sections that are above, listed logicalNot among the missing operations, and refreshed the test list, which was three suites out of thirteen. RST validated with docutils: no errors or warnings. mikepapadim /mikepapadim changed the title ~~ cuTile CUDA Tile cuTile execution path to the CUDA backend as TileContext API~~ cuTile Java CUDA Tile cuTile execution path to the CUDA backend as TileContext API Sep 14, 2026 event-31093504847 Learn more about bidirectional Unicode characters https://github.co/hiddenchars