TRSV accuray ci test fail in hygon

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=85
Running unit tests for tests/test_trsv.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 686 items

tests/test_trsv.py ...............................................................................................................................................F

=================================== FAILURES ===================================
________________________ test_accuracy_dtrsv[0-0-0-1-0] ________________________
triton.compiler.errors.UnsupportedLanguageConstruct: at 13:7:
a_ptr,
s,
LDA,
offs,
row_mask,
TRANS: tl.constexpr,
UNIT: tl.constexpr,
LOWER_EFF: tl.constexpr,
BLOCK_N: tl.constexpr,
ROWLOAD: tl.constexpr,
):
if UNIT and LOWER_EFF and TRANS == 0 and ROWLOAD:
^
chained boolean operators (A or B or C) are not supported; use parentheses to split the chain.

The above exception was the direct cause of the following exception:

n = 1, uplo = 0, trans = 0, diag = 0, lda_extra = 0

@pytest.mark.dtrsv
@pytest.mark.parametrize("n, lda_extra", TRSV_CASES)
@pytest.mark.parametrize("uplo", UPLOS)
@pytest.mark.parametrize("trans", REAL_TRANS)
@pytest.mark.parametrize("diag", DIAGS)
def test_accuracy_dtrsv(n, uplo, trans, diag, lda_extra):
check_fp64_support()
dtype = torch.float64
lda = max(1, n + lda_extra)
A_col, A_row = make_triangular(n, lda, uplo, diag, dtype, flag_blas.device)
x = trsv_randn((max(n, 1),), dtype, flag_blas.device)
ref_x = trsv_reference(uplo, trans, diag, n, A_col, lda, x, 1)
> flag_blas.dtrsv(uplo, trans, diag, n, A_row, lda, x, 1)

tests/test_trsv.py:385:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
src/flag_blas/runtime/backend/_hygon/ops/trsv.py:609: in dtrsv
return _common_dtrsv(public_uplo, public_trans, diag, n, A, lda, x, incx)
src/flag_blas/ops/level2/trsv.py:4816: in dtrsv
dtrsv_fwd_fused_kernel[inv_grid](
.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)
.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:295: in compile
module = src.make_ir(options, codegen_fns, module_map, context)
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _

self =
options = HIPOptions(num_warps=4, waves_per_eu=1, num_stages=0, 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')
codegen_fns = {'min_dot_size': . at 0x7f2e31f4ab00>}
module_map = {'triton.language.extra.libdevice': }
context =

def make_ir(self, options, codegen_fns, module_map, context):
> return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns,
module_map=module_map)
E triton.compiler.errors.CompilationError: at 25:11:
E ):
E pid = tl.program_id(0)
E offs = tl.arange(0, BLOCK_N)
E BK: tl.constexpr = BLOCK_N * CHUNK
E koffs = tl.arange(0, BK)
E row_start = pid * BLOCK_N
E row_end = tl.minimum(row_start + BLOCK_N, n)
E size = row_end - row_start
E rows = row_start + offs
E rows = tl.max_contiguous(tl.multiple_of(rows, BLOCK_N), BLOCK_N)
E row_mask = offs < size
E dinv = _dtrsv_diag_block_inv(
E ^

.venv/lib/python3.10/site-packages/triton/compiler/compiler.py:122: CompilationError
=========================== short test summary info ============================
FAILED tests/test_trsv.py::test_accuracy_dtrsv[0-0-0-1-0] - triton.compiler.errors.CompilationError: at 25:11:
):
pid = tl.program_id(0)
offs = tl.arange(0, BLOCK_N)
BK: tl.constexpr = BLOCK_N * CHUNK
koffs = tl.arange(0, BK)
row_start = pid * BLOCK_N
row_end = tl.minimum(row_start + BLOCK_N, n)
size = row_end - row_start
rows = row_start + offs
rows = tl.max_contiguous(tl.multiple_of(rows, BLOCK_N), BLOCK_N)
row_mask = offs < size
dinv = _dtrsv_diag_block_inv(
^
!!!!!!!!!!!!!!!!!!!!!!!!!! stopping after 1 failures !!!!!!!!!!!!!!!!!!!!!!!!!!!
================== 1 failed, 143 passed in 1218.81s (0:20:18) ==================
Error: E size = row_end - row_start
E rows = row_start + offs
E rows = tl.max_contiguous(tl.multiple_of(rows, BLOCK_N), BLOCK_N)
E row_mask = offs < size
E dinv = _dtrsv_diag_block_inv(
E ^

.venv/lib/python3.10/site-packages/triton/compiler/compiler.py:122: CompilationError
=========================== short test summary info ============================
FAILED tests/test_trsv.py::test_accuracy_dtrsv[0-0-0-1-0] - triton.compiler.errors.CompilationError: at 25:11:
):
pid = tl.program_id(0)
offs = tl.arange(0, BLOCK_N)
BK: tl.constexpr = BLOCK_N * CHUNK
koffs = tl.arange(0, BK)
row_start = pid * BLOCK_N
row_end = tl.minimum(row_start + BLOCK_N, n)
size = row_end - row_start
rows = row_start + offs
rows = tl.max_contiguous(tl.multiple_of(rows, BLOCK_N), BLOCK_N)
row_mask = offs < size
dinv = _dtrsv_diag_block_inv(
^
!!!!!!!!!!!!!!!!!!!!!!!!!! stopping after 1 failures !!!!!!!!!!!!!!!!!!!!!!!!!!!
================== 1 failed, 143 passed in 1218.81s (0:20:18) ==================

flag_blas import OK

Start with tests/test_trsv.py, especially test_accuracy_dtrsv, and trace the failing call into src/flag_blas/runtime/backend/_hygon/ops/trsv.py and src/flag_blas/ops/level2/trsv.py. Run the targeted TRSV accuracy test in the Hygon environment and inspect the reported Triton construct. Done means the failing case compiles and passes without regressing the surrounding TRSV tests.

Written by the indexing model from the issue text.