Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
18 changes: 18 additions & 0 deletions .github/workflows/ci_sim.yml
Original file line number Diff line number Diff line change
Expand Up @@ -210,6 +210,7 @@ jobs:
rm -rf "${PYPTO_WORKSPACE}"
rm -rf "${PYPTO_RUN_WORKSPACE}"
rm -rf "${PTO_ISA_ROOT}"
rm -rf "${GITHUB_WORKSPACE}/.work/camodel"

- name: Prepare LLVM source
shell: bash
Expand Down Expand Up @@ -453,6 +454,23 @@ jobs:
run_pypto_fa_ptoas_smoke a2a3sim
run_pypto_int8_codegen_smoke a2a3sim

- name: Prepare quiet camodel
shell: bash
run: |
set -euo pipefail
readarray -t camodel_candidates < <(
find "${ASCEND_HOME_PATH}" -type d -path '*/simulator/dav_3510/lib' | sort
)
if [[ "${#camodel_candidates[@]}" -eq 0 ]]; then
echo "ERROR: cannot find dav_3510 camodel lib under ${ASCEND_HOME_PATH}" >&2
exit 1
fi
SIM_LIB_DIR="$(python3 scripts/prepare_quiet_camodel.py \
--source-dir "${camodel_candidates[0]}" \
--output-dir "${GITHUB_WORKSPACE}/.work/camodel")"
echo "SIM_LIB_DIR=${SIM_LIB_DIR}" >> "${GITHUB_ENV}"
echo "SIM_LIB_DIR=${SIM_LIB_DIR}"

- name: Run VPTO SIM validation
if: ${{ true }}
shell: bash
Expand Down
103 changes: 103 additions & 0 deletions scripts/prepare_quiet_camodel.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,103 @@
#!/usr/bin/env python3
# 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.

"""Prepare a shared quiet camodel directory.

The output directory is a simulator lib directory view: non-config entries are
symlinked from the original camodel lib directory, and config.json is copied and
patched to reduce simulator log/dump I/O.
"""

import argparse
import fcntl
import json
import os
import shutil
import sys


def parse_args():
parser = argparse.ArgumentParser(description=__doc__)
parser.add_argument("--source-dir", required=True, help="Original camodel simulator lib directory.")
parser.add_argument("--output-dir", required=True, help="Quiet camodel output directory.")
return parser.parse_args()


def patch_config(config_path):
with open(config_path, "r", encoding="utf-8") as handle:
config = json.load(handle)
config.setdefault("LOG", {})["flush_level"] = 6
config.setdefault("LOG", {})["core_enable_mask"] = ["0x0"]
wrapper = config.setdefault("WRAPPER", {})
wrapper["adapter_log_file_level"] = 6
wrapper["aic_wrap_log_file_level"] = 6
wrapper["cosim_log_file_level"] = 6
wrapper["cosim_log_flush_level"] = 6
wrapper["cosim_log_scr_level"] = 6
with open(config_path, "w", encoding="utf-8") as handle:
json.dump(config, handle, indent=4)
handle.write("\n")


def prepare_quiet_camodel(source_dir, quiet_dir):
source_dir = os.path.realpath(source_dir)
if not os.path.isdir(source_dir):
raise FileNotFoundError(f"camodel source dir is invalid: {source_dir}")

os.makedirs(quiet_dir, exist_ok=True)
lock_path = os.path.join(quiet_dir, ".quiet-camodel.lock")
with open(lock_path, "w", encoding="utf-8") as lock_file:
fcntl.flock(lock_file, fcntl.LOCK_EX)
return prepare_quiet_camodel_locked(source_dir, quiet_dir)
Comment on lines +54 to +57

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

Opening the lock file in write ("w") mode truncates the file before acquiring the lock. If another process is currently holding the lock or waiting for it, this truncation can cause race conditions or interfere with its operations. It is a safer best practice to open lock files in append ("a") mode to avoid truncating the file.

Suggested change
lock_path = os.path.join(quiet_dir, ".quiet-camodel.lock")
with open(lock_path, "w", encoding="utf-8") as lock_file:
fcntl.flock(lock_file, fcntl.LOCK_EX)
return prepare_quiet_camodel_locked(source_dir, quiet_dir)
lock_path = os.path.join(quiet_dir, ".quiet-camodel.lock")
with open(lock_path, "a", encoding="utf-8") as lock_file:
fcntl.flock(lock_file, fcntl.LOCK_EX)
return prepare_quiet_camodel_locked(source_dir, quiet_dir)



