Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
52 commits
Select commit Hold shift + click to select a range
c9cb855
Route SuperNPUBench tileop smoke toward current Linx clang
Jun 21, 2026
5cba0ef
Enable Linx direct-boot smoke for SuperNPUBench TAdd
Jun 21, 2026
5b0b076
Exercise SuperNPUBench MatMul in Linx direct-boot bring-up
Jun 21, 2026
833ad2c
Promote TSUB into Linx direct-boot AI smoke
Jun 21, 2026
88a8c4e
Promote TAND into Linx direct-boot AI smoke
Jun 21, 2026
274da5c
Promote TOR into Linx direct-boot AI smoke
Jun 21, 2026
3d53373
Promote TAdds into Linx direct-boot AI smoke
Jun 21, 2026
ab6b727
Promote TSubs into Linx direct-boot AI smoke
Jun 21, 2026
b40395a
Promote TMul into Linx direct-boot AI smoke
Jun 21, 2026
e2ad0fe
Promote TMuls into Linx direct-boot AI smoke
Jun 21, 2026
84c2116
Promote TMax into Linx direct-boot AI smoke
Jun 21, 2026
e05638f
Promote TMaxs into Linx direct-boot AI smoke
Jun 21, 2026
e216516
Promote TAbs into Linx direct-boot AI smoke
Jun 21, 2026
abd66f4
Promote TCopyIn and TCopyOut into Linx direct-boot AI smoke
Jun 21, 2026
7b08d3a
Promote TCopy into Linx direct-boot AI smoke
Jun 21, 2026
f217094
Promote TReshape into Linx direct-boot AI smoke
Jun 21, 2026
86fd3e1
Promote TTrans into Linx direct-boot AI smoke
Jun 21, 2026
ab66556
Promote TPad into Linx direct-boot AI smoke
Jun 21, 2026
891f3a2
Promote TCI into Linx direct-boot AI smoke
Jun 21, 2026
8d053a6
Promote TExpandScalar into Linx direct-boot AI smoke
Jun 21, 2026
41b9930
Promote TExpandRow and TExpandCol into Linx direct-boot AI smoke
Jun 21, 2026
89a3ea9
Promote TRowSum into Linx direct-boot AI smoke
Jun 21, 2026
1120221
Promote TRowMax into Linx direct-boot AI smoke
Jun 21, 2026
5a7a592
Promote row expand reductions into Linx direct-boot AI smoke
Jun 21, 2026
d60feb6
Promote TCmp into Linx direct-boot AI smoke
Jun 21, 2026
6465e4f
Promote TAdd_mask into Linx direct-boot AI smoke
Jun 21, 2026
7fa355b
Promote TDiv into Linx direct-boot AI smoke
Jun 21, 2026
8ce3069
Promote TDivs into Linx direct-boot AI smoke
Jun 21, 2026
c79def5
Promote TRem into Linx direct-boot AI smoke
Jun 21, 2026
dcc09c3
Promote TCvt into Linx direct-boot AI smoke
Jun 21, 2026
6fae79f
Promote TRecip into Linx direct-boot AI smoke
Jun 21, 2026
1ac6112
Promote TSqrt into Linx direct-boot AI smoke
Jun 21, 2026
7af0cd1
Promote MatMacc into Linx direct-boot AI smoke
Jun 21, 2026
c78a331
Promote matrix test smokes into Linx model lane
Jun 21, 2026
2f7a281
Promote TExp into Linx model lane
Jun 21, 2026
ba4e194
Expose MatMul_e4m3's real Linx contract blocker
Jun 21, 2026
b311630
Add linx blockisa llvm musl toolchain 2026-06-22
Jun 22, 2026
fab0235
Expose the real SuperNPUBench MX direct-boot boundary
Jun 22, 2026
aaaa99a
Promote duplicate tileop smokes into model bring-up
Jun 22, 2026
4443e67
Promote duplicate matrix tileop smokes
Jun 22, 2026
f74e129
Promote remaining duplicate tileop smokes
Jun 22, 2026
f9c41dc
Drop stale lowercase tileop manifest row
Jun 22, 2026
d72799a
Expose GELU vector-runtime blocker under Linx
Jun 22, 2026
d6b8f46
Expose real data-object runtime blockers
Jun 22, 2026
f8718f0
Expose SIMD control runtime contract
Jun 22, 2026
0d1fcb6
Promote bounded Linx control lookup smoke
Jun 22, 2026
f569d74
Promote bounded hash-probe control smoke
Jun 23, 2026
a54fc74
Promote MatMul e4m3 through Linx smoke
Jun 23, 2026
43dbc02
Make Linx benchmarks navigable as active NPU suites
Jun 24, 2026
e680825
Merge pull request #2 from PTO-ISA/codex/benchmark-npu-navigation
zhoubot Jun 24, 2026
f104e71
Rename tile memory APIs to load and store
Jun 24, 2026
8e2e1b5
Clarify benchmark navigation and portable samples
Jun 24, 2026
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
  •  
  •  
  •  
