論文の概要: Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4

• arxiv url: http://arxiv.org/abs/2608.10103v1
• Date: Mon, 10 Aug 2026 18:13:00 GMT
• ステータス: 翻訳完了
• システム内更新日: 2026-08-12 19:14:45.783592
• Title: Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4
• Title（参考訳）: 手書きPTXテンソルコアGEMMカーネル:NVIDIA L4のマルチ精度研究
• Abstract要約: コアカーネルは非同期データ移動から構築された低レベルのPTXパイプラインに依存している。 ほとんどのアプリケーションコードは、WMMA C++ APIを通じて間接的にCoresにアクセスする。 WMMAを手書きのPTXに置き換えるのはいつ頃なのか?
• 参考スコア（独自算出の注目度）: 0.7265954868708899
• License: http://arxiv.org/licenses/nonexclusive-distrib/1.0/
• Abstract: High-performance Tensor Core kernels rely on a low-level PTX pipeline built from asynchronous data movement with cp.async, warp-level matrix loads with ldmatrix, and matrix multiply-accumulate operations with mma.sync. However, most application code accesses Tensor Cores indirectly through the WMMA C++ API. This paper asks a focused, practical question: when does replacing WMMA with hand-written PTX actually pay off? To answer this question, we conduct a controlled, single-GPU study on an NVIDIA L4 GPU (Ada, SM89), comparing double-buffered WMMA baselines with a family of hand-written PTX GEMM kernels across FP16, INT8, and INT4 arithmetic and square problem sizes from $N=512$ to $N=8192$. Every kernel is profiled with Nsight Compute across the full metric set, and PTX speedups are reported relative to the corresponding same-precision WMMA baseline. Hand-written PTX provides no end-to-end speedup for FP16, because its instruction-level gains are offset by operand-packing overhead. In contrast, the PTX kernels achieve consistent speedups of 1.4x-1.8x for INT8, driven primarily by lower instruction counts and better global-memory coalescing, and 2.9x-4.3x for INT4, where native mma.sync.m16n8k64.s4 execution avoids the software-emulated sequence used by the WMMA path. Relative to the FP16 WMMA baseline, the best quantized kernels reach 34.4x (INT8) and 98.7x (INT4) at $N=8192$. Across these experiments, occupancy is a poor predictor of throughput. For large matrices, performance instead tracks memory-system behavior -- particularly global-load coalescing and DRAM-active cycles -- more closely than Tensor Core utilization. These results identify the precisions and operating regimes in which the additional complexity of hand-written PTX is justified.
• Abstract（参考訳）: 高性能Tensor Coreカーネルは、cp.asyncによる非同期データ移動、ldmatrixによるワープレベルの行列負荷、mma.syncによる行列乗算演算による低レベルのPTXパイプラインに依存している。 しかし、ほとんどのアプリケーションコードはWMMA C++ APIを通じて間接的にTensor Coresにアクセスする。 WMMAを手書きのPTXに置き換えるのはいつ頃なのか? この質問に答えるために、我々はNVIDIA L4 GPU(Ada, SM89)の制御されたシングルGPU研究を行い、FP16, INT8, INT4演算と2乗問題サイズをN=512$からN=8192$の2バッファのWMMAベースラインと手書きのPTX GEMMカーネルのファミリを比較した。 各カーネルは全メトリックセットにわたってNsight Computeでプロファイルされ、PTXスピードアップは対応する同一精度のWMMAベースラインに対して報告される。 手書きPTXは、命令レベルのゲインがオペランドパッケージのオーバーヘッドによってオフセットされるため、FP16のエンドツーエンドのスピードアップを提供しない。 対照的にPTXカーネルは、命令数が少なく、より優れたグローバルメモリの合体によって駆動されるINT8の1.4x-1.8xと、ネイティブの mma.sync.m16n8k64.s4 がWMMAパスで使用されるソフトウェアエミュレーションシーケンスを避けるINT4の2.9x-4.3xという一貫したスピードアップを実現している。 FP16 WMMAベースラインとは対照的に、最高の量子化されたカーネルは34.4x (INT8) と98.7x (INT4) で$N=8192$である。 これらの実験全体では、占有力はスループットの予測に乏しい。 大きな行列の場合、パフォーマンスは代わりにメモリシステムの振る舞い、特にグローバルな負荷の合体とDRAMのアクティブなサイクルを追跡します。 これらの結果は、手書きPTXの複雑さが正当化される精度と操作条件を特定する。

関連論文リスト

• Realizing Native INT8 Compute for Diffusion Transformers on Consumer GPUs: A Fused INT8 GEMM Kernel for Ideogram 4.0 [0.08599681538174887]
  拡散変圧器の後のINT8(W8A8)量子化は、速度最適化として広く展開されている。 消費者向けのAmpere GPUでは、FP8やNF4よりも遅いことがしばしばある。 われわれはこのギャップをTriton INT8 GEMMで埋める。 したがって、INT8は最も遅い変種から最も速いものへと移行する。
  論文  参考訳（メタデータ） (2026-06-12T16:19:03Z)
• APEX4: Efficient Pure W4A4 LLM Inference via Intra-SM Compute Rebalancing [7.218138615728212]
  W4A4の量子化はINT4 Coreの完全な利用を約束するが、Coresのグループの分散化オーバーヘッドは、既存のシステムを混合精度のフォールバックへと駆り立てている。 本稿では,SM内計算バランスがこのボトルネックをいかに支配するかについて,最初の系統的研究を行う。
  論文  参考訳（メタデータ） (2026-06-07T18:01:55Z)
