[Bug] [LLVM] FP16 asin, acos, asinh and acosh fail code generation as unknown intrinsics

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

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

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

評価

難易度
3/5
見積もり時間
1〜2日
初心者へのやさしさ
52/100
issue の種類
バグ
明瞭さ
おおむね明確
活発さ
活発
技術スタック
cpp, python
領域
compilers

調査の方向性

First run the supplied reproduction for each operator, then read src/target/intrin_rule.h, especially FloatSuffix and the DispatchPureExtern registrations. Trace how the LLVM pipeline lowers these FP16 calls, using the FP32 and explicit-promotion cases as controls. Done means direct FP16 calls to asin, acos, asinh, and acosh compile for LLVM and pass the supplied numerical checks.

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

説明

needs-triage type: bug
Expected behavior

Accepted FP16 expressions for these floating-point operators should lower for LLVM, for example by promoting the calculation to FP32 and casting the result back to FP16. The explicit promotion in the script below compiles and produces the expected FP16 values.

Actual behavior

The default LLVM compilation pipeline leaves the FP16 call unlowered and fails in code generation:

unknown intrinsic ir.Op(... name="tirx.asinh" ...)

The same error occurs for tirx.asin, tirx.acos and tirx.acosh. Each direct FP16 case failed in three fresh Python processes. Both direct FP32 and explicit FP16-input/FP32-calculation/FP16-output controls compiled in three fresh processes per operator and checked three legal inputs per process against NumPy.

No custom schedule, disabled pass, or StmtSimplify ablation is required. Inputs for acosh are in [1.25, 4]; inputs for asin and acos are in [-0.75, 0.75], so these controls stay inside their mathematical domains.

The source registrations use DispatchPureExtern. FloatSuffix in src/target/intrin_rule.h selects function names for float32 and float64, and returns an empty name for float16. This appears to explain why the FP16 intrinsic reaches LLVM code generation without a supported lowering.

Environment
  • Apache TVM 0.26.dev0+source.8f328e8 (source snapshot)
  • Ubuntu 22.04.5, Linux 6.8.0-52-generic, x86-64
  • Python 3.11.16, LLVM 15.0.7
  • Target llvm; no GPU required
Steps to reproduce

Save as repro.py and run:

# Fails with an unknown tirx.asinh intrinsic:
python repro.py asinh

# Both compile and verify three inputs:
python repro.py asinh --promote
python repro.py asinh --dtype float32

# asin, acos and acosh reproduce with the same arguments.

import argparse
import numpy as np
import tvm
from tvm import te, tirx

p = argparse.ArgumentParser()
p.add_argument("op", choices=["asin", "acos", "asinh", "acosh"])
p.add_argument("--dtype", choices=["float16", "float32"], default="float16")
p.add_argument("--promote", action="store_true")
args = p.parse_args()
x = te.placeholder((8,), dtype=args.dtype, name="x")
op = getattr(tirx, args.op)
def compute(i):
    value = x[i].astype("float32") if args.promote else x[i]
    result = op(value)
    return result.astype(args.dtype) if args.promote else result
y = te.compute((8,), compute, name="y")
mod = tvm.IRModule({"main": te.create_prim_func([x, y])})
executable = tvm.compile(mod, target="llvm")
for seed in [42, 9001, 20261103]:
    rng = np.random.default_rng(seed)
    lo, hi = (1.25, 4.0) if args.op == "acosh" else (-0.75, 0.75)
    values = rng.uniform(lo, hi, (8,)).astype(args.dtype)
    actual = tvm.runtime.empty((8,), args.dtype, tvm.cpu())
    executable["main"](tvm.runtime.tensor(values, tvm.cpu()), actual)
    expected = getattr(np, "arc" + args.op[1:])(
        values.astype("float64")
    ).astype(args.dtype)
    rtol, atol = (2e-3, 5e-4) if args.dtype == "float16" else (3e-6, 1e-7)
    np.testing.assert_allclose(actual.numpy(), expected, rtol=rtol, atol=atol)
print("three legal inputs passed", flush=True)
Triage
  • needs-triage
  • type: bug
主要言語
Python
スター
13.9k
フォーク
4k
平均マージ
8時間 5分
マージ済み PR（30日）
170

環境構築

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

はじめの一歩

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

apache/tvm のほかの issue

apache/tvm の issue をすべて見る

似ている issue

Python の issue をもっと見る

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

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