1 change: 1 addition & 0 deletions .gitattributes
Original file line number Diff line number Diff line change
@@ -0,0 +1 @@
*.tar.gz filter=lfs diff=lfs merge=lfs -text
221 changes: 137 additions & 84 deletions README.md

Large diffs are not rendered by default.

12 changes: 12 additions & 0 deletions archive/outdated/README.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,12 @@
# Outdated Archive

This directory preserves superseded or unusable material that should not be the default benchmark navigation surface. Nothing here was deleted; the table below records why each item moved and where active work should happen instead.

| Archived path | Rationale | Replacement |
| --- | --- | --- |
| [`tests/other/tileop_api`](tests/other/tileop_api) | Legacy duplicate TileOP API surface with committed generated logs under `script/checknum_true`. | [`../../benchmarks/api/tileop`](../../benchmarks/api/tileop) |
| [`tests/other/py_api`](tests/other/py_api) | Older Python API duplicate. Active Python correctness material is kept outside the benchmark tree. | [`../../tests/py_api`](../../tests/py_api) |
| [`tests/accelerator/v220`](tests/accelerator/v220) | Superseded legacy NPU validation surface, not part of the active Linx benchmark catalog. | [`../../benchmarks/npu`](../../benchmarks/npu) |
| [`tests/accelerator/v310`](tests/accelerator/v310) | Superseded legacy NPU validation surface, not part of the active Linx benchmark catalog. | [`../../benchmarks/npu`](../../benchmarks/npu) |

Archive files may retain historical path references because they document the old layout. Do not add new active benchmark cases here.
Original file line number Diff line number Diff line change
Expand Up @@ -7,12 +7,12 @@
· 文件路径:JanusCoreBench/test/golden_cmp/py_api/src/

· 操作说明:

1. 如果是添加一个新的运算方式(如 texp),则需要新建一个 HPP 文件。
2. 如果是同一运算方式的不同属性(如不同的矩阵尺寸或 tile 大小),则直接在对应的 HPP 文件中添加。

· 标准函数格式:

