PaperlayerView exact revision on arXiv

Source-linked AI summary

Annotation-Driven Migration of CUDA Programs to Tenstorrent Blackhole

Ayumi Ohno, Shinya Takamaeda-Yamazaki

arXiv:2610.02658v1cs.DC

TL;DR

CUDA leaves physical placement and scheduling implicit, while Blackhole migration requires mapping data across cores and coordinating compute with data movement within each core. This paper presents an MLIR compiler that derives those mappings for an affine CUDA subset and exposes alternative per-core realizations through annotations. Its best measured configurations outperform the compiler’s defaults, with the preferred choices depending on surrounding policies, precision, and loop schedule.

  • Problem

    CUDA leaves physical data placement, movement, and scheduling implicit, while Blackhole migration requires deriving spatial tile dataflow and a decoupled per-core realization.

    Source: §I · p. 1

  • Method

    An MLIR-based compiler maps affine CUDA programs with per-core working sets fitting in L1 to core placement, communication, tile computation, and coordinated compute/DM kernels.

    Source: §I · p. 1

  • Results

    4.2× for BF16 gaussian and 2.1–2.2× for FP32 gaussian and jacobi_2d are the best measured speedups over the compiler’s default realizations.

    Source: §I · p. 1, §VII · p. 10

  • Takeaways & Limitations

    The measured effect of an annotation depends on surrounding policies, and the preferred realization also varies with precision and loop schedule.

    Source: §I · p. 2, §VII · p. 10

    The compiler targets an affine CUDA subset whose per-core working sets fit in local L1.

    Source: §I · p. 1

Abstract

from arXiv · show

Tenstorrent Blackhole combines distributed local memories, explicit inter-core communication, and Tensix cores decoupling data movement from computation. CUDA offers a substantial HPC software base but leaves physical data placement and scheduling largely implicit. Migrating CUDA HPC kernels to Blackhole requires spatial mapping across cores and per-core coordination of compute and data-movement kernels. We present an MLIR-based compiler deriving data placement, inter-core communication, and tile computation from a statically shaped affine CUDA subset. Declarative annotations express choices not fixed by the source, including compute/DM operation placement and streaming granularity. We evaluate it on Gaussian elimination, five-point stencil, and symmetric rank-k update. Relative to our compiler's default realizations, the best measured configurations achieve a 4.2x speedup for BF16 Gaussian and 2.1-2.2x for FP32 Gaussian and Jacobi. A choice's performance impact can reverse with surrounding policies, precision, and loop schedule, motivating comparison of alternative per-core realizations rather than independent policy selection.

I. INTRODUCTION

Blackhole migration requires deriving both cross-core dataflow and the division and coordination of work between each core’s compute and data-movement kernels. The compiler exposes these choices as annotations and evaluates alternative realizations across three CUDA workloads.

A. Tenstorrent Blackhole Architecture

  • Blackhole’s separate compute and data-movement processors make per-core coordination an additional mapping problem beyond distributing data and computation across cores.The processors exchange data through circular buffers in L1 memory.Source: §I · p. 1

D. Prior Work on HPC Acceleration for Tenstorrent

  • Prior Tenstorrent HPC studies optimize specific kernels, while TileLoom targets tile-based workloads; a general compilation path for varied affine HPC patterns remains open.The related-work discussion includes stencil and FFT studies – and TileLoom.Source: §I · p. 1, D. Prior Work on HPC Acceleration for… · pp. 2–3

B. Compiler Infrastructure

  • The compiler derives core placement, communication, tile operations, compute/DM partitioning, and buffer synchronization for an affine CUDA subset whose per-core working sets fit in L1.It emits D2M rather than the standard TTIR-to-TTNN path, which does not expose the low-level data-movement and synchronization choices used in optimized HPC kernels.Source: §I · p. 1, B. Compiler Infrastructure · p. 2

I. INTRODUCTION

  • 4.2× for BF16 gaussian and 2.1–2.2× for FP32 gaussian and jacobi_2d are the best measured speedups over the compiler’s default realizations.The preferred realization depends on surrounding policies, arithmetic precision, and loop schedule.Source: §I · pp. 1–2

C. Target Workloads and Scope

  • The evaluation covers gaussian, jacobi_2d, and syrk, with frontend adaptations for gaussian and syrk and 32 × 32 CUDA thread blocks across workloads.Evaluated per-core working sets fit within each core’s 1.5 MiB L1.Source: C. Target Workloads and Scope · p. 2

III. MAPPING SIMT TO A DECOUPLED ACCELERATOR