def prepare_quiet_camodel_locked(source_dir, quiet_dir):
source_marker = os.path.join(quiet_dir, ".quiet-camodel-source")
config_path = os.path.join(quiet_dir, "config.json")
if os.path.isfile(source_marker):
with open(source_marker, "r", encoding="utf-8") as handle:
existing_source = handle.read().strip()
if existing_source != source_dir:
raise RuntimeError(
f"output dir already points to {existing_source}, cannot reuse it for {source_dir}"
)
if os.path.isfile(config_path):
return os.path.abspath(quiet_dir)

for name in os.listdir(source_dir):
src = os.path.join(source_dir, name)
dst = os.path.join(quiet_dir, name)
if name == "config.json":
shutil.copy2(src, dst)
os.chmod(dst, os.stat(dst).st_mode | 0o200)
continue
if os.path.lexists(dst):
continue
try:
os.symlink(src, dst)
except FileExistsError:
pass

if not os.path.isfile(config_path):
raise FileNotFoundError(f"camodel config.json not found under {source_dir}")
patch_config(config_path)
with open(source_marker, "w", encoding="utf-8") as handle:
handle.write(source_dir + "\n")
return os.path.abspath(quiet_dir)


def main():
args = parse_args()
quiet_dir = os.path.abspath(args.output_dir)
print(prepare_quiet_camodel(args.source_dir, quiet_dir))
return 0


if __name__ == "__main__":
sys.exit(main())
86 changes: 86 additions & 0 deletions test/tilelang_st/npu/a5/src/st/smoke/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -0,0 +1,86 @@
# 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.

cmake_minimum_required(VERSION 3.16)
project(tilelang_st)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_POSITION_INDEPENDENT_CODE ON)

# CMake 3.27+ may ask the linker to emit dependency files via
# `--dependency-file`. bisheng/cce-ld does not support that flag, so disable
# linker-generated link dependencies for this standalone ST build.
if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.27)
set(CMAKE_LINK_DEPENDS_USE_LINKER FALSE)
endif()

set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin)
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)

# --------------------------------------------------------------------------
# PTOAS binary — passed by run_st.py via -DPTOAS_BIN=...
# --------------------------------------------------------------------------
if(NOT DEFINED PTOAS_BIN)
message(FATAL_ERROR "PTOAS_BIN is not set. Pass -DPTOAS_BIN=/path/to/ptoas to cmake.")
endif()

# --------------------------------------------------------------------------
# ASCEND environment
# --------------------------------------------------------------------------
if(NOT DEFINED ENV{ASCEND_HOME_PATH})
message(FATAL_ERROR "Cannot find ASCEND_HOME_PATH, please run set_env.sh.")
else()
set(ASCEND_HOME_PATH $ENV{ASCEND_HOME_PATH})
endif()

set(PTO_ISA_ROOT "${CMAKE_CURRENT_LIST_DIR}/../../../../../../../../pto-isa" CACHE PATH "Path to pto-isa repo")
set(PTO_TILELANG_ST_COMMON_DIR
"${CMAKE_CURRENT_LIST_DIR}/../common")
set(ASCEND_DRIVER_PATH /usr/local/Ascend/driver)

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

The Ascend driver path is hardcoded to /usr/local/Ascend/driver. If the driver is installed in a non-standard directory, the build will fail. It is better to allow overriding this path via a CMake cache variable or environment variable.

if(NOT DEFINED ASCEND_DRIVER_PATH)
    set(ASCEND_DRIVER_PATH /usr/local/Ascend/driver CACHE PATH "Path to Ascend driver directory")
endif()


set(CMAKE_COMPILER bisheng)
set(CMAKE_C_COMPILER ${CMAKE_COMPILER})
set(CMAKE_CXX_COMPILER ${CMAKE_COMPILER})

add_compile_options(
-D_FORTIFY_SOURCE=2
-O2 -std=c++17
-Wno-macro-redefined -Wno-ignored-attributes -Wno-unknown-attributes
-fstack-protector-strong
-fPIC
)
add_link_options(
-s
-Wl,-z,relro
-Wl,-z,now
)

set(CMAKE_CCE_COMPILE_OPTIONS
-xcce
-fPIC
-Xhost-start -Xhost-end
"SHELL:-mllvm -cce-aicore-stack-size=0x8000"
"SHELL:-mllvm -cce-aicore-function-stack-size=0x8000"
"SHELL:-mllvm -cce-aicore-record-overflow=true"
"SHELL:-mllvm -cce-aicore-addr-transform"
"SHELL:-mllvm -cce-aicore-dcci-insert-for-scalar=false"
)

set(CMAKE_CPP_COMPILE_OPTIONS
-xc++
"SHELL:-include stdint.h"
"SHELL:-include stddef.h"
)

include_directories(
${ASCEND_HOME_PATH}/include
${ASCEND_DRIVER_PATH}/kernel/inc
)

add_subdirectory(testcase)
Loading
Loading