diff --git a/docs/repro/l8_widen/.gitignore b/docs/repro/l8_widen/.gitignore new file mode 100644 index 0000000000..17aa483ab4 --- /dev/null +++ b/docs/repro/l8_widen/.gitignore @@ -0,0 +1 @@ +outputs/ diff --git a/docs/repro/l8_widen/README.md b/docs/repro/l8_widen/README.md new file mode 100644 index 0000000000..de17ae126a --- /dev/null +++ b/docs/repro/l8_widen/README.md @@ -0,0 +1,80 @@ +# L=8 ui8→ui16 VMI widen repro + +Fused quant→dequant needs a compact **L=8** unsigned widen (`ui8→ui16`) on the +scale-factor path. At PTOAS tag **vmi-v0.1.3**, that shape lowers through MI +and CCE, but VMI legalization leaves residual `pto.vmi.extui` and fails. Kernels +that need this step (for example fused `per_token_cast_and_cast_back`) fall back +to composing separate cast kernels or using larger L=256 tiles until L=8 VMI +widen is fixed. + +Remote: https://github.com/learning-chip/PTOAS.git (branch `zjw/l8_widen_issue`) + +## What breaks + +| Fixture | Role | vmi-v0.1.3 result | +|---------|------|-------------------| +| `fixtures/current_slow_vmi.pto` | L=256 VMI widen (works today) | PASS | +| `fixtures/desired_vmi.pto` | L=8 group-slot VMI widen (needed) | **FAIL** — illegal `extui` | +| `fixtures/target_mi.pto` | Small-L MI shape (`PAT_VL8`, `vcvt` EVEN) | PASS | +| `fixtures/reference_asc_cce.asc` | Hand-written CCE baseline | PASS | + +L=256 widen reference: `test/lit/vmi_new/vmi_to_vpto_integer_casts.pto`. + +## Layout + +``` +l8_widen/ + fixtures/ IR and ASC inputs copied from gap 06 + scripts/ env setup and compile check + outputs/ generated by check script (gitignored) +``` + +## Prerequisites + +- CANN 9.0 (`bisheng`, Ascend headers) +- PTOAS built at tag **vmi-v0.1.3** on branch `zjw/l8_widen_issue` + - needs `pto-test-opt`, `ptoas`, and `ptoas_runtime_staging` in `build/` + +## Run + +From this directory (`docs/repro/l8_widen/`): + +```bash +source scripts/env.sh +./scripts/check_l8_widen.sh +``` + +The script lowers/compiles each fixture and prints PASS or FAIL. Full log: +`outputs/check_l8_widen/compile_results.txt`. + +## Check results (2026-07-27, vmi-v0.1.3) + +``` +=== current_slow_vmi.pto (VMI L=256 ui8→ui16 lower) === +PASS: current_slow_vmi.pto lowers to VPTO (L=256 vcvt EVEN/ODD) + +=== desired_vmi.pto (VMI L=8 ui8→ui16 lower) === +FAIL: desired_vmi.pto lower exit 1 (residual pto.vmi.extui on vmi-v0.1.3) + error: failed to legalize operation 'pto.vmi.extui' that was explicitly marked illegal + error: VMI-RESIDUAL-OP: failed to convert all VMI ops/types to VPTO + +=== target_mi.pto (MI PAT_VL8 vcvt EVEN + ui16 store) === +PASS: ptoas --emit-vpto target_mi.pto +PASS: ptoas --emit-vpto-llvm-ir target_mi.pto +PASS: ptoas device object compile for target_mi.pto + +=== reference_asc_cce.asc (bisheng --cce-aicore-only) === +PASS: reference_asc_cce.asc -> non-empty .o (CCE PAT_VL8 widen baseline) +``` + +## Done when + +VMI L=8 `ui8→ui16` in `fixtures/desired_vmi.pto` lowers to the same MI ops as +`fixtures/target_mi.pto` (no residual `extui`), so fused quant round-trip can +match ASC bandwidth without compose fallback or L=256 workarounds. + +## Pin + +- PTOAS tag: **vmi-v0.1.3** +- Branch: `zjw/l8_widen_issue` +- `PTOAS_ROOT` defaults to the repo root (`../../..` from this repro directory). diff --git a/docs/repro/l8_widen/fixtures/current_slow_vmi.pto b/docs/repro/l8_widen/fixtures/current_slow_vmi.pto new file mode 100644 index 0000000000..e32826e4cd --- /dev/null +++ b/docs/repro/l8_widen/fixtures/current_slow_vmi.pto @@ -0,0 +1,20 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Current slow path: L=256 ui8→ui16 vcvt (see test/lit/vmi_new/vmi_to_vpto_integer_casts.pto). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ui8_to_ui16_l256( + %input: !pto.vmi.vreg<256xui8, #pto.vmi.layout>) + -> !pto.vmi.vreg<256xui16, #pto.vmi.layout> { + %wide = pto.vmi.vcvt %input + : !pto.vmi.vreg<256xui8, #pto.vmi.layout> + -> !pto.vmi.vreg<256xui16, #pto.vmi.layout> + return %wide : !pto.vmi.vreg<256xui16, #pto.vmi.layout> + } +} diff --git a/docs/repro/l8_widen/fixtures/current_slow_vmi.py b/docs/repro/l8_widen/fixtures/current_slow_vmi.py new file mode 100644 index 0000000000..5f15167bc0 --- /dev/null +++ b/docs/repro/l8_widen/fixtures/current_slow_vmi.py @@ -0,0 +1,27 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``current_slow_vmi.pto`` — L=256 ui8→ui16 ``vcvt``. + +Canonical IR: ``current_slow_vmi.pto``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def ui8_to_ui16_l256(): + ub = pto.alloc_tile(shape=[1, 256], dtype=pto.ui8) + off = pto.const(0, dtype=pto.index) + inp = pto.vmi.vload(ub.as_ptr(), off, size=256) + wide = pto.vmi.vcvt(inp, pto.ui16) + _ = wide + + +if __name__ == "__main__": + print(ui8_to_ui16_l256.compile().mlir_text()) diff --git a/docs/repro/l8_widen/fixtures/desired_vmi.pto b/docs/repro/l8_widen/fixtures/desired_vmi.pto new file mode 100644 index 0000000000..453480165a --- /dev/null +++ b/docs/repro/l8_widen/fixtures/desired_vmi.pto @@ -0,0 +1,20 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Idiomatic small-L widen: L=8 ui8→ui16 (needed for fused quant→dequant UE8M0). + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ui8_to_ui16_l8( + %input: !pto.vmi.vreg<8xui8, #pto.vmi.layout>) + -> !pto.vmi.vreg<8xui16, #pto.vmi.layout> { + %wide = pto.vmi.vcvt %input + : !pto.vmi.vreg<8xui8, #pto.vmi.layout> + -> !pto.vmi.vreg<8xui16, #pto.vmi.layout> + return %wide : !pto.vmi.vreg<8xui16, #pto.vmi.layout> + } +} diff --git a/docs/repro/l8_widen/fixtures/desired_vmi.py b/docs/repro/l8_widen/fixtures/desired_vmi.py new file mode 100644 index 0000000000..c0ce08b2f4 --- /dev/null +++ b/docs/repro/l8_widen/fixtures/desired_vmi.py @@ -0,0 +1,31 @@ +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +"""Python DSL form of ``desired_vmi.pto`` — L=8 ui8→ui16 widen. + +Canonical IR: ``desired_vmi.pto``. VMI legalization leaves residual ``extui``. +""" + +from ptodsl import pto + + +@pto.jit(target="a5", backend="vpto", mode="explicit") +def ui8_to_ui16_l8(): + # 32B-aligned tile; load/convert only L=8 lanes (compact SF widen). + ub = pto.alloc_tile(shape=[1, 32], dtype=pto.ui8) + off = pto.const(0, dtype=pto.index) + inp = pto.vmi.vload(ub.as_ptr(), off, size=8) + wide = pto.vmi.vcvt(inp, pto.ui16) + _ = wide + + +if __name__ == "__main__": + try: + print(ui8_to_ui16_l8.compile().mlir_text()) + except Exception as exc: # noqa: BLE001 — document the gap + print(f"desired emit/lower may fail (L=8 widen): {exc}") diff --git a/docs/repro/l8_widen/fixtures/lowered_vpto.pto b/docs/repro/l8_widen/fixtures/lowered_vpto.pto new file mode 100644 index 0000000000..e60f96b4a8 --- /dev/null +++ b/docs/repro/l8_widen/fixtures/lowered_vpto.pto @@ -0,0 +1,9 @@ +module attributes {pto.kernel_kind = #pto.kernel_kind, pto.target_arch = "a5"} { + func.func @ui8_to_ui16_l256(%arg0: !pto.vreg<256xui8>) -> (!pto.vreg<128xui16>, !pto.vreg<128xui16>) { + %0 = pto.pset_b8 "PAT_ALL" : !pto.mask + %1 = pto.vcvt %arg0, %0 {part = "EVEN"} : !pto.vreg<256xui8>, !pto.mask -> !pto.vreg<128xui16> + %2 = pto.vcvt %arg0, %0 {part = "ODD"} : !pto.vreg<256xui8>, !pto.mask -> !pto.vreg<128xui16> + return %1, %2 : !pto.vreg<128xui16>, !pto.vreg<128xui16> + } +} + diff --git a/docs/repro/l8_widen/fixtures/reference_asc_cce.asc b/docs/repro/l8_widen/fixtures/reference_asc_cce.asc new file mode 100644 index 0000000000..37249243ec --- /dev/null +++ b/docs/repro/l8_widen/fixtures/reference_asc_cce.asc @@ -0,0 +1,21 @@ +// ASC/CCE working baseline for gap 06: small-L ui8→ui16 widen (PART_EVEN). +// Compiles with bisheng --cce-aicore-only (see README). +#include "kernel_operator.h" + +__simd_vf__ inline void vf_ui8_widen_l8(__ubuf__ uint8_t* src, __ubuf__ uint16_t* dst) { + // L=8-shaped mask: the compact SF widen VMI wants on the fused path. + vector_bool v = pset_b16(PAT_VL8); + vector_u8 x; + vlds(x, src, 0, NORM); + vector_u16 y; + vcvt(y, x, v, PART_EVEN, MODE_ZEROING); + vsts(y, dst, 0, NORM_B16, v); +} + +extern "C" __global__ __vector__ void ref_gap06_ui8_widen(__gm__ uint8_t* /*gm_in*/, + __gm__ uint16_t* /*gm_out*/) { + AscendC::InitSocState(); + __ubuf__ uint8_t* ub8 = (__ubuf__ uint8_t*)0; + __ubuf__ uint16_t* ub16 = (__ubuf__ uint16_t*)(ub8 + 256); + vf_ui8_widen_l8(ub8, ub16); +} diff --git a/docs/repro/l8_widen/fixtures/target_mi.pto b/docs/repro/l8_widen/fixtures/target_mi.pto new file mode 100644 index 0000000000..22b9e16a05 --- /dev/null +++ b/docs/repro/l8_widen/fixtures/target_mi.pto @@ -0,0 +1,29 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// Desired MI: small-L (L=8) ui8→ui16 widen via vcvt PART_EVEN, then ui16 store. +// Fused quant→dequant needs this shape; VMI currently illegalizes to extui. + +module attributes {pto.target_arch = "a5", pto.kernel_kind = #pto.kernel_kind} { + func.func @ui8_to_ui16_widen_l8_mi( + %src: !pto.ptr, + %dst: !pto.ptr, + %off: index) attributes {pto.kernel} { + pto.vecscope { + %m8 = pto.pset_b8 "PAT_VL8" : !pto.mask + %in = pto.vlds %src[%off] {dist = "NORM"} + : !pto.ptr -> !pto.vreg<256xui8> + %even = pto.vcvt %in, %m8 {part = "EVEN"} + : !pto.vreg<256xui8>, !pto.mask -> !pto.vreg<128xui16> + %m16 = pto.pset_b16 "PAT_VL8" : !pto.mask + pto.vsts %even, %dst[%off], %m16 + : !pto.vreg<128xui16>, !pto.ptr, !pto.mask + } + return + } +} diff --git a/docs/repro/l8_widen/scripts/check_l8_widen.sh b/docs/repro/l8_widen/scripts/check_l8_widen.sh new file mode 100755 index 0000000000..1ebd082f30 --- /dev/null +++ b/docs/repro/l8_widen/scripts/check_l8_widen.sh @@ -0,0 +1,131 @@ +#!/usr/bin/env bash +# Compile/lower L=8 ui8→ui16 widen fixtures; print PASS/FAIL for each path. +set -euo pipefail +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +REPRO_ROOT="$(cd "${SCRIPT_DIR}/.." && pwd)" +FIXTURES="${REPRO_ROOT}/fixtures" +OUT="${REPRO_ROOT}/outputs/check_l8_widen" +LOG="${OUT}/compile_results.txt" +VMI_PASSES=(-vmi-lower-unified-to-legacy -vmi-to-vpto) + +# shellcheck disable=SC1091 +source "${SCRIPT_DIR}/env.sh" + +mkdir -p "${OUT}" +: > "${LOG}" + +BISHENG="${BISHENG:-${ASCEND_HOME_PATH}/tools/bisheng_compiler/bin/bisheng}" +ASCEND="${ASCEND_HOME_PATH}" +NPU_ARCH="${ASCEND_NPU_ARCH:-dav-3510}" +PTOAS_FLAGS=(--cann-output-version=9.0.0 --pto-arch=a5 --pto-backend=vpto) + +if ! command -v pto-test-opt >/dev/null 2>&1; then + echo "ERROR: pto-test-opt not found; build PTOAS at tag vmi-v0.1.3 first." | tee -a "${LOG}" + exit 1 +fi +if ! command -v ptoas >/dev/null 2>&1; then + echo "ERROR: ptoas not found; build PTOAS at tag vmi-v0.1.3 first." | tee -a "${LOG}" + exit 1 +fi + +echo "l8_widen compile check — $(date -u +%Y-%m-%dT%H:%M:%SZ)" | tee -a "${LOG}" +echo "PTOAS_ROOT=${PTOAS_ROOT}" | tee -a "${LOG}" +echo "pto-test-opt=$(command -v pto-test-opt)" | tee -a "${LOG}" +echo "ptoas=$(command -v ptoas)" | tee -a "${LOG}" +echo | tee -a "${LOG}" + +# --- current_slow_vmi.pto: L=256 ui8→ui16 VMI (known-good widen shape) --- +echo "=== current_slow_vmi.pto (VMI L=256 ui8→ui16 lower) ===" | tee -a "${LOG}" +CURRENT_OUT="${OUT}/current_slow_vmi_vpto.pto" +set +e +pto-test-opt "${FIXTURES}/current_slow_vmi.pto" "${VMI_PASSES[@]}" -o "${CURRENT_OUT}" \ + > "${OUT}/current_slow_vmi.log" 2>&1 +rc=$? +set -e +if [ "${rc}" -eq 0 ] && [ -s "${CURRENT_OUT}" ]; then + echo "PASS: current_slow_vmi.pto lowers to VPTO (L=256 vcvt EVEN/ODD)" | tee -a "${LOG}" +else + echo "FAIL: current_slow_vmi.pto lower exit ${rc}" | tee -a "${LOG}" + tail -20 "${OUT}/current_slow_vmi.log" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +# --- desired_vmi.pto: L=8 ui8→ui16 VMI (illegalizes to extui on vmi-v0.1.3) --- +echo "=== desired_vmi.pto (VMI L=8 ui8→ui16 lower) ===" | tee -a "${LOG}" +DESIRED_OUT="${OUT}/desired_vmi_vpto.pto" +set +e +pto-test-opt "${FIXTURES}/desired_vmi.pto" "${VMI_PASSES[@]}" -o "${DESIRED_OUT}" \ + > "${OUT}/desired_vmi.log" 2>&1 +rc=$? +set -e +if [ "${rc}" -eq 0 ] && [ -s "${DESIRED_OUT}" ]; then + echo "PASS: desired_vmi.pto lowers to VPTO (L=8 widen gap may be closed)" | tee -a "${LOG}" +else + echo "FAIL: desired_vmi.pto lower exit ${rc} (residual pto.vmi.extui on vmi-v0.1.3)" | tee -a "${LOG}" + rg -n "extui|illegal|VMI-RESIDUAL" "${OUT}/desired_vmi.log" | tail -5 | tee -a "${LOG}" || tail -15 "${OUT}/desired_vmi.log" | tee -a "${LOG}" +fi +echo | tee -a "${LOG}" + +# --- target_mi.pto: small-L MI shape (ptoas lower + device object) --- +echo "=== target_mi.pto (MI PAT_VL8 vcvt EVEN + ui16 store) ===" | tee -a "${LOG}" +TARGET_VPTO="${OUT}/target_mi_vpto.pto" +TARGET_LL="${OUT}/target_mi.ll" +TARGET_OBJ="${OUT}/target_mi.o" + +set +e +ptoas "${PTOAS_FLAGS[@]}" --emit-vpto "${FIXTURES}/target_mi.pto" -o "${TARGET_VPTO}" \ + > "${OUT}/target_mi_emit_vpto.log" 2>&1 +vpto_rc=$? +set -e +if [ "${vpto_rc}" -ne 0 ]; then + echo "FAIL: ptoas --emit-vpto target_mi.pto exit ${vpto_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/target_mi_emit_vpto.log" | tee -a "${LOG}" +else + echo "PASS: ptoas --emit-vpto target_mi.pto" | tee -a "${LOG}" + set +e + ptoas "${PTOAS_FLAGS[@]}" --emit-vpto-llvm-ir "${FIXTURES}/target_mi.pto" -o "${TARGET_LL}" \ + > "${OUT}/target_mi_emit_llvm.log" 2>&1 + llvm_rc=$? + set -e + if [ "${llvm_rc}" -ne 0 ]; then + echo "FAIL: ptoas --emit-vpto-llvm-ir target_mi.pto exit ${llvm_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/target_mi_emit_llvm.log" | tee -a "${LOG}" + else + echo "PASS: ptoas --emit-vpto-llvm-ir target_mi.pto" | tee -a "${LOG}" + set +e + ptoas "${PTOAS_FLAGS[@]}" "${FIXTURES}/target_mi.pto" -o "${TARGET_OBJ}" \ + > "${OUT}/target_mi_ptoas.log" 2>&1 + obj_rc=$? + set -e + if [ "${obj_rc}" -eq 0 ] && [ -s "${TARGET_OBJ}" ]; then + echo "PASS: ptoas device object compile for target_mi.pto" | tee -a "${LOG}" + else + echo "FAIL: ptoas device object compile exit ${obj_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/target_mi_ptoas.log" | tee -a "${LOG}" + fi + fi +fi +echo | tee -a "${LOG}" + +# --- reference_asc_cce.asc: hand-written CCE baseline --- +echo "=== reference_asc_cce.asc (bisheng --cce-aicore-only) ===" | tee -a "${LOG}" +REF_OBJ="${OUT}/reference_asc_cce.o" +set +e +"${BISHENG}" -O2 -fPIC -std=c++17 --npu-arch="${NPU_ARCH}" --cce-aicore-only -c \ + "${FIXTURES}/reference_asc_cce.asc" -o "${REF_OBJ}" \ + -I"${ASCEND}/include" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/impl" \ + -I"${ASCEND}/compiler/tikcpp/tikcfw/interface" \ + > "${OUT}/reference_asc_cce.log" 2>&1 +ref_rc=$? +set -e +if [ "${ref_rc}" -eq 0 ] && [ -s "${REF_OBJ}" ]; then + echo "PASS: reference_asc_cce.asc -> non-empty .o (CCE PAT_VL8 widen baseline)" | tee -a "${LOG}" +else + echo "FAIL: reference_asc_cce.asc compile exit ${ref_rc}" | tee -a "${LOG}" + tail -20 "${OUT}/reference_asc_cce.log" | tee -a "${LOG}" +fi + +echo | tee -a "${LOG}" +echo "Full log: outputs/check_l8_widen/compile_results.txt" | tee -a "${LOG}" diff --git a/docs/repro/l8_widen/scripts/env.sh b/docs/repro/l8_widen/scripts/env.sh new file mode 100755 index 0000000000..99af46705e --- /dev/null +++ b/docs/repro/l8_widen/scripts/env.sh @@ -0,0 +1,28 @@ +#!/usr/bin/env bash +# CANN + PTOAS env for l8_widen compile checks. +set -euo pipefail +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +REPRO_ROOT="$(cd "${SCRIPT_DIR}/.." && pwd)" +ASCEND_HOME_PATH="${ASCEND_HOME_PATH:-${ASCEND_TOOLKIT_HOME:-/usr/local/Ascend/cann-9.0.0}}" + +# PTOAS repo root: l8_widen -> repro -> docs -> repo (three levels up). +PTOAS_ROOT="${PTOAS_ROOT:-$(cd "${REPRO_ROOT}/../../.." && pwd)}" + +# shellcheck disable=SC1091 +source "${ASCEND_HOME_PATH}/bin/setenv.bash" +if [ -n "${PTOAS_ROOT}" ] && [ -f "${PTOAS_ROOT}/scripts/ptoas_env.sh" ]; then + export PTOAS_ENV_SKIP_SMOKE_TEST="${PTOAS_ENV_SKIP_SMOKE_TEST:-1}" + # shellcheck disable=SC1091 + source "${PTOAS_ROOT}/scripts/ptoas_env.sh" + for _tool_dir in \ + "${PTOAS_ROOT}/build/tools/pto-test-opt" \ + "${PTOAS_ROOT}/build/tools/ptoas"; do + if [ -d "${_tool_dir}" ]; then + export PATH="${_tool_dir}:${PATH}" + fi + done +fi + +export ASCEND_NPU_ARCH="${ASCEND_NPU_ARCH:-dav-3510}" +export PTOAS_HOST_TARGET_CPU="${PTOAS_HOST_TARGET_CPU:-znver3}" +unset CCC_OVERRIDE_OPTIONS || true