The compiler maps CUDA arrays and computation onto Blackhole’s distributed cores, then realizes each core’s work across compute and data-movement engines. Declarative policies expose alternatives whose performance can interact, and lowering proceeds through inspectable policy checkpoints.

A. From CUDA Constructs to Blackhole

  • Each core owns a shard in persistent L1, and affine accesses to other cores’ data identify NoC transfers under the owner-computes rule.The current scope requires per-core working sets to fit in local L1; larger working sets also need DRAM-transfer and tile-reuse planning.Source: A. From CUDA Constructs to Blackhole · p. 3
  • CUDA thread indices map to tile positions; reductions, guards, and differing access maps become tile operations, masks, and within-tile transformations.The default realization assigns arithmetic to compute and data manipulation to DM kernels, with L1 buffers and explicit synchronization for overlap.Source: A. From CUDA Constructs to Blackhole · p. 3

B. Per-Core Realization Design Space

  • A guard may use a compute-side mask or DM-side selective copy, while row/column operations and streaming granularity present additional realization choices.Column-wise DM operations can use strided scalar accesses or transpose to enable bulk transfers; finer streaming may increase overlap while also increasing synchronization and fragmenting transfers.Source: B. Per-Core Realization Design Space · p. 3
  • Declarative annotations specify policy values and application sites, allowing combinations to be lowered without handwritten kernels despite interacting performance effects.Table I summarizes the exposed policy dimensions.Source: B. Per-Core Realization Design Space · p. 3, Table I · p. 3

A. Four Compiler States and Three Policy Checkpoints

  • TilePlan, TileFlow, and Task are policy checkpoints where users inspect IR, add annotations, and resume lowering toward explicit physical assignments and D2M IR.The four states are TilePlan, TileFlow, Task, and Physical; subsequent lowering is derived from the initial specification and annotations.Source: A. Four Compiler States and Three… · pp. 3–4, Figure 2 · p. 4

B. C1: TilePlan — Tiled Iteration and Computation

C1 loop scheduling applies reduction blocking and loop reordering while accesses remain affine, enabling dependence analysis to validate the transformations.

1) Loop Scheduling:

  • Loop-schedule annotations apply reduction blocking and loop reordering after flat loads and stores are normalized into rank-2 affine accesses.The transformations preserve affine accesses and guards, allowing MLIR dependence analysis to validate loop reordering.Source: 1) Loop Scheduling · p. 4
  • The syrk schedule uses strip-mining and interchange to organize the reduction while retaining its accumulator in L1.The example’s C1 forms are shown schematically in Listing 1.Source: 1) Loop Scheduling · p. 4
  • In syrk, strip-mining groups operand reads into 1 × B blocks, while interchange exposes S × B reads and keeps an S × S accumulator in L1 across reduction blocks.K is the number of tile-level reduction steps, B the reduction-tile block size, and S the output tiles per shard dimension.Source: 1) Loop Scheduling · p. 4

2) Automatic Lowering:

The compiler lowers CUDA work into tile operations and temporal loops, then applies semantics-preserving predicate placement choices. For Jacobi’s west operand, it also presents alternative ways to apply the halo patch in a transposed basis.

2) Automatic Lowering:

  • The lowering rebases affine accesses and guards to destination coordinates, maps loads and stores to tile reads and writes, and represents conditions as tile-region predicates.Sequential and reduction loops become temporal tile loops, and host loops around kernel launches become device-side temporal loops.Source: 2) Automatic Lowering · p. 4
  • Fusion forwards Gaussian’s elimination factor directly to the matrix update, eliminating the intermediate global-memory storage.The factor m[i, t] = A[i, t]/A[t, t] would otherwise pass between the two CUDA kernels through global memory.Source: 2) Automatic Lowering · p. 4
  • Jacobi’s west-neighbor access prevents fusion because the second kernel could overwrite A before the first kernel reads it.The logical read represents the global access (i, j) → (i, j −1).Source: 2) Automatic Lowering · p. 4

3) Predicate Placement Policy:

  • C1 assigns each shared boundary predicate to one of Jacobi’s two writes, after checking affine predicates, shard coordinates, and tile accesses for equivalence.The policy requires B to be an internal intermediate used only by the two kernels; both placements preserve A’s boundary values.Source: 3) Predicate Placement Policy · pp. 4–5
  • Fig. 3 contrasts applying the west-operand halo as a column patch after transposing back with applying it as a row patch in the transposed basis.Both realizations express the local column shift as a row shift in the transposed basis.Source: Figure 3 · p. 5

