Skip to content

[TEST] Cover Gluon Blackwell tensor memory ops - #458

Merged
Jokeren merged 5 commits into
mainfrom
test/gluon-blackwell-tensor-memory-ops
Jun 26, 2026
Merged

Jokeren merged 5 commits into
mainfrom
test/gluon-blackwell-tensor-memory-ops

Conversation

@Jokeren

@Jokeren Jokeren commented Jun 20, 2026

Copy link
Copy Markdown
Member

Summary

Adds CPU regression coverage for Blackwell tensor-memory roundtrip operations, stacked on the Hopper WGMMA PR.

The test covers:

  • blackwell.allocate_tensor_memory
  • TensorMemoryLayout
  • tensor-memory descriptor store
  • tensor-memory descriptor load
  • explicit tensor-memory register layout conversion through get_reg_layout(num_warps=...)

The simulated output is compared against the original torch CPU input tile.

This also makes Gluon builtin patching avoid injecting _semantic into helper functions that are marked builtin but do not accept that keyword, which is needed by Blackwell tensor-memory layout helpers.

Validation

  • PYTHONPATH=/mnt/keren/triton-viz pytest tests/end_to_end/test_gluon_blackwell_tensor_memory_ops.py -q
  • PYTHONPATH=/mnt/keren/triton-viz pytest tests/end_to_end/test_gluon_blackwell_tensor_memory_ops.py tests/end_to_end/test_gluon_wgmma_ops.py tests/end_to_end/test_gluon_tma_im2col_ops.py tests/end_to_end/test_gluon_blackwell_tma_ops.py tests/end_to_end/test_gluon_tma_ops.py tests/end_to_end/test_gluon_async_copy_ops.py tests/end_to_end/test_gluon_core_ops.py tests/end_to_end/test_gluon.py tests/end_to_end/test_core.py tests/unit/test_adapters.py -q
  • pre-commit run --files triton_viz/core/simulation/gluon.py tests/end_to_end/test_gluon_blackwell_tensor_memory_ops.py

@Jokeren
Jokeren force-pushed the test/gluon-hopper-wgmma-ops branch from 2d4a108 to afd6898 Compare June 25, 2026 13:11
Base automatically changed from test/gluon-hopper-wgmma-ops to main June 25, 2026 13:22
@Jokeren
Jokeren force-pushed the test/gluon-blackwell-tensor-memory-ops branch from d199fc5 to 8a915c5 Compare June 25, 2026 14:15
@github-actions

github-actions Bot commented Jun 25, 2026 •

Copy link
Copy Markdown

Performance Benchmark

Benchmark main (min) PR (min) Change Samples
gemm 0.105s 0.105s -0.1% 20 / 20
gemm_oob 0.118s 0.117s -0.4% 20 / 20
indirect_load 0.022s 0.022s -0.7% 20 / 20
nested_loop 0.235s 0.234s -0.6% 20 / 20
block_pointer_loop_advance 0.206s 0.205s -0.6% 20 / 20
liger_jsd 0.137s 0.138s +0.2% 20 / 20
flaggems_layernorm 0.398s 0.395s -0.8% 20 / 20
swiglu 0.172s 0.172s -0.2% 20 / 20
cross_entropy 0.980s 0.985s +0.5% 20 / 20
fused_linear_jsd 0.209s 0.210s +0.5% 20 / 20
Total 2.583s 2.582s -0.0% N/A

Iterations: 1 warmup + 20 measured
Samples are shown as main / PR; long pytest benchmarks may use fewer samples.

@Jokeren
Jokeren force-pushed the test/gluon-blackwell-tensor-memory-ops branch from 2754b31 to 1c0b764 Compare June 25, 2026 14:38
@Jokeren
Jokeren force-pushed the test/gluon-blackwell-tensor-memory-ops branch from d9be516 to 389a784 Compare June 25, 2026 15:30
@Jokeren
Jokeren force-pushed the test/gluon-blackwell-tensor-memory-ops branch 8 times, most recently from bbf9c2a to 4d83822 Compare June 26, 2026 12:48
@Jokeren
Jokeren marked this pull request as ready for review June 26, 2026 13:08

@chatgpt-codex-connector chatgpt-codex-connector Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 4d83822c4a

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

Comment on lines 1394 to 1399
return gluon_blackwell.TensorMemoryLayout(
tuple(block),
int(col_stride),
[] if cga_layout is None else [list(basis) for basis in cga_layout],
bool(two_ctas),
bool(fp4_padded),
)

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Badge Preserve FP4 padding when rebuilding tensor-memory layouts

When a Gluon kernel uses a Blackwell tensor-memory layout with fp4_padded=True, _to_ir passes that flag into this builder, but this constructor call now drops it and always returns a layout constructed without the FP4 padding state. That makes the interpreter treat FP4-padded tensor-memory descriptors as ordinary layouts, so any later layout-dependent register conversion or descriptor behavior can diverge from real Gluon semantics. Please propagate fp4_padded here, using a version-aware constructor path if needed for older Triton builds.

Useful? React with 👍 / 👎.

@Jokeren
Jokeren force-pushed the test/gluon-blackwell-tensor-memory-ops branch from 4d83822 to 60c2748 Compare June 26, 2026 13:21
@Jokeren
Jokeren merged commit 044936e into main Jun 26, 2026
4 checks passed
@Jokeren
Jokeren deleted the test/gluon-blackwell-tensor-memory-ops branch June 26, 2026 14:36
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant