GEMV fail

Run source .venv/bin/activate
WARNING: /opt/hyhal/lib/cmake/rocm_smi doesn't exist, please install the latest driver(version >= 6.3.15)
The rocm_smi_lib will be used in default
WARNING: /opt/hyhal/lib/cmake/rocm_smi doesn't exist, please install the latest driver(version >= 6.3.15)
The rocm_smi_lib will be used in default
Sourced Hygon DTK environment: /opt/dtk-26.04/env.sh
LD_LIBRARY_PATH=/opt/dtk-26.04/dcc/gcvm/lib:/opt/dtk-26.04/hip/lib:/opt/dtk-26.04/llvm/lib:/opt/dtk-26.04/lib:/opt/dtk-26.04/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:/opt/dtk-26.04/dushmem/lib:/opt/dtk-26.04/opencl/lib:/opt/dtk-26.04/.hyhal/rocm_smi/lib:/opt/dtk-26.04/dcc/gcvm/lib:/opt/dtk-26.04/hip/lib:/opt/dtk-26.04/llvm/lib:/opt/dtk-26.04/lib:/opt/dtk-26.04/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:/opt/dtk-26.04/dushmem/lib:/opt/dtk-26.04/opencl/lib:/opt/dtk-26.04/.hyhal/rocm_smi/lib:/opt/hyhal/lib/rocprofiler:/opt/hyhal/lib/criu/:/opt/hyhal/lib
Environment configured for vendor: hygon (BLAS_VENDOR=hygon)
PR_ID=107
Running unit tests for tests/test_gemv.py
[correctness] reference backend: hipBLAS (--ref cuda)
============================= test session starts ==============================
platform linux -- Python 3.10.20, pytest-9.1.1, pluggy-1.6.0
rootdir: /home/secure/actions-runner/_work/FlagBLAS/FlagBLAS
configfile: pytest.ini
plugins: md-report-0.8.0
collected 1644 items

tests/test_gemv.py .....ssss............ssss......loc("/home/secure/actions-runner/_work/FlagBLAS/FlagBLAS/src/flag_blas/ops/level2/gemv.py":1252:38): error: operand #5 does not dominate this use
F

=================================== FAILURES ===================================
_________________ test_zgemv_splitk_stride_tail[0.0-1-4096-0] __________________

m = 1, n = 4096, trans = 0, beta = 0.0

@pytest.mark.zgemv
@HYGON_ONLY
@pytest.mark.parametrize(
"m,n,trans",
[
(1, 4096, CUBLAS_OP_N),
(4096, 1, CUBLAS_OP_T),
(4096, 1, CUBLAS_OP_C),
],
)
@pytest.mark.parametrize("beta", [0.0, 0.5 + 0.25j])
def test_zgemv_splitk_stride_tail(m, n, trans, beta):
dtype, alpha, incx, incy = torch.complex128, 1.5 + 0.5j, 2, 2
sentinel_x = 37.0 + 11.0j
sentinel_y = 41.0 + 13.0j
A = torch.randn((m, n), dtype=dtype, device=flag_blas.device)
len_x, len_y = (n, m) if trans == CUBLAS_OP_N else (m, n)
x = torch.full(
(1 + (len_x - 1) * incx + 2,),
sentinel_x,
dtype=dtype,
device=flag_blas.device,
)
x_values = torch.randn(len_x, dtype=dtype, device=flag_blas.device)
x[0 : len_x * incx : incx] = x_values
y = torch.full(
(1 + (len_y - 1) * incy + 2,),
sentinel_y,
dtype=dtype,
device=flag_blas.device,
)
y_values = torch.randn(len_y, dtype=dtype, device=flag_blas.device)
y[0 : len_y * incy : incy] = y_values
if beta == 0.0:
y[0 : len_y * incy : incy] = complex(float("nan"), float("nan"))

if trans == CUBLAS_OP_N:
op_A = A
elif trans == CUBLAS_OP_T:
op_A = A.T
else:
op_A = A.mH
expected = to_reference(y.clone(), upcast=True)
product = op_A.to("cpu") @ x_values.to("cpu")
old_y = y_values.to("cpu")
expected[0 : len_y * incy : incy] = alpha * product + beta * old_y

> flag_blas.zgemv(trans, m, n, alpha, A, n, x, incx, beta, y, incy)