1) Automatic Lowering:

The compiler derives tile-level access realizations from affine accesses and annotations that select compute or data-movement operations. These policies include Gaussian reciprocal and guard realization choices, as well as alternative halo-conversion placements for Jacobi.

1) Automatic Lowering:

  • Affine accesses are factored into source core, local tile coordinate, and within-tile map, yielding local reads or transfers and tile operations such as broadcasts and shifts.Jacobi’s west operand becomes a within-tile shift patched at the tile boundary; guarded writes become masked overlays that preserve the prior destination value on inactive lanes.Source: 1) Automatic Lowering · p. 5

2) Policies:

  • DM-side reciprocal computation evaluates Gaussian’s reciprocal at the pivot cell and broadcasts the scalar instead of computing the reciprocal over the full tile.This is the annotated arithmetic-placement realization.Source: 2) Policies · p. 5
  • A guard-realization annotation can rewrite a guarded update by sinking its select into a designated factor, with the rewrite assuming the other factor is finite on disabled lanes.The example pairs a row mask with the broadcast-column factor; a further annotation can move the fill to DM before broadcasting.Source: 2) Policies · p. 5
  • For Jacobi’s west operand, the transposed-basis annotation applies the transposed halo as a row copy before transposing back.This contrasts with applying the halo as a column patch after transposing back.Source: 2) Policies · p. 5
  • Sender-side halo conversion materializes a converted view once per time step, allowing local reads and remote deliveries to reuse it under the current global-semaphore scheme.The annotation selects whether the halo transpose is performed at the sender or receiver.Source: 2) Policies · p. 6
  • Guard realization may place each global boundary mask in compute as a masked select or in DM as a conditional lane copy; evaluation pairs it with the corresponding kernel write to reduce compute–DM handoffs.The evaluated pairings are compute realization with B ← stencil(A), and DM realization with A ← B.Source: 2) Policies · p. 6

3) Derived Optimizations:

The compiler derives communication, reuse, and transpose simplifications after applying C2 annotations.

3) Derived Optimizations:

  • C2 processing splits partly local transfers, requests rectangular deliveries as frames, identifies multicast opportunities, reuses Gaussian’s A[t, j] frame across the outer loop, and simplifies transpose pairs.A[t, j] varies with inner-loop j but not outer-loop i, so the compiler prepares its frame before the i loop; transpose simplification combines T(a)+T(b) into T(a+b) and cancels adjacent transpose pairs.Source: 3) Derived Optimizations · p. 6

D. C3: Task — Pipelining Compute and Data Movement

C3 partitions operations by engine and inserts logical L1 buffers where data crosses between compute and data movement. Streaming granularity controls producer–consumer synchronization frequency, balancing overlap against batching opportunities.

D. C3: Task — Pipelining Compute and Data Movement

  • After C2, operations are labeled compute or DM, and logical L1 buffers are emitted at transfer destinations and for DM-constructed operands.Values used only within compute remain tile values, while exchanged values require logical L1 buffers; C3 also accounts for physical 16×16-face boundaries when materializing DM-side shifts.Source: D. C3: Task — Pipelining Compute and… · p. 6
  • DM operations default to DM0, while lane annotations can assign operations to DM0 or DM1 subject to checks against unsafe synchronized-buffer sharing and conflicting persistent-L1 accesses.For Jacobi, north/south operands may use one lane and west/east operands the other because they read separate views of A.Source: D. C3: Task — Pipelining Compute and… · p. 6
  • Column streaming synchronizes once per tile and permits finer-grained overlap, whereas row streaming synchronizes once per shard row and batches tileflow.roll local-copy work across columns.Row streaming is rejected when a sequence crosses compute to DM and back within the column loop, because synchronization is required at each iteration.Source: D. C3: Task — Pipelining Compute and… · p. 6

E. Physical — Resource Planning and D2M Emission

Physicalization turns logical resources into D2M IR, allocating L1-backed circular buffers and enforcing cross-core dependencies. Dst allocation and Task emission complete the per-core and dispatch plan.

