Hacktoberfest 2026：メンテナが10月に向けて印を付けた、オープンで初心者向けの issue。 Hacktoberfest の issue を見る

Multi-tensor swizzle kernels fail with "too many resources requested for launch" (missing __launch_bounds__)

クローズ
  • 何か問題がありますか？
#3,621 コメント 0 件 リアクション 0 件 担当者 0 名 GitHub で見る

メンテナーはふだん 2 日以内に返信

関連するプルリクエストがすでにマージされています。

  • #3622 @ravimajeti による — マージ済み

評価

難易度
2/5
見積もり時間
1〜3時間
初心者へのやさしさ
25/100
issue の種類
バグ
明瞭さ
明確に書かれている
活発さ
活発
技術スタック
cpp
領域
backend

調査の方向性

The four multi-tensor kernels live in transformer_engine/common/swizzle/swizzle.cu and are launched from launch_multi_tensor_swizzle_scaling_factors (the line in the report's stack trace); #2076 added launch_bounds(TB_DIM * TB_DIM) to the single-tensor kernels, so compare those declarations first. Done means the row, row and col variants stay within 64 registers at 1024 threads and the MultiTensorSwizzleTestSuite cases pass when run via tests/cpp/build/operator/test_operator --gtest_filter='MultiTensorSwizzleTestSuite'. Note PR #3622 is already open against this issue, so coordinate before starting.

索引モデルが issue の本文から書いたものです。

説明

Describe the bug

The multi-tensor swizzle kernels in transformer_engine/common/swizzle/swizzle.cu (multi_tensor_swizzle_{row,col}_scaling_kernel and multi_tensor_unswizzle_{row,col}_scaling_kernel) are launched with 1024 threads per block (dim3 block_size(TB_DIM, TB_DIM)), but unlike every other swizzle kernel they have no __launch_bounds__(TB_DIM * TB_DIM). Without it, ptxas may give a thread more than 64 registers. 1024 threads × more than 64 registers exceeds the 64K registers available to a block, so the launch fails with:

CUDA Error: too many resources requested for launch

#2076 added __launch_bounds__ to the single-tensor swizzle kernels; the multi-tensor kernels added a week earlier in #2019 were not included.

Whether it fails depends on the GPU arch and CUDA version, because the register count does. Measured with cuobjdump --dump-resource-usage on swizzle.cu compiled with the TE build flags (registers per thread; > 64 cannot launch with 1024 threads):

Kernel CUDA 12.8 sm_90 CUDA 12.8 sm_100 CUDA 12.8 sm_120 CUDA 13.4 sm_90 CUDA 13.4 sm_100 CUDA 13.4 sm_120
multi_tensor_swizzle_row_scaling_kernel<int4> 89 89 96 61 50 56
multi_tensor_swizzle_row_scaling_kernel<int2> 32 61 71 32 61 70
multi_tensor_swizzle_col_scaling_kernel<int4> 89 99 99 64 56 61
other multi-tensor (un)swizzle variants ≤ 56 ≤ 40 ≤ 48 ≤ 55 ≤ 40 ≤ 48

So on sm_120 it fails with both CUDA versions, and with CUDA 12.8 (the minimum for Blackwell) it should also fail on sm_90 and sm_100.

Steps/Code to reproduce bug

RTX 5090 (CC 12.0), TE built with NVTE_CUDA_ARCHS=120, CUDA 12.8:

tests/cpp/build/operator/test_operator --gtest_filter='*MultiTensorSwizzleTestSuite*'

3 cases fail, all routed to multi_tensor_swizzle_row_scaling_kernel<int4>:

[  FAILED  ] OperatorTest/MultiTensorSwizzleTestSuite.TestMultiTensorSwizzle/n2_M128_K1024_row
[  FAILED  ] OperatorTest/MultiTensorSwizzleTestSuite.TestMultiTensorSwizzle/n3_M256_K4096_row
[  FAILED  ] OperatorTest/MultiTensorSwizzleTestSuite.TestMultiTensorSwizzle/n2_M128_K8192_row
C++ exception with description "transformer_engine/common/swizzle/swizzle.cu:1360 in function launch_multi_tensor_swizzle_scaling_factors: CUDA Error: too many resources requested for launch" thrown in the test body.

(n2_M128_K1024_row is meant to cover the narrow-K kernel, which needs 128 KB of shared memory. On a 99 KiB GPU the dispatcher correctly falls back to the regular multi-tensor kernel, which then fails to launch.)

The row<int2> and col<int4> variants are not reached by the current test shapes, e.g. {2, 128, 4352, true} (row, vec_load_size = 2) and {2, 512, 4096, false} (col, vec_load_size = 4) would cover them.

Expected behavior

The multi-tensor swizzle succeeds for all shapes, like the single-tensor path.

Proposed fix

Add __launch_bounds__(TB_DIM * TB_DIM) to the four multi-tensor kernels, matching the other swizzle kernels, and add test shapes that reach the row<int2> and col<int4> variants. I'm happy to open a PR.

Environment overview

  • Environment location: Docker on vast.ai
  • Method of Transformer Engine install: from source (main at 5759fa0f)
  • Docker image: Ubuntu 24.04 with /venv/main

Environment details

  • OS version: Ubuntu 24.04
  • PyTorch version: 2.11.0+cu128
  • Python version: 3.12
  • Transformer Engine version: 2.21.0.dev0+5759fa0f
  • CUDA version: 12.8 (register table also with 13.4)
  • CUDNN version: 9.19

Device details

  • GPU model: NVIDIA GeForce RTX 5090 (CC 12.0), driver 580.95.05
主要言語
Python
スター
3.6k
フォーク
862
平均マージ
4日 1時間
マージ済み PR（30日）
57

環境構築

はじめの一歩

  1. issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
  2. 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
  3. リポジトリをフォークし、ブランチを切って変更します。
  4. issue 番号を参照したプルリクエストを送ります。

NVIDIA/TransformerEngine のほかの issue

NVIDIA/TransformerEngine の issue をすべて見る

似ている issue

Python の issue をもっと見る

新しい issue をメールで受け取る

初心者向けの GitHub issue を短くまとめたダイジェスト。