已关闭
[Bug-Report|缺陷反馈]: master 分支在 CANN-9.0.0 社区版编译 Ascend950 算子失败 #313
xujiachen8创建于  7月16日关闭于  8月5日
xujiachen8成员
7月16日 创建

Thanks for sending an issue! Please fill in the following template to help quickly solve your problem.

Describe the current behavior / 问题描述 (Mandatory / 必填)

Ascend950 环境:
当前 master 分支使用 CANN-9.0.0 社区正式版(hiascend官网取包: https://www.hiascend.com/cann/download )编译会失败,看日志还是 tensor_api 仓引入的问题;
测试 cann0708-weekly 构建包可以正常编译,算子功能正常。构建包地址:https://ascend.devcloud.huaweicloud.com/artifactory/cann-run-mirror/software/master/20260708000326093/

image.png

Environment / 环境信息 (Mandatory / 必填)

硬件:Ascend950PR
image.png

软件:
Ubuntu 24.04 X86
CANN-9.0.0 社区版
image.png

Steps to reproduce the issue / 重现步骤 (Mandatory / 必填)

bash build.sh [--pkg] --soc=ascend950 [--ops=sbmv,spmv,symv,trmv,tbmv --run]

Describe the expected behavior / 预期结果 (Mandatory / 必填)

编译通过

(xjc) root@szblue-c01-22-3u:~/repos/gitcode/ops-blas_dev# bash build.sh --soc=ascend950 --ops=sbmv,spmv,symv,trmv,tbmv --run
[INFO] ctrmv: skipped (no arch35 implementation)
BUILD_OPS=ssbmv,sspmv,ssymv,strmv,stbmv, RUN_TEST=ON, ENABLE_PACKAGE=FALSE
-- The C compiler identification is GNU 13.3.0
-- The CXX compiler identification is GNU 13.3.0
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Check for working C compiler: /usr/bin/cc - skipped
-- Detecting C compile features
-- Detecting C compile features - done
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /usr/bin/c++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- SOC_VERSION=ascend950, NPU_ARCH=dav-3510
-- SOC_ARCH_DIRS=arch35
-- ASC_DEVKIT_MAJOR=9, ASC_DEVKIT_MINOR=0, ENABLE_BLASLT_MXFP8=FALSE, ENABLE_BLAS_TRMM=FALSE, ENABLE_BLAS_DGMM=FALSE
-- Skipping blasLt MXFP8/MXFP4 (requires asc-devkit >= 9.1, current 9.0)
-- System processer: x86_64
-- CMAKE_ASC_COMPILER: /home/xjc/CANN-9.0.0-Community/cann-9.0.0/bin/bisheng
-- ASCEND_CANN_PACKAGE_LINUX_PATH: /home/xjc/CANN-9.0.0-Community/cann-9.0.0/x86_64-linux
-- CMAKE_ASC_LLD_LINKER: /home/xjc/CANN-9.0.0-Community/cann-9.0.0/x86_64-linux/ccec_compiler/bin/ld.lld
-- Skipping blas trmm/strmm (requires asc-devkit >= 9.1, current 9.0)
-- Skipping blas dgmm/sdgmm (requires asc-devkit >= 9.1, current 9.0)
-- Configuring done (2.9s)
-- Generating done (0.1s)
-- Build files have been written to: /home/xjc/repos/gitcode/ops-blas_dev/build
[  1%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_handle.cpp.o
[  2%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_logger.cpp.o
[  2%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_layout.cpp.o
[  3%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_matmul.cpp.o
[  3%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_matmul_algo.cpp.o
[  3%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_transform.cpp.o
[  5%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_transform_desc.cpp.o
[  5%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_matmul_preference.cpp.o
[  7%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/internal/src/aclblaslt_algo_heuristic.cpp.o
[  7%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_version.cpp.o
[  7%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/api/aclblaslt_matmul_desc.cpp.o
[  8%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/internal/src/aclblaslt_handle_impl.cpp.o
[  8%] Building ASC object CMakeFiles/ops_blas.dir/blas/common/helper/aclblas_auxiliary.cpp.o
[  9%] Building ASC object CMakeFiles/ops_blas.dir/blas/asum/arch35/sasum_host.cpp.o
[ 10%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/matmul/engine/matmul_engine.cpp.o
[ 11%] Building ASC object CMakeFiles/ops_blas.dir/blas/common/helper/aclblas_logger_manager.cpp.o
[ 12%] Building ASC object CMakeFiles/ops_blas.dir/blas/copy/arch35/scopy_host.cpp.o
[ 15%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/internal/src/aclblaslt_matmul_problem.cpp.o
[ 15%] Building ASC object CMakeFiles/ops_blas.dir/blas/axpy/arch35/saxpy_host.cpp.o
[ 15%] Building ASC object CMakeFiles/ops_blas.dir/blas/asum/arch35/sasum_kernel.cpp.o
[ 15%] Building ASC object CMakeFiles/ops_blasLt.dir/blasLt/matmul/fp32/arch35/matmul_fp32_kernel.cpp.o
[ 15%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/matmul/epilogue/arch35/epilogue_alpha_beta_host.cpp.o
[ 17%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/internal/src/aclblaslt_capsule.cpp.o
[ 17%] Building ASC object CMakeFiles/ops_blas.dir/blas/axpy/arch35/saxpy_kernel.cpp.o
[ 18%] Building ASC object CMakeFiles/ops_blas.dir/blas/axpy_ex/arch35/axpy_ex_kernel.cpp.o
[ 19%] Building ASC object CMakeFiles/ops_blas.dir/blas/dotex/arch35/dotex_kernel.cpp.o
[ 19%] Building ASC object CMakeFiles/ops_blas.dir/blas/axpy_ex/arch35/axpy_ex_host.cpp.o
[ 19%] Building ASC object CMakeFiles/ops_blasLt.dir/blasLt/matmul/epilogue/arch35/epilogue_alpha_beta_kernel.cpp.o
[ 21%] Building ASC object CMakeFiles/ops_blasLt.dir/blasLt/matrix_transform/arch35/matrix_transform_kernel.cpp.o
[ 21%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/matrix_transform/arch35/matrix_transform_host.cpp.o
[ 21%] Building ASC object CMakeFiles/ops_blas.dir/blas/dotex/arch35/dotex_host.cpp.o
[ 22%] Building ASC object CMakeFiles/ops_blas.dir/blas/dot/arch35/sdot_host.cpp.o
[ 23%] Building ASC object CMakeFiles/ops_blas.dir/blas/gbmv/arch35/sgbmv_host.cpp.o
[ 24%] Building ASC object CMakeFiles/ops_blas.dir/blas/dot/arch35/sdot_kernel.cpp.o
[ 24%] Building ASC object CMakeFiles/ops_blas.dir/blas/copy/arch35/scopy_kernel.cpp.o
[ 25%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/internal/src/aclblaslt_logger_impl.cpp.o
[ 25%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_batched_ex/arch35/gemm_batched_ex_host.cpp.o
[ 26%] Building ASC object CMakeFiles/ops_blas.dir/blas/gels_batched/arch35/sgels_batched_host.cpp.o
[ 26%] Building ASC object CMakeFiles/ops_blas.dir/blas/gbmv/arch35/sgbmv_kernel.cpp.o
[ 27%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm/arch35/gemm_ex_host.cpp.o
[ 27%] Building ASC object CMakeFiles/ops_blas.dir/blas/gels_batched/arch35/sgels_batched_kernel.cpp.o
[ 28%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_ex/arch35/sgemm_ex_kernel.cpp.o
[ 29%] Building CXX object CMakeFiles/ops_blasLt.dir/blasLt/matmul/fp32/arch35/matmul_fp32_host.cpp.o
[ 29%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_batched_ex/arch35/gemm_batched_ex_kernel.cpp.o
[ 30%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm/arch35/gemm_ex_kernel.cpp.o
[ 31%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_ex/arch35/sgemm_ex_host.cpp.o
[ 31%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_grouped_batched/arch35/gemm_grouped_batched_host.cpp.o
[ 32%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemv_batched/arch35/gemv_batched_host.cpp.o
[ 32%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_grouped_batched_ex/arch35/gemm_grouped_batched_ex_host.cpp.o
[ 34%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_grouped_batched_ex/arch35/gemm_grouped_batched_ex_epilogue_kernel.cpp.o
[ 34%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_grouped_batched_ex/arch35/gemm_grouped_batched_ex_kernel.cpp.o
[ 34%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemv/arch35/sgemv_kernel.cpp.o
[ 35%] Building ASC object CMakeFiles/ops_blas.dir/blas/geqrf_batched/arch35/sgeqrf_batched_host.cpp.o
[ 36%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemv_strided_batched/arch35/gemv_strided_batched_kernel.cpp.o
[ 37%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemm_grouped_batched/arch35/gemm_grouped_batched_kernel.cpp.o
[ 37%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemv_strided_batched/arch35/gemv_strided_batched_host.cpp.o
[ 38%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemv_batched/arch35/gemv_batched_kernel.cpp.o
[ 39%] Building ASC object CMakeFiles/ops_blas.dir/blas/gemv/arch35/sgemv_host.cpp.o
[ 39%] Building ASC object CMakeFiles/ops_blas.dir/blas/geqrf_batched/arch35/sgeqrf_batched_kernel.cpp.o
[ 40%] Building ASC object CMakeFiles/ops_blas.dir/blas/ger/arch35/sger_host.cpp.o
[ 42%] Building ASC object CMakeFiles/ops_blas.dir/blas/getri_batched/arch35/getri_batched_host.cpp.o
[ 42%] Building ASC object CMakeFiles/ops_blas.dir/blas/ger/arch35/sger_kernel.cpp.o
[ 42%] Building ASC object CMakeFiles/ops_blas.dir/blas/iamin/arch35/isamin_kernel.cpp.o
[ 42%] Building ASC object CMakeFiles/ops_blas.dir/blas/getri_batched/arch35/getri_batched_kernel.cpp.o
[ 43%] Building ASC object CMakeFiles/ops_blas.dir/blas/getrf_batched/arch35/sgetrf_batched_kernel.cpp.o
[ 43%] Building ASC object CMakeFiles/ops_blas.dir/blas/getrf_batched/arch35/sgetrf_batched_host.cpp.o
[ 45%] Building ASC object CMakeFiles/ops_blas.dir/blas/getrs_batched/arch35/sgetrs_batched_kernel.cpp.o
[ 45%] Building ASC object CMakeFiles/ops_blas.dir/blas/iamax/arch35/isamax_kernel.cpp.o
[ 46%] Building ASC object CMakeFiles/ops_blas.dir/blas/getrs_batched/arch35/sgetrs_batched_host.cpp.o
[ 46%] Building ASC object CMakeFiles/ops_blas.dir/blas/iamax/arch35/isamax_host.cpp.o
[ 47%] Building ASC object CMakeFiles/ops_blas.dir/blas/matinv_batched/arch35/init_ptr_array_kernel.cpp.o
[ 48%] Building ASC object CMakeFiles/ops_blas.dir/blas/iamin/arch35/isamin_host.cpp.o
[ 49%] Building ASC object CMakeFiles/ops_blas.dir/blas/matinv_batched/arch35/copy_atoa_kernel.cpp.o
[ 51%] Building ASC object CMakeFiles/ops_blas.dir/blas/nrm2/arch35/snrm2_ex_kernel.cpp.o
[ 51%] Building ASC object CMakeFiles/ops_blas.dir/blas/matinv_batched/arch35/matinv_batched_host.cpp.o
[ 51%] Building ASC object CMakeFiles/ops_blas.dir/blas/nrm2/arch35/snrm2_kernel.cpp.o
[ 52%] Building ASC object CMakeFiles/ops_blas.dir/blas/nrm2/arch35/snrm2_ex_host.cpp.o
[ 53%] Building ASC object CMakeFiles/ops_blas.dir/blas/rot/arch35/srot_host.cpp.o
[ 53%] Building ASC object CMakeFiles/ops_blas.dir/blas/nrm2/arch35/snrm2_host.cpp.o
[ 53%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotg/arch35/srotg_host.cpp.o
[ 54%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotex/arch35/rotex_kernel.cpp.o
[ 55%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotg/arch35/srotg_kernel.cpp.o
[ 56%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotex/arch35/rotex_host.cpp.o
[ 57%] Building ASC object CMakeFiles/ops_blas.dir/blas/rot/arch35/srot_kernel.cpp.o
[ 57%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotm/arch35/srotm_host.cpp.o
[ 58%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotm/arch35/srotm_kernel.cpp.o
[ 60%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotmg/arch35/srotmg_kernel.cpp.o
[ 60%] Building ASC object CMakeFiles/ops_blas.dir/blas/rotmg/arch35/srotmg_host.cpp.o
[ 60%] Building ASC object CMakeFiles/ops_blas.dir/blas/sbmv/arch35/ssbmv_kernel.cpp.o
[ 60%] Building ASC object CMakeFiles/ops_blas.dir/blas/sbmv/arch35/ssbmv_host.cpp.o
[ 60%] Building ASC object CMakeFiles/ops_blas.dir/blas/scal/arch35/sscal_kernel.cpp.o
[ 61%] Building ASC object CMakeFiles/ops_blas.dir/blas/scal/arch35/sscal_host.cpp.o
[ 62%] Building ASC object CMakeFiles/ops_blas.dir/blas/spmv/arch35/sspmv_kernel.cpp.o
[ 64%] Building ASC object CMakeFiles/ops_blas.dir/blas/scalex/arch35/scalex_host.cpp.o
[ 64%] Building ASC object CMakeFiles/ops_blas.dir/blas/scalex/arch35/scalex_kernel.cpp.o
[ 64%] Building ASC object CMakeFiles/ops_blas.dir/blas/spmv/arch35/sspmv_host.cpp.o
[ 65%] Building ASC object CMakeFiles/ops_blas.dir/blas/spr/arch35/sspr_host.cpp.o
[ 65%] Building ASC object CMakeFiles/ops_blas.dir/blas/spr2/arch35/sspr2_host.cpp.o
[ 66%] Building ASC object CMakeFiles/ops_blas.dir/blas/spr2/arch35/sspr2_kernel.cpp.o
[ 67%] Building ASC object CMakeFiles/ops_blas.dir/blas/spr/arch35/sspr_kernel.cpp.o
[ 67%] Building ASC object CMakeFiles/ops_blas.dir/blas/stbsv/arch35/stbsv_kernel.cpp.o
[ 69%] Building ASC object CMakeFiles/ops_blas.dir/blas/stbsv/arch35/stbsv_kernel_simt.cpp.o
[ 69%] Building ASC object CMakeFiles/ops_blas.dir/blas/stbsv/arch35/stbsv_host.cpp.o
[ 70%] Building ASC object CMakeFiles/ops_blas.dir/blas/swap/arch35/sswap_host.cpp.o
[ 71%] Building ASC object CMakeFiles/ops_blas.dir/blas/symm/arch35/ssymm_host.cpp.o
[ 72%] Building ASC object CMakeFiles/ops_blas.dir/blas/symm/arch35/ssymm_kernel.cpp.o
[ 72%] Building ASC object CMakeFiles/ops_blas.dir/blas/symv/arch35/ssymv_host.cpp.o
[ 72%] Building ASC object CMakeFiles/ops_blas.dir/blas/swap/arch35/sswap_kernel.cpp.o
[ 73%] Building ASC object CMakeFiles/ops_blas.dir/blas/symv/arch35/ssymv_kernel.cpp.o
[ 74%] Building ASC object CMakeFiles/ops_blas.dir/blas/syr/arch35/syr_host.cpp.o
[ 75%] Building ASC object CMakeFiles/ops_blas.dir/blas/syr2/arch35/syr2_host.cpp.o
[ 75%] Building ASC object CMakeFiles/ops_blas.dir/blas/syr/arch35/syr_kernel.cpp.o
[ 76%] Building ASC object CMakeFiles/ops_blas.dir/blas/syr2/arch35/syr2_kernel.cpp.o
[ 76%] Building ASC object CMakeFiles/ops_blas.dir/blas/tbmv/arch35/stbmv_fallback_kernel.cpp.o
[ 77%] Building ASC object CMakeFiles/ops_blas.dir/blas/tbmv/arch35/stbmv_host.cpp.o
[ 77%] Building ASC object CMakeFiles/ops_blas.dir/blas/tpmv/arch35/stpmv_host.cpp.o
[ 78%] Building ASC object CMakeFiles/ops_blas.dir/blas/tpmv/arch35/stpmv_kernel_simt.cpp.o
[ 79%] Building ASC object CMakeFiles/ops_blas.dir/blas/tbmv/arch35/stbmv_simd_fastpath_kernel.cpp.o
[ 80%] Building ASC object CMakeFiles/ops_blas.dir/blas/tpsv/arch35/stpsv_host.cpp.o
[ 80%] Building ASC object CMakeFiles/ops_blas.dir/blas/tpsv/arch35/stpsv_kernel.cpp.o
[ 81%] Building ASC object CMakeFiles/ops_blas.dir/blas/tpsv/arch35/stpsv_kernel_simt.cpp.o
[ 82%] Building ASC object CMakeFiles/ops_blas.dir/blas/tpttr/arch35/stpttr_host.cpp.o
[ 82%] Building ASC object CMakeFiles/ops_blas.dir/blas/tpttr/arch35/stpttr_kernel.cpp.o
[ 83%] Building ASC object CMakeFiles/ops_blas.dir/blas/trsm/arch35/strsm_kernel.cpp.o
[ 84%] Building ASC object CMakeFiles/ops_blas.dir/blas/trmv/arch35/strmv_kernel.cpp.o
[ 84%] Building ASC object CMakeFiles/ops_blas.dir/blas/trsm/arch35/strsm_blocked_kernel.cpp.o
[ 85%] Building ASC object CMakeFiles/ops_blas.dir/blas/trmv/arch35/strmv_host.cpp.o
[ 86%] Building ASC object CMakeFiles/ops_blas.dir/blas/trsm/arch35/strsm_host.cpp.o
[ 86%] Building ASC object CMakeFiles/ops_blas.dir/blas/trsv/arch35/strsv_host.cpp.o
[ 87%] Building ASC object CMakeFiles/ops_blas.dir/blas/trsv/arch35/strsv_kernel.cpp.o
[ 88%] Building ASC object CMakeFiles/ops_blas.dir/blas/trttp/arch35/strttp_host.cpp.o
[ 88%] Building ASC object CMakeFiles/ops_blas.dir/blas/trttp/arch35/strttp_kernel.cpp.o
WARNING: kernel type of __global__ func: strsm_right_kernel(unsigned char*, unsigned char*, unsigned char*, StrsmTilingData) is not marked. auto type derivate may be failed.
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/copy_atom_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/cube/copy_gm2l1.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/copy.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/routing.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/copy_impl/dn2nz.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/copy_impl/instruction.h:25:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:63:70: error: use of undeclared identifier 'asc_get_phy_buf_addr'
     return MakeLocationMemPtr<PtrPattern>(reinterpret_cast<pointer>(asc_get_phy_buf_addr(0) + arg)); 
                                                                     ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/copy_atom_impl.h:27:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/cube/copy_l0c2gm.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/l0c_to_gm/copy.h:27:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/utils/l0c2out_utils.h:222:51: error: use of undeclared identifier 'asc_get_phy_buf_addr'
    deqTensorTempBuf = reinterpret_cast<uint64_t>(asc_get_phy_buf_addr(0));
                                                  ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:25:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/axpy/axpy_vf.h:66:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:26:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/binary/binary_vf.h:62:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:27:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/binary_scalar/binary_scalar_vf.h:65:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:28:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/cast/cast_vf.h:67:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:29:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/dual/dual_vf.h:66:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:30:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/unary/unary_vf.h:60:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
WARNING: kernel type of __global__ func: MatrixTransformKernel_e5m2e5m2e4m3(unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, MatrixTransformTilingData) is not marked. auto type derivate may be failed.
WARNING: kernel type of __global__ func: MatrixTransformKernel_e5m2e5m2e5m2(unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, MatrixTransformTilingData) is not marked. auto type derivate may be failed.
WARNING: kernel type of __global__ func: MatrixTransformKernel_e4m3e4m3e5m2(unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, unsigned char*, MatrixTransformTilingData) is not marked. auto type derivate may be failed.
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:294:40: error: no matching function for call to 'MakeMemPtr'
        auto tensorAL0Cur = MakeTensor(MakeMemPtr<Location::L0A, T>(l0OffsetCur), layoutAL0Cur);
                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L0A, DataType = float, Addr = unsigned long, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L0A, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:296:40: error: no matching function for call to 'MakeMemPtr'
        auto tensorBL0Cur = MakeTensor(MakeMemPtr<Location::L0B, T>(l0OffsetCur), layoutBL0Cur);
                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L0B, DataType = float, Addr = unsigned long, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L0B, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:310:44: error: expected body of lambda expression
    auto loadL1FromGM = [&](uint64_t kOff) __aicore__ {
                                           ^
/home/xjc/CANN-9.0.0-Community/cann-9.0.0/x86_64-linux/tikcpp/tikcfw/impl/utils/sys_macros.h:39:20: note: expanded from macro '__aicore__'
#define __aicore__ [aicore]
                   ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:310:44: error: use of undeclared identifier 'aicore'
/home/xjc/CANN-9.0.0-Community/cann-9.0.0/x86_64-linux/tikcpp/tikcfw/impl/utils/sys_macros.h:39:21: note: expanded from macro '__aicore__'
#define __aicore__ [aicore]
                    ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:343:40: error: no matching function for call to 'MakeMemPtr'
        auto tensorBL1Cur = MakeTensor(MakeMemPtr<Location::L1, float>(l1OffsetB),
                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L1, DataType = float, Addr = unsigned long, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L1, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:361:33: error: no matching function for call to 'MakeMemPtr'
    auto tensorL0C = MakeTensor(MakeMemPtr<Location::L0C, float>(0),
                                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L0C, DataType = float, Addr = int, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L0C, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/copy_atom_impl.h:27:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/cube/copy_l0c2gm.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/l0c_to_gm/copy.h:27:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/utils/l0c2out_utils.h:212:9: error: use of undeclared identifier 'asc_set_l0c2gm_channel_para'
        asc_set_l0c2gm_channel_para(channelPara);
        ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/utils/l0c2out_utils.h:170:9: note: in instantiation of function template specialization 'AscendC::Te::SetRegisterInstr::SetChannelPara<unsigned long>' requested here
        SetChannelPara<uint64_t>(srcNZC0Stride);
        ^
15 errors generated.
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/copy_atom_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/cube/copy_gm2l1.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/copy.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/routing.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/copy_impl/dn2nz.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/gm_to_l1/copy_impl/instruction.h:25:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:63:70: error: use of undeclared identifier 'asc_get_phy_buf_addr'
     return MakeLocationMemPtr<PtrPattern>(reinterpret_cast<pointer>(asc_get_phy_buf_addr(0) + arg)); 
                                                                     ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/copy_atom_impl.h:27:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/cube/copy_l0c2gm.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/l0c_to_gm/copy.h:27:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/utils/l0c2out_utils.h:222:51: error: use of undeclared identifier 'asc_get_phy_buf_addr'
    deqTensorTempBuf = reinterpret_cast<uint64_t>(asc_get_phy_buf_addr(0));
                                                  ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:25:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/axpy/axpy_vf.h:66:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:26:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/binary/binary_vf.h:62:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:27:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/binary_scalar/binary_scalar_vf.h:65:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:28:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/cast/cast_vf.h:67:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:29:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/dual/dual_vf.h:66:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/transform_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/transform/transform_vf_set.h:30:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/vector/unary/unary_vf.h:60:37: error: use of undeclared identifier 'asc_get_vf_len'
        uint16_t VECTOR_REG_WIDTH = asc_get_vf_len();
                                    ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:294:40: error: no matching function for call to 'MakeMemPtr'
        auto tensorAL0Cur = MakeTensor(MakeMemPtr<Location::L0A, T>(l0OffsetCur), layoutAL0Cur);
                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L0A, DataType = float, Addr = unsigned long, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L0A, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:296:40: error: no matching function for call to 'MakeMemPtr'
        auto tensorBL0Cur = MakeTensor(MakeMemPtr<Location::L0B, T>(l0OffsetCur), layoutBL0Cur);
                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L0B, DataType = float, Addr = unsigned long, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L0B, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:310:44: error: expected body of lambda expression
    auto loadL1FromGM = [&](uint64_t kOff) __aicore__ {
                                           ^
/home/xjc/CANN-9.0.0-Community/cann-9.0.0/x86_64-linux/tikcpp/tikcfw/impl/utils/sys_macros.h:39:20: note: expanded from macro '__aicore__'
#define __aicore__ [aicore]
                   ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:310:44: error: use of undeclared identifier 'aicore'
/home/xjc/CANN-9.0.0-Community/cann-9.0.0/x86_64-linux/tikcpp/tikcfw/impl/utils/sys_macros.h:39:21: note: expanded from macro '__aicore__'
#define __aicore__ [aicore]
                    ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:343:40: error: no matching function for call to 'MakeMemPtr'
        auto tensorBL1Cur = MakeTensor(MakeMemPtr<Location::L1, float>(l1OffsetB),
                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L1, DataType = float, Addr = unsigned long, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L1, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:361:33: error: no matching function for call to 'MakeMemPtr'
    auto tensorL0C = MakeTensor(MakeMemPtr<Location::L0C, float>(0),
                                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:60:24: note: candidate template ignored: substitution failure [with PtrPattern = AscendC::Te::Location::L0C, DataType = float, Addr = int, $3 = 0]
__aicore__ inline auto MakeMemPtr(Addr arg)
                       ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:68:34: note: candidate template ignored: requirement 'IsMemPtrIterator<float, void>::value' was not satisfied [with PtrPattern = AscendC::Te::Location::L0C, Iterator = float]
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/tensor/pointer_pattern.h:74:34: note: candidate template ignored: invalid explicitly-specified argument for 2nd template parameter
__aicore__ inline constexpr auto MakeMemPtr(Iterator iterator)
                                 ^
In file included from /home/xjc/repos/gitcode/ops-blas_dev/blas/trsm/arch35/strsm_blocked_kernel.cpp:15:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/tensor.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/include/tensor_api/algorithm/copy.h:24:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/algorithm/copy_impl.h:25:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/copy_atom_impl.h:27:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/atom/cube/copy_l0c2gm.h:26:
In file included from /home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/l0c_to_gm/copy.h:27:
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/utils/l0c2out_utils.h:212:9: error: use of undeclared identifier 'asc_set_l0c2gm_channel_para'
        asc_set_l0c2gm_channel_para(channelPara);
        ^
/home/xjc/repos/gitcode/ops-blas_dev/build/third_party/ops-tensor/include/tensor_api/impl/tensor_api/arch/cube/utils/l0c2out_utils.h:170:9: note: in instantiation of function template specialization 'AscendC::Te::SetRegisterInstr::SetChannelPara<unsigned long>' requested here
        SetChannelPara<uint64_t>(srcNZC0Stride);
        ^
15 errors generated.
gmake[2]: *** [CMakeFiles/ops_blas.dir/build.make:694: CMakeFiles/ops_blas.dir/blas/trsm/arch35/strsm_blocked_kernel.cpp.o] Error 1
gmake[1]: *** [CMakeFiles/Makefile2:251: CMakeFiles/ops_blas.dir/all] Error 2
gmake[1]: *** Waiting for unfinished jobs....
[ 88%] Linking CXX shared library libops_blasLt.so
[ 88%] Built target ops_blasLt
gmake: *** [Makefile:136: all] Error 2

Special notes for this issue/备注 (Optional / 选填)

likedislike
Xxujiachen8成员
7月16日 添加了label:bug-report
wangzitao
wangzitao成员
7月16日 评论:

您好,感谢反馈,问题确认和修复中。

likedislike
wangzitaowangzitao成员
7月16日 将 eternityk 设为负责人
wangzitaowangzitao成员
7月16日 将 wangzitao_leo 设为负责人
Xxujiachen8成员
7月16日 修改了issue 的描述
Xxujiachen8成员
7月16日 修改了issue 的描述
wangzitaowangzitao成员
7月16日 修改了issue 的描述
wangzitaowangzitao成员
7月24日 移除了负责人 eternityk
LLambdaMiX成员
8月4日 关联了pull request:fix(trsm,gemm): 为 strsm_blocked_kernel 和 gemm 添加 CANN-9.0.0 版本守卫
CANN-robotCANN-robot成员
8月5日 关闭了 issue
CANN-robotCANN-robot成员
8月5日 添加了label:resolved