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

Đang mở
  • Có gì không ổn?
#20,836 0 bình luận 0 reaction 0 người được giao Xem trên GitHub

Maintainer thường phản hồi trong vòng 1 ngày

Chưa có ai nhận issue này.

Đánh giá

Độ khó
4/5
Thời gian dự kiến
3-5 ngày
Mức phù hợp với người mới
48/100
Loại issue
Lỗi
Độ rõ ràng
Đặc tả rõ ràng
Mức độ hoạt động
Sôi nổi
Công nghệ
python
Lĩnh vực
backend, compilers

Hướng nghiên cứu

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.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.

Mô tả

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
Ngôn ngữ chính
Python
Star
13.9k
Fork
4k
Merge trung bình
8 giờ 2 phút
Pull request đã merge (30 ngày)
180

Chuẩn bị môi trường

Dự án này không cung cấp dev container, Dockerfile hay hướng dẫn đóng góp, nên bạn cần tự thiết lập môi trường: hãy bắt đầu từ README và xem hướng dẫn đóng góp lần đầu của chúng tôi để biết các bước chung.

Bắt đầu từ đâu

  1. Đọc hết issue, rồi đọc hướng dẫn đóng góp của dự án.
  2. Bình luận trên issue rằng bạn sẽ nhận — tránh hai người làm cùng một việc.
  3. Fork repository và làm thay đổi trên một nhánh.
  4. Mở pull request có tham chiếu số hiệu của issue.

Issue khác của apache/tvm

Tất cả issue của apache/tvm

Issue tương tự

Thêm issue về Python

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.