tests/test_gemv.py:927:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
src/flag_blas/runtime/backend/_hygon/ops/gemv.py:398: in zgemv
return common_zgemv(
src/flag_blas/ops/level2/gemv.py:2715: in zgemv
zgemv_n_splitk_kernel[grid_sk](
.venv/lib/python3.10/site-packages/triton/runtime/jit.py:339: in
return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)
src/flag_blas/utils/libentry.py:743: in run
kernel = self.fn.run(*args, **kwargs)
src/flag_blas/utils/libentry.py:471: in run
best_config, timings = self.policy(
src/flag_blas/utils/libentry.py:418: in policy
return policy_impl(fn, configs, args, kwargs)
src/flag_blas/utils/libentry.py:572: in default_policy
timings: Dict[triton.Config, float] = {
src/flag_blas/utils/libentry.py:573: in
config: bench_fn(config) for config in configs
src/flag_blas/utils/libentry.py:467: in bench
ret = self._bench(*args, config=config, **kwargs)
.venv/lib/python3.10/site-packages/triton/runtime/autotuner.py:131: in _bench
return do_bench(kernel_call, warmup=self.num_warmups, rep=self.num_reps, quantiles=(0.5, 0.2, 0.8))
.venv/lib/python3.10/site-packages/triton/testing.py:119: in do_bench
fn()
.venv/lib/python3.10/site-packages/triton/runtime/autotuner.py:115: in kernel_call
self.fn.run(
.venv/lib/python3.10/site-packages/triton/runtime/jit.py:680: in run
kernel = self.compile(
.venv/lib/python3.10/site-packages/triton/compiler/compiler.py:301: in compile
next_module = compile_ir(module, metadata)
.venv/lib/python3.10/site-packages/triton/backends/hcu/compiler.py:355: in
stages["ttgir"] = lambda src, metadata: self.make_ttgir(src, metadata, options)
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _

mod =
metadata = {'hash': '2708eeff54c02ca70fb477052701c657244f6c6e9988e74a968230b6eb68617b', 'target': GPUTarget(backend='hip', arch='gfx936', warp_size=64), 'num_warps': 4, 'waves_per_eu': 1, ...}
options = HIPOptions(num_warps=4, waves_per_eu=1, num_stages=2, reorder_instr=1, num_ctas=1, num_ldmatrixes=0, enable_mmacfuse=0...=1, allow_flush_denorm=False, max_num_imprecise_acc_default=0, backend_name='hip', instruction_sched_variant='default')

@staticmethod
def make_ttgir(mod, metadata, options):
pm = ir.pass_manager(mod.context)
pm.enable_debug()
passes.ttir.add_convert_to_ttgpuir(pm, f"hip:{options.arch}", options.num_warps, options.warp_size,
options.num_ctas)
pm.run(mod)
pm = ir.pass_manager(mod.context)
pm.enable_debug()
passes.ttgpuir.add_coalesce(pm)
passes.ttgpuir.add_remove_layout_conversions(pm)
passes.ttgpuir.add_optimize_thread_locality(pm)
hcu.passes.ttgpuir.add_accelerate_matmul(pm, options.arch, options.matrix_instr_nonkdim, options.kpack,
options.num_ldmatrixes, options.enable_mmacfuse)
passes.ttgpuir.add_remove_layout_conversions(pm)
hcu.passes.ttgpuir.add_hcu_accelerate_flash_attention(pm)
if options.reorder_instr != 0:
hcu.passes.ttgpuir.add_reorder_instructions(pm)
passes.common.add_cse(pm)
# if options.arch == "gfx936" or options.arch == "gfx928":
# hcu.passes.ttgpuir.add_optimize_epilogue(pm)
passes.ttgpuir.add_optimize_dot_operands(pm, True)
use_global_to_local = os.getenv("TRITON_ENABLE_GLOBAL_TO_LOCAL", "0") == "1"
use_muti_pipeline = os.getenv("TRITON_ENABLE_MUTI_PIPELINE", "0") == "1"
if use_global_to_local:
if use_muti_pipeline:
hcu.passes.ttgpuir.add_hcu_stream_pipeline(pm, options.num_stages, 0, 0, True)
else:
passes.ttgpuir.add_pipeline(pm, options.num_stages)
else:
use_new_pipeliner = os.getenv("TRITON_HIP_USE_NEW_STREAM_PIPELINE", "1") == "1"
if hcu.has_matrix_core_feature(options.arch):
if use_new_pipeliner:
# In the old pipeliner we only support num_stages = 0/1, which means something
# different than the NVIDIA side. In the new pipeliner we unify the num_stages
# interpretation. Default to use 2 stages if not explicitly set.
num_stages = options.num_stages if options.num_stages != 0 else 2
hcu.passes.ttgpuir.add_stream_pipelinev2(pm, num_stages)
else:
if options.num_stages == 0:
hcu.passes.ttgpuir.add_stream_pipeline(pm)
passes.common.add_canonicalizer(pm)
hcu.passes.ttgpuir.insert_instruction_sched_hints(pm)
passes.ttgpuir.add_optimize_dot_operands(pm, True)
passes.ttgpuir.add_remove_layout_conversions(pm)
passes.ttgpuir.add_reduce_data_duplication(pm)
if use_global_to_local:
pass
else:
if use_new_pipeliner or options.num_stages != 0:
hcu.passes.ttgpuir.add_reorder_instructions(pm)
hcu.passes.ttgpuir.add_canonicalize_pointers(pm)
if os.environ.get("HCUGCN_USE_BUFFER_OPS", "0") == "1":
#hcu.passes.ttgpuir.add_canonicalize_pointers(pm)
passes.common.add_canonicalizer(pm)
hcu.passes.ttgpuir.add_convert_to_buffer_ops(pm)
passes.common.add_canonicalizer(pm)
passes.common.add_cse(pm)
passes.common.add_symbol_dce(pm)
if use_muti_pipeline:
hcu.passes.ttgpuir.add_update_async_wait_count(pm)
> pm.run(mod)
E RuntimeError: PassManager::run failed

.venv/lib/python3.10/site-packages/triton/backends/hcu/compiler.py:228: RuntimeError
=========================== short test summary info ============================
FAILED tests/test_gemv.py::test_zgemv_splitk_stride_tail[0.0-1-4096-0] - RuntimeError: PassManager::run failed
!!!!!!!!!!!!!!!!!!!!!!!!!! stopping after 1 failures !!!!!!!!!!!!!!!!!!!!!!!!!!!
=================== 1 failed, 23 passed, 8 skipped in 9.97s ====================
Error: passes.ttgpuir.add_remove_layout_conversions(pm)
passes.ttgpuir.add_reduce_data_duplication(pm)
if use_global_to_local:
pass
else:
if use_new_pipeliner or options.num_stages != 0:
hcu.passes.ttgpuir.add_reorder_instructions(pm)
hcu.passes.ttgpuir.add_canonicalize_pointers(pm)
if os.environ.get("HCUGCN_USE_BUFFER_OPS", "0") == "1":
#hcu.passes.ttgpuir.add_canonicalize_pointers(pm)
passes.common.add_canonicalizer(pm)
hcu.passes.ttgpuir.add_convert_to_buffer_ops(pm)
passes.common.add_canonicalizer(pm)
passes.common.add_cse(pm)
passes.common.add_symbol_dce(pm)
if use_muti_pipeline:
hcu.passes.ttgpuir.add_update_async_wait_count(pm)
> pm.run(mod)
E RuntimeError: PassManager::run failed

.venv/lib/python3.10/site-packages/triton/backends/hcu/compiler.py:228: RuntimeError
=========================== short test summary info ============================
FAILED tests/test_gemv.py::test_zgemv_splitk_stride_tail[0.0-1-4096-0] - RuntimeError: PassManager::run failed
!!!!!!!!!!!!!!!!!!!!!!!!!! stopping after 1 failures !!!!!!!!!!!!!!!!!!!!!!!!!!!
=================== 1 failed, 23 passed, 8 skipped in 9.97s ====================

flag_blas import OK
Error: Process completed with exit code 1.

intorduced by pr107

Start by running the failing parametrized case in tests/test_gemv.py: test_zgemv_splitk_stride_tail[0.0-1-4096-0]. Then inspect src/flag_blas/ops/level2/gemv.py around line 1252 and the zgemv path around line 2715, with PR107 as context for the regression. Done means this case completes without the Triton compiler error and the GEMV tests pass.

Written by the indexing model from the issue text.