E. Physical — Resource Planning and D2M Emission

  • Circular-buffer depth defaults to one, rises to two for streamed delivery, and determines capacity and L1 allocation from buffer shape.Non-streamed delivery uses depth two only when explicitly annotated and compiler legality checks pass.Source: E. Physical — Resource Planning and… · p. 7
  • Grid-wide semaphore barriers protect persistent L1 operands by preceding remote reads and subsequent owner-core updates.The barriers enforce cross-core RAW and WAR dependences, respectively.Source: E. Physical — Resource Planning and… · p. 7
  • Dst allocation scopes each output-tile loop, assigns live values by linear scan, and spills values crossing scope boundaries to L1.A Dst-resident matmul accumulator stays in one bank across reduction; an L1-resident accumulator is emitted as linalg.generic for downstream D2M subdivision.Source: E. Physical — Resource Planning and… · p. 7
  • Each Task phase becomes d2m.generic with DM0, DM1, and compute as concurrent threads; host–DRAM transfers and layout conversion use separate dispatches.Physicalization realizes synchronization with CB reserve/push and wait/pop operations.Source: E. Physical — Resource Planning and… · pp. 6–7

F. Scope

The policy measurements show that realization and schedule choices have workload- and precision-dependent effects, while the implementation and comparisons remain bounded by specific kernel, hardware, and measurement scopes.

F. Scope

  • The implementation targets affine, statically shaped rank-2 kernels with positive extents and 32 × 32 blocks; extents and output tile grids must divide evenly, and shared memory and __syncthreads() are unsupported.Cross-core transfers cannot yet be split into streamed remote loads because barriers synchronize around entire persistent L1 buffers.Source: F. Scope · p. 7

1) Targets:

  • Blackhole tests use a p150a and A100; GPU baselines retain original iteration structures with selected optimizations, including alternating Jacobi buffers and remapped Gaussian threads.Blackhole configurations use the same spatial mapping within each workload and differ in checkpoint policies.Source: 1) Targets · p. 7

2) Measurement:

  • Device times average five measurements, excluding host submission and transfers; policy studies time d2m.generic with operands resident in L1, while A100 comparisons fold DRAM transfers into one dispatch.Peak L1 denotes the statically planned per-core footprint, including persistent buffers and temporary CBs.Source: 2) Measurement · p. 7

B. Incremental Effects of Realization Policies

  • 3.19× BF16 and 1.93× FP32: moving Gaussian masks to DM gives the largest step speedups; DM-side A[i, t]/A[t, t] adds 1.07× at FP32.Sinking the row mask also improves performance at both precisions; the FP32 DM-side arithmetic choice gives the lowest measured time.Source: 1) gaussian · p. 8
  • 1.49× BF16 and 1.29× FP32: a transposed-basis Jacobi halo patch replaces strided column patches with contiguous row copies.Sender-side conversion reduces tile transposes from five to two per output tile; a second DM lane gives 1.53× at BF16 but little benefit at FP32. Moving both boundary masks to DM slows both precisions, while row streaming with DM masks gives 1.14× at BF16; the corresponding FP32 configuration exceeds L1 capacity.Source: 2) jacobi_2d · p. 8
  • 5.57×: strip-mining syrk’s reduction groups 10 tiles into shaped deliveries and blocked matmul, reducing remote loads per core from 20,000 to 2,000.Moving the block loop outward reduces the count to 20 and provides 5.36×; depth-two delivery adds 1.27× but raises peak L1 usage to 1404 KiB.Source: 3) syrk · p. 8

C. Interactions among Realization Policies

Realization choices can reverse in performance as surrounding policies, precision, and loop schedules change. The A100 comparison shows that Blackhole and A100 encounter different bottlenecks across workloads.

C. Interactions among Realization Policies

1) Multiplication and broadcast placement in gaussian:

  • DM-side scaling increases gaussian time by 67% at BF16 but reduces it by 6% at FP32 when only row broadcast is on compute.With row streaming, the slower stage determines execution time: DM becomes the BF16 bottleneck, while moving scaling to DM relieves FP32 compute.Source: 1) Multiplication and broadcast… · p. 8, Figure 4 · p. 9
  • With only column broadcast on compute, DM-side scaling increases FP32 time by 5%.Compute-side broadcast fuses into a binary operation, while compute-side scaling fuses into a multiplication outside the column loop.Source: 1) Multiplication and broadcast… · p. 8, Figure 4 · p. 9

3) DM-lane assignment across syrk schedules:

  • Splitting syrk multicast deliveries across two DM lanes reduces time by 48% with the reduction innermost but increases it by 7% after interchange.Interchange reduces the delivery count to 20 per core, making coordination overhead outweigh the benefit of greater transaction-level concurrency.Source: 3) DM-lane assignment across syrk… · p. 9, Figure 4 · p. 9