· 文件头需要包含必要的头文件。
· 声明变量和函数名称时,需注意命名规范。
· 函数声明后,需将函数与模块绑定(m.def)。
Expand Down Expand Up @@ -47,9 +47,9 @@ void texp1_py(py::array_t<float> dst_py, py::array_t<float> src_py) {

tile_shape d0, d1;

TCOPYIN(d0, s0);
TLOAD(d0, s0);
TEXP(d1, d0);
TCOPYOUT(res, d1);
TSTORE(res, d1);
}
}
}
Expand All @@ -72,14 +72,14 @@ void bind_texp(py::module_& m) {
步骤说明:

1. 添加文件头:

· 在文件开头添加包含新 HPP 文件的头文件路径。
```
#include "src/texp1.hpp"
```

2. 添加绑定内容:

· 在模块中绑定新函数。
```
py::module_ _api = m.def_submodule("_api", "API module");
Expand All @@ -93,7 +93,7 @@ void bind_texp(py::module_& m) {
步骤说明:

1. 在 cases 中添加新测试用例:

· 按照以下格式添加新函数的属性。
```
{
Expand All @@ -109,7 +109,7 @@ void bind_texp(py::module_& m) {
· **output_shapes**:输出矩阵的形状。

2. 在 op_map 中添加操作映射:

· 按照以下格式添加新操作的映射。
```
"texp": [
Expand All @@ -129,10 +129,10 @@ void bind_texp(py::module_& m) {

在 /JanusCoreBench/test 路径下,执行以下命令:
```
make clean
make TESTCASE=tileop_py PLAT=cpu PY_LIB=on
python3 golden_cmp/golden_cmp.py -i tadd1
make clean
make TESTCASE=tileop_py PLAT=cpu PY_LIB=on
python3 golden_cmp/golden_cmp.py -i tadd1
```
其中,PLAT 和 PY_LIB 的值可以根据需要进行修改。具体的可选项可参考 common 文件夹下的 Makefile 。其中 -i 后面跟着的是函数的名称,具体的函数名可以参考 config.json 文件中的内容。
其中,PLAT 和 PY_LIB 的值可以根据需要进行修改。具体的可选项可参考 common 文件夹下的 Makefile 。其中 -i 后面跟着的是函数的名称,具体的函数名可以参考 config.json 文件中的内容。
之后print出的对比结果中,在最后两行会显示loss(误差)以及是否pass or fail

Original file line number Diff line number Diff line change
Expand Up @@ -18,18 +18,18 @@ void tadd_py(float* dst, float* src0, float* src1){
int offset = i * (tile_row * gm_col) + j * tile_col;
gm_shape s0(src0 + offset);
gm_shape s1(src1 + offset);
gm_shape res(dst + offset);
gm_shape res(dst + offset);

tile_shape d0, d1, d2;
TCOPYIN(d0, s0);
TCOPYIN(d1, s1);
TLOAD(d0, s0);
TLOAD(d1, s1);
TADD(d2, d0, d1);
TCOPYOUT(res, d2);
TSTORE(res, d2);
}
}
}

#ifdef __cpu_sim__
#ifdef __cpu_sim__
void bind_tadd(py::module_& m) {
m.def("tadd", [](py::array_t<float> dst_py, py::array_t<float> src0_py, py::array_t<float> src1_py){
float* dst = static_cast<float*>(dst_py.request().ptr);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -23,9 +23,9 @@ void texp_py(float* dst, float* src) {

tile_shape d0, d1;

TCOPYIN(d0, s0);
TLOAD(d0, s0);
TEXP(d1, d0);
TCOPYOUT(res, d1);
TSTORE(res, d1);
}
}
}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -23,10 +23,10 @@ void tmax_py(float* dst, float* src0, float* src1){

tile_shape d0, d1, d2;

TCOPYIN(d0, s0);
TCOPYIN(d1, s1);
TLOAD(d0, s0);
TLOAD(d1, s1);
TMAX(d2, d1, d0);
TCOPYOUT(res, d2);
TSTORE(res, d2);
}
}
}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -18,18 +18,18 @@ void tsub_py(float* dst, float* src0, float* src1) {
int offset = i * (tile_row * gm_col) + j * tile_col;
gm_shape s0(src0 + offset);
gm_shape s1(src1 + offset);
gm_shape res(dst + offset);
gm_shape res(dst + offset);

tile_shape d0, d1, d2;
TCOPYIN(d0, s0);
TCOPYIN(d1, s1);
TLOAD(d0, s0);
TLOAD(d1, s1);
TSUB(d2, d0, d1);
TCOPYOUT(res, d2);
TSTORE(res, d2);
}
}
}

#ifdef __cpu_sim__
#ifdef __cpu_sim__
void bind_tsub(py::module_& m) {
m.def("tsub", [](py::array_t<float> dst_py, py::array_t<float> src0_py, py::array_t<float> src1_py){
float* dst = static_cast<float*>(dst_py.request().ptr);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -8,13 +8,12 @@ make TESTCASE=TAdd_mask
make TESTCASE=TAdd
make TESTCASE=TAdds
make TESTCASE=TCopy
make TESTCASE=TCopyIn
make TESTCASE=TCopyOut
make TESTCASE=TLoad
make TESTCASE=TStore
make TESTCASE=TCvt
make TESTCASE=TDiv
make TESTCASE=TDivs
make TESTCASE=test_MatMacc
make TESTCASE=test_matmul
make TESTCASE=test_MatMul
make TESTCASE=TExp
make TESTCASE=TExpandCol
Expand All @@ -33,4 +32,4 @@ make TESTCASE=TRowSumExpand
make TESTCASE=TSqrt
make TESTCASE=TSub
make TESTCASE=TSubs
make TESTCASE=TTrans
make TESTCASE=TTrans
Original file line number Diff line number Diff line change
@@ -1,16 +1,34 @@
#ifndef DATA_H
#define DATA_H

#ifdef __linx
#include <stddef.h>
#include <stdint.h>
extern "C" void exit(int);
extern "C" void free(void *);
extern "C" void *malloc(size_t);
extern "C" int printf(const char *, ...);
#else
#include <iostream>
#include <cmath>
#endif
#include "common/type.hpp"

#ifdef __linx
static constexpr float s_fp32 = 0.1f;
static constexpr __half s_fp16 = __half(0.0f);
static constexpr int8_t s_i8 = 1;
static constexpr int16_t s_i16 = 1;
static constexpr int32_t s_i32 = 1;
static constexpr int64_t s_i64 = 1;
#else
float s_fp32 = 0.1;
__half s_fp16 = 0.1;
int8_t s_i8 = 1;
int16_t s_i16 = 1;
int32_t s_i32 = 1;
int64_t s_i64 = 1;
#endif

template <typename T> void init_src_uint(T *aar, uint16_t size) {
for (uint16_t i = 0; i < size; i++) {
Expand All @@ -36,7 +54,13 @@ void init_src_int8(int8_t *aar, uint16_t size) {

template <typename T> void init_src_fp(T *aar, uint16_t size) {
for (uint16_t i = 0; i < size; i++) {
#ifdef __linx
const float x = (i + 1) / 100.0f;
const float x2 = x * x;
aar[i] = x * (1.0f - x2 / 6.0f + (x2 * x2) / 120.0f);
#else
aar[i] = sin((i + 1) / 100.0f);
#endif
}
}

Expand Down Expand Up @@ -81,22 +105,37 @@ template <typename T> void init_rows_fp(T *aar, uint16_t row, uint16_t col) {
}

template <typename T> void OutArray(const T *aar, size_t size) {
#ifdef __linx
(void)aar;
(void)size;
#else
for (uint16_t i = 0; i < size; i++) {
std::cout << aar[i] << " ";
}
std::cout << std::endl;
#endif
}
void OutArray(const int8_t *aar, size_t size) {
#ifdef __linx
(void)aar;
(void)size;
#else
for (uint16_t i = 0; i < size; i++) {
std::cout << static_cast<int32_t>(aar[i]) << " ";
}
std::cout << std::endl;
#endif
}
void OutArray(const __half *aar, size_t size) {
#ifdef __linx
(void)aar;
(void)size;
#else
for (uint16_t i = 0; i < size; i++) {
std::cout << static_cast<__fp16>(aar[i]) << " ";
}
std::cout << std::endl;
#endif
}

// check memory allocation
Expand Down Expand Up @@ -181,4 +220,4 @@ template <typename T> void check_mem_alloc(const T *p) {
free(d2); \
free(d3);

#endif
#endif
Loading