• FairyFuse: Multiplication-Free LLM Inference on CPUs via Fused Ternary Kernels [2.339485034137561]
  提案するFairyFuseは,コモディティCPU上での乗算不要な実行を可能にする推論システムである。 ルーフライン解析により、16倍の重み圧縮により、メモリバウンドGEMVは帯域幅制限CPUの計算系にシフトすることが示された。 エンドツーエンドでは、FairyFuseは1つのIntel Xeon 8558Pで毎秒32.4トークンを獲得し、ラマを上回っている。
  論文  参考訳（メタデータ） (2026-04-22T02:42:10Z)
• Scaling DoRA: High-Rank Adaptation via Factored Norms and Fused Kernels [83.99688944263843]
  DoRA(Weight-De Low-Rank Adaptation)は、LoRAを方向から分離することで拡張する。 d_in = 8192 とランク r = 384 では、単一のモジュールのノルムは bf16 で512MB の過渡的なワーキングメモリを必要とする。 因子ノルムは、二乗ノルムを O(d_out r + r2) 中間体を通して計算可能な基底、交差、およびグラマー項に分解し、密積を除去する。
  論文  参考訳（メタデータ） (2026-03-23T17:57:24Z)
• INT v.s. FP: A Comprehensive Study of Fine-Grained Low-bit Quantization Formats [51.72056104795248]
  NvidiaのBlackwellアーキテクチャのような現代のAIハードウェアは、低精度浮動小数点(FP)フォーマットをますます受け入れている。 本稿では,FPフォーマットと整数(INT)フォーマットのトレードオフを系統的に検討する。 FPは粗粒度量子化に優れるが、きめ細かい(ブロックワイド)レベルでの比較はよりニュアンスが高い。
  論文  参考訳（メタデータ） (2025-10-29T15:11:53Z)
• Speeding Up MACE: Low-Precision Tricks for Equivarient Force Fields [51.95157731126864]
  機械学習力場は高い計算コストで正確な分子動力学(MD)を提供することができる。 この論文は、計算ボトルネックを特定し、低精度の実行ポリシーを評価することで、MACEを安価かつ高速にすることを目的としている。
  論文  参考訳（メタデータ） (2025-10-23T14:02:34Z)
• Bridging the Gap Between Promise and Performance for Microscaling FP4 Quantization [77.67818998672516]
  本研究は,MXFP4とNVFP4の学習後量子化に関する総合的研究である。 本稿では,従来のGPTQ量子化アルゴリズムの変種であるMicro-Rotated-GPTQ(MR-GPTQ)を紹介する。 MR-GPTQは最先端の精度で一致または性能が向上することを示す。
  論文  参考訳（メタデータ） (2025-09-27T09:22:21Z)
• APT-LLM: Exploiting Arbitrary-Precision Tensor Core Computing for LLM Acceleration [5.075697428779204]
  大規模言語モデル(LLM)は、AIアプリケーションに革命をもたらしたが、その膨大な計算要求は、デプロイメントとリアルタイムのパフォーマンスを著しく制限している。 これは主にGPU Coreの限定的なサポート、非効率なメモリ管理、非フレキシブルなカーネル最適化が原因である。 本稿では,任意の精度のLLM,すなわちAPT-LLMに対する包括的加速法を提案する。
  論文  参考訳（メタデータ） (2025-08-26T14:48:29Z)
• MicroMix: Efficient Mixed-Precision Quantization with Microscaling Formats for Large Language Models [3.305409455598179]
  量子化は大規模言語モデル(LLM)の推論を著しく加速する 近年のウェイトアクティベーション量子化の進歩は、主にウェイトとアクティベーションの両方をINT4フォーマットにマッピングすることに焦点を当てている。 マイクロスケーリング(MX)データフォーマットをベースとした混合精度量子化アルゴリズムと行列乗算カーネルであるMicroMixを提案する。
  論文  参考訳（メタデータ） (2025-08-04T12:22:39Z)
• MARLIN: Mixed-Precision Auto-Regressive Parallel Inference on Large Language Models [58.3342517278868]
  本稿では,Mixed-precision AutoRegressive LINearカーネルの設計について述べる。 バッチサイズは16-32までサポートでき、量子化のスピードアップが最大 (4times$) になる。 MarLINは非同期メモリアクセス、複雑なタスクスケジューリング、パイプライン化といったテクニックを組み合わせてこれを実現している。
  論文  参考訳（メタデータ） (2024-08-21T16:10:41Z)
• vTensor: Flexible Virtual Tensor Management for Efficient LLM Serving [53.972175896814505]
  大規模言語モデル(LLM)は様々なドメインで広く使われ、数百万の日次要求を処理する。 大規模言語モデル(LLM)は様々なドメインで広く使われ、数百万の日次要求を処理する。
  論文  参考訳（メタデータ） (2024-07-22T14:37:58Z)

関連論文リストは本サイト内にある論文のタイトル・アブストラクトから自動的に作成しています。

指定された論文の情報です。
本サイトの運営者は本サイト（すべての情報・翻訳含む）の品質を保証せず、本サイト（すべての情報・翻訳含む）を使用して発生したあらゆる結果について一切の責任を負いません。