[Bug] [Relax] CUDA FuseTIR rejects add + strided_slice with int32 TIR buffer extents

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

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

まだ誰も着手していません。

評価

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

調査の方向性

Start by running the supplied repro.py with the default CUDA pipeline and confirm the FuseTIRInner buffer type assertion; compare the int32 and int64 controls described in the issue. Then trace the buffer type comparison in FuseTIRInner and the fusion path for R.strided_slice. Done when the int32 CUDA graph compiles and the existing controls continue to behave as reported.

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

説明

needs-triage type: bug
Expected behavior

The default CUDA pipeline should compile the valid Relax graph below. It calls an elementwise addition PrimFunc and takes a full-range slice of the result. The result should be x + y, with shape (16,) and dtype float32.

Using int32 constants for TIR buffer extents and loop indices should not make this graph fail during fusion.

Actual behavior

tvm.compile fails in FuseTIRInner with:

Check failed: (ffi::StructuralEqual()((*it).second.type(), new_buf.type())) is false:
Inconsistent buffers C and z mapped to the same relax var: z

This reproduces with the default pipeline and no disabled passes. No AlterOpImpl invocation is needed.

The following controls were checked in three independent processes:

TIR indices/extents Graph Target Result
int32 add followed by full-range slice CUDA Same FuseTIR assertion
int32 add followed by full-range slice LLVM Compiles and returns x + y
int64 add followed by full-range slice CUDA Compiles and returns x + y
int32 add without the slice CUDA Compiles and returns x + y

Each successful control was also checked with three random input pairs. Explicitly changing the TIR buffer extents, loop extent, and block-axis extent to T.int64(16) avoids the error. This suggests an index-dtype mismatch at the buffer-type comparison in FuseTIR.

Environment
TVM: 0.26.dev0+source.8f328e8 (source snapshot)
Python: 3.11.16
OS: Ubuntu 22.04.5 LTS, x86_64
LLVM: 15.0.7
CUDA toolkit: 12.4, V12.4.131
GPU: NVIDIA RTX A6000, sm_86
NVIDIA driver: 550.120
Steps to reproduce

Save the following as repro.py and run python repro.py with the TVM build above:

import tvm
from tvm.script import ir as I, relax as R, tirx as T

@I.ir_module
class Before:
    @T.prim_func(private=True, s_tir=True)
    def add(A: T.Buffer((16,), "float32"),
            B: T.Buffer((16,), "float32"),
            C: T.Buffer((16,), "float32")):
        for i in range(16):
            with T.sblock("add"):
                vi = T.axis.spatial(16, i)
                C[vi] = A[vi] + B[vi]

    @R.function
    def main(x: R.Tensor((16,), "float32"),
             y: R.Tensor((16,), "float32")) -> R.Tensor((16,), "float32"):
        R.func_attr({"global_symbol": "main"})
        with R.dataflow():
            z = R.call_tir(Before.add, (x, y), out_ty=R.Tensor((16,), "float32"))
            out = R.strided_slice(z, axes=[0], begin=[0], end=[16], strides=[1])
            R.output(out)
        return out

target = tvm.target.Target({"kind": "cuda", "arch": "sm_86",
                            "host": {"kind": "llvm", "mcpu": "generic"}})
with target, tvm.transform.PassContext(opt_level=2):
    tvm.compile(Before, target=target)

The script raises the assertion above during compilation. It was verified unchanged in three fresh Python processes.

Triage
  • needs-triage
主要言語
Python
スター
13.9k
フォーク
4k
平均マージ
8時間 2分
マージ済み PR（30日）
180

環境構築

このプロジェクトには開発コンテナ、Dockerfile、コントリビューションガイドがありません。まず README を読み、一般的な手順ははじめてのコントリビューションガイドを参照してください。

はじめの一歩

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

apache/tvm のほかの issue

apache/tvm の issue をすべて見る

似ている issue

Python の issue をもっと見る

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

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