D. Comparison with the A100

  • Blackhole is 2.94× slower than the A100 for gaussian and 2.98× slower for jacobi_2d.These comparisons include device-memory traffic but exclude host–device transfers.Source: D. Comparison with the A100 · p. 9, Table V · p. 9
  • Cross-core communication and CB synchronization account for 7% of gaussian and 13% of jacobi_2d execution time, so neither dominates.Serialized preparation adds 7% and 10%, respectively; jacobi_2d writeback adds another 7%. Most DM-side tile construction is hidden behind compute.Source: D. Comparison with the A100 · p. 9, Figure 5 · p. 9
  • In the evaluated FP32 path, operand staging and SFPU execution serialize on the math thread, limiting throughput below the arithmetic peak.Each SFPU operation serially performs arithmetic, Dst-to-register transfers, and address updates.Source: D. Comparison with the A100 · p. 9
  • For syrk, Blackhole sustains 179 TFLOP/s, or 65% of its scaled peak, when the workload is L1-resident; DRAM–L1 transfers raise time from 0.365 to 0.708 ms.The A100 completes the device-resident workload in 0.363 ms at 180 TFLOP/s, or 58% of its BF16 peak; overlapping transfers remains a primary Blackhole opportunity.Source: D. Comparison with the A100 · pp. 9–10

E. Numerical Accuracy

The paper evaluates numerical error against FP64 references and finds markedly larger errors at BF16, including a systematic Jacobi bias. It also motivates feedback-driven exploration because the annotation space is large and interacting choices affect performance.

E. Numerical Accuracy

  • At FP32, relative Frobenius error is 9.3 × 10^-7 for gaussian and 9.5 × 10^-6 for jacobi_2d after 1000 steps.The errors are measured against an FP64 reference on the gaussian offdiagonal and jacobi_2d interior.Source: E. Numerical Accuracy · p. 10
  • At BF16, error rises to 4.7×10^-3 for gaussian, 1.68 for jacobi_2d, and 1.0 × 10^-2 for syrk.The jacobi_2d bias is attributed in part to 0.2 not being exactly representable in BF16; operation and accumulation precision could be selectively promoted at a cost to Dst capacity and execution time.Source: E. Numerical Accuracy · p. 10

F. Toward Feedback-Driven Annotation Selection

  • The evaluated BF16 annotation space contains 336 gaussian and 256 jacobi_2d configurations, motivating comparison of alternative realizations rather than fixed per-annotation costs.All counted configurations executed successfully; the counts exclude some unlikely but legal choices.Source: F. Toward Feedback-Driven Annotation… · p. 10
  • The proposed feedback-driven system has an LLM propose annotation sets while the compiler checks legality and performs low-level lowering.Measured performance guides subsequent proposals.Source: F. Toward Feedback-Driven Annotation… · p. 10

VI. RELATED WORK

  • Unlike Booth and HetGPU [21], this compiler targets affine CUDA HPC kernels and preserves flexibility in their per-core realizations.The related-work discussion also distinguishes TT-Lang and Dato, which expose accelerator execution or dataflow through programming models rather than starting from existing CUDA code.Source: §VI · p. 10

VII. CONCLUSION

  • The fastest measured configurations improve performance by 4.2× for BF16 gaussian and 2.1–2.2× for FP32 gaussian and jacobi_2d against the compiler’s defaults.The conclusion also notes that annotation effects depend on surrounding policies.Source: §VII · p. 10

APPENDIX: ARTIFACT DESCRIPTION/ARTIFACT

The artifact provides compiler snapshots, executable Blackhole configurations, and A100 comparison sources, along with workflows for reproducing timings and accuracy checks. It does not include the CUDA-to-D2M frontend.

Artifact Description (AD): Artifact Identification

  • The artifact includes a Tenstorrent backend fork, MLIR snapshots, compiled Blackhole programs, and A100 comparison sources.The compiled programs run with upstream ttrt; the A100 sources use ping-pong buffers for Jacobi, coalesced Gaussian accesses, and cuBLAS for SYRK.Source: Artifact Description (AD): Artifact… · p. 12
  • The artifact cannot regenerate its snapshots from CUDA or apply new frontend annotations because it omits the CUDA-to-D2M frontend.It supports inspecting compiler representations, running measured programs, and recompiling from final D2M with the backend fork.Source: Artifact Description (AD): Artifact… · p. 12

Artifact Evaluation (AE): Reproducibility of Experiments

  • Reproduction requires a p150a and specified software revisions; the workflow covers device timings, numerical errors, and BF16 simulator comparisons.A timing measurement takes about half a minute; the A100 experiments require an A100-PCIE-40GB with the specified driver and CUDA Toolkit.Source: Artifact Evaluation (AE):… · p. 12
Loading 2610.02658v1…