diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md index be3dd01..c29c08d 100644 --- a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md @@ -158,7 +158,7 @@ EOF #克隆代码仓库 git clone https://gitlink.org.cn/metax-maca/op_optimization.git #切换到fused moe目录下benchmark项目 -cd op_optimization/基于AI\ Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark +cd op_optimization/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark ``` @@ -333,23 +333,6 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2 本节的冒烟提交只用于验证函数接口、索引逻辑和提交流程；性能优化请在正确性通过后再进行。 -#### 提交前准备 - -##### 代码准备 - -建议在工作目录下保留一个候选版本目录，例如： - -```bash -mkdir -p oj/problem_1_fused_moe -``` - -本教程建议把 OJ 提交源码先保存为： - -```text -oj/problem_1_fused_moe/solution001.cu -``` - -真正提交时，只需要把这个文件里的内容复制到 XPU-OJ 提交框。 ##### 账号准备 @@ -400,11 +383,10 @@ Candidate 就是一次可复现的候选方案。 这样后续多次打榜时，不会忘记哪一版代码对应哪一次提交结果。 -### 步骤 7：实现 Fused MoE GEMM CUDA MACA 冒烟代码 +### 步骤 7：以 CUDA MACA 语言为例实现冒烟代码 -本节以当前 XPU-OJ 题目 **1. Fused MoE GEMM** 为例。 +本节以 CUDA MACA 语言为例介绍冒烟代码的编写方式，选手提交阶段可自行选择 Triton、CUDA MACA 或 TileLang 作为实现语言。 -题目要求你提交一份 CUDA 源码，并提供固定的 C 符号。评测程序会调用这个符号，并检查你是否把结果正确写入 `out`。 #### 接口约定 @@ -446,7 +428,7 @@ token(r) = token_ids[r] / topk expert(r) = expert_ids[r / 128] ``` -再说得直白一点： +注意： * `token_ids` 不是直接拿来当 `a` 的行号，要先除以 `topk`； @@ -734,7 +716,7 @@ extern "C" void run_kernel( 3. 进入比赛页面； [![image6](https://origin.picgo.net/2026/06/23/image6047f2ac4bd2a0f08.png)](https://www.picgo.net/image/image6.4ScJM4) -4. 找到题目： ```text1. Fused MoE GEMM``` +4. 找到题目： ```text1. Fused MoE i8 tn``` 5. 点击题目进入详情页； diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/scripts/build_fused_moe_i8_tn_pybind.sh b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/build_fused_moe_i8_tn_pybind.sh similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/scripts/build_fused_moe_i8_tn_pybind.sh rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/build_fused_moe_i8_tn_pybind.sh diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/scripts/run_fused_moe_i8_tn_benchmark.sh b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_benchmark.sh similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/scripts/run_fused_moe_i8_tn_benchmark.sh rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_benchmark.sh diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/scripts/run_fused_moe_i8_tn_pybind_test.sh b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_pybind_test.sh similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/scripts/run_fused_moe_i8_tn_pybind_test.sh rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_pybind_test.sh diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_bf16_tn/CMakeLists.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/CMakeLists.txt similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_bf16_tn/CMakeLists.txt rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/CMakeLists.txt diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/CMakeLists.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/CMakeLists.txt similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/CMakeLists.txt rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/CMakeLists.txt diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/benchmark_results.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/benchmark_results.md similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/benchmark_results.md rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/benchmark_results.md diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/attr.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/attr.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/attr.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/attr.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/buffer_info.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/buffer_info.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/buffer_info.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/buffer_info.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/cast.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/cast.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/cast.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/cast.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/chrono.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/chrono.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/chrono.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/chrono.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/common.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/common.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/common.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/common.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/complex.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/complex.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/complex.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/complex.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/README.txt b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/README.txt similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/README.txt rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/README.txt diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/pybind11_conduit_v1.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_conduit_v1.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/pybind11_conduit_v1.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_conduit_v1.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/pybind11_platform_abi_id.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_platform_abi_id.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/pybind11_platform_abi_id.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_platform_abi_id.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/wrap_include_python_h.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/wrap_include_python_h.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/conduit/wrap_include_python_h.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/wrap_include_python_h.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/critical_section.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/critical_section.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/critical_section.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/critical_section.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/argument_vector.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/argument_vector.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/argument_vector.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/argument_vector.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/class.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/class.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/class.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/class.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/common.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/common.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/common.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/common.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/cpp_conduit.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/cpp_conduit.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/cpp_conduit.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/cpp_conduit.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/descr.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/descr.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/descr.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/descr.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/dynamic_raw_ptr_cast_if_possible.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/dynamic_raw_ptr_cast_if_possible.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/dynamic_raw_ptr_cast_if_possible.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/dynamic_raw_ptr_cast_if_possible.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/exception_translation.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/exception_translation.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/exception_translation.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/exception_translation.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/function_record_pyobject.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_record_pyobject.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/function_record_pyobject.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_record_pyobject.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/function_ref.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_ref.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/function_ref.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_ref.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/holder_caster_foreign_helpers.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/holder_caster_foreign_helpers.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/holder_caster_foreign_helpers.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/holder_caster_foreign_helpers.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/init.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/init.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/init.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/init.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/internals.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/internals.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/internals.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/internals.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/native_enum_data.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/native_enum_data.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/native_enum_data.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/native_enum_data.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/pybind11_namespace_macros.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/pybind11_namespace_macros.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/pybind11_namespace_macros.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/pybind11_namespace_macros.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/struct_smart_holder.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/struct_smart_holder.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/struct_smart_holder.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/struct_smart_holder.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/type_caster_base.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/type_caster_base.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/type_caster_base.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/type_caster_base.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/typeid.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/typeid.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/typeid.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/typeid.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/using_smart_holder.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/using_smart_holder.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/using_smart_holder.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/using_smart_holder.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/value_and_holder.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/value_and_holder.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/detail/value_and_holder.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/value_and_holder.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen/common.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/common.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen/common.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/common.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen/matrix.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/matrix.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen/matrix.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/matrix.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen/tensor.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/tensor.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eigen/tensor.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/tensor.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/embed.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/embed.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/embed.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/embed.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eval.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eval.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/eval.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eval.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/functional.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/functional.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/functional.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/functional.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/gil.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/gil.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/gil_safe_call_once.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_safe_call_once.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/gil_safe_call_once.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_safe_call_once.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/gil_simple.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_simple.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/gil_simple.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_simple.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/iostream.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/iostream.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/iostream.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/iostream.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/native_enum.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/native_enum.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/native_enum.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/native_enum.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/numpy.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/numpy.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/numpy.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/numpy.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/operators.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/operators.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/operators.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/operators.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/options.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/options.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/options.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/options.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/pybind11.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/pybind11.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/pybind11.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/pybind11.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/pytypes.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/pytypes.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/pytypes.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/pytypes.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/stl.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/stl.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/stl.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/stl.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/stl/filesystem.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/stl/filesystem.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/stl/filesystem.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/stl/filesystem.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/stl_bind.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/stl_bind.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/stl_bind.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/stl_bind.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/subinterpreter.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/subinterpreter.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/subinterpreter.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/subinterpreter.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/trampoline_self_life_support.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/trampoline_self_life_support.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/trampoline_self_life_support.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/trampoline_self_life_support.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/type_caster_pyobject_ptr.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/type_caster_pyobject_ptr.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/type_caster_pyobject_ptr.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/type_caster_pyobject_ptr.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/typing.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/typing.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/typing.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/typing.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/warnings.h b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/warnings.h similarity index 100% rename from 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/third_party/pybind11/include/pybind11/warnings.h rename to 基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/warnings.h diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_bf16_tn/Makefile b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_bf16_tn/Makefile deleted file mode 100644 index 1a343b6..0000000 --- a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_bf16_tn/Makefile +++ /dev/null @@ -1,28 +0,0 @@ -MACA_PATH ?= /opt/maca -MXCC := $(MACA_PATH)/mxgpu_llvm/bin/mxcc -ROOT_DIR := $(abspath $(CURDIR)/../..) -BUILD_DIR := $(CURDIR)/build -SRC := $(CURDIR)/src/fused_moe_bf16_tn_example.cpp -BIN := $(BUILD_DIR)/fused_moe_bf16_tn_example - -.PHONY: all build run clean - -all: build - -build: $(BIN) - -$(BIN): $(SRC) - mkdir -p $(BUILD_DIR) - $(MXCC) -std=c++17 -xmaca \ - -I$(ROOT_DIR)/include \ - -I$(MACA_PATH)/include \ - $(SRC) \ - -L$(MACA_PATH)/lib \ - -lmcruntime \ - -o $(BIN) - -run: $(BIN) - $(BIN) - -clean: - rm -rf $(BUILD_DIR) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/Makefile b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/Makefile deleted file mode 100644 index 869704f..0000000 --- a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/Makefile +++ /dev/null @@ -1,27 +0,0 @@ -MACA_PATH ?= /opt/maca -MXCC := $(MACA_PATH)/mxgpu_llvm/bin/mxcc -BUILD_DIR := $(CURDIR)/build -SRC := $(CURDIR)/src/fused_moe_i8_tn_example.cpp -BIN := $(BUILD_DIR)/fused_moe_i8_tn_example - -.PHONY: all build run clean - -all: build - -build: $(BIN) - -$(BIN): $(SRC) - mkdir -p $(BUILD_DIR) - $(MXCC) -std=c++17 -xmaca \ - -I$(CURDIR)/src \ - -I$(MACA_PATH)/include \ - $(SRC) \ - -L$(MACA_PATH)/lib \ - -lmcruntime \ - -o $(BIN) - -run: $(BIN) - $(BIN) - -clean: - rm -rf $(BUILD_DIR) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o deleted file mode 100644 index ac296d8..0000000 Binary files a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o and /dev/null differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so deleted file mode 100644 index e55d11c..0000000 Binary files a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so and /dev/null differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/benchmark_fused_moe_i8_tn.cpython-312.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/benchmark_fused_moe_i8_tn.cpython-312.pyc deleted file mode 100644 index d8d9f6f..0000000 Binary files a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/benchmark_fused_moe_i8_tn.cpython-312.pyc and /dev/null differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-310.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-310.pyc deleted file mode 100644 index 8f49683..0000000 Binary files a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-310.pyc and /dev/null differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-312.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-312.pyc deleted file mode 100644 index 3092674..0000000 Binary files a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/fused_moe_i8_tn_triton.cpython-312.pyc and /dev/null differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-310.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-310.pyc deleted file mode 100644 index 14ca558..0000000 Binary files a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-310.pyc and /dev/null differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-312.pyc b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-312.pyc deleted file mode 100644 index 073dbb6..0000000 Binary files a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/standalone/fused_moe_i8_tn/python/__pycache__/test_fused_moe_i8_tn_pybind.cpython-312.pyc and /dev/null differ diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/starter/示例冒烟代码-CUDA b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/starter/示例冒烟代码-CUDA new file mode 100644 index 0000000..95bbbc1 --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/starter/示例冒烟代码-CUDA @@ -0,0 +1,457 @@ +#include + +#include + +#include + + + +// xcore1000's CUDA-compatible compiler does not expose NVIDIA's __dp4a. + +// This is a correctness-first replacement: each int32 stores four signed + +// int8 values in little-endian byte order. + +__device__ inline int32_t signed_byte(uint32_t x) { + + x &= 0xffu; + + return (int32_t)(x ^ 0x80u) - 128; + +} + + + +__device__ inline int32_t dp4a_compat(int32_t a, int32_t b, int32_t acc) { + + uint32_t ua = (uint32_t)a; + + uint32_t ub = (uint32_t)b; + + acc += signed_byte(ua) * signed_byte(ub); + + acc += signed_byte(ua >> 8) * signed_byte(ub >> 8); + + acc += signed_byte(ua >> 16) * signed_byte(ub >> 16); + + acc += signed_byte(ua >> 24) * signed_byte(ub >> 24); + + return acc; + +} + + + +__global__ void w8a8_moe_gemm_kernel( + + const int8_t* __restrict__ a, + + const int8_t* __restrict__ b_col_major, + + const float* __restrict__ scale_a, + + const float* __restrict__ scale_b, + + const float* __restrict__ moe_weights, + + const int32_t* __restrict__ token_ids, + + const int32_t* __restrict__ expert_ids, + + int K, int N, int topk, + + __nv_bfloat16* __restrict__ out) + +{ + + int n_base = blockIdx.x * 128; + + int m_base = blockIdx.y * 128; + + int expert = expert_ids[blockIdx.y]; + + + + int tid = threadIdx.x; + + int warp_id = tid / 32; + + int lane_id = tid & 31; + + + + int warp_y = warp_id / 2; + + int warp_x = warp_id & 1; + + int my = lane_id / 8; + + int mx = lane_id & 7; + + + + int m_idx[8]; + + int n_idx[8]; + +#pragma unroll + + for (int i = 0; i < 8; ++i) { + + m_idx[i] = warp_y * 32 + my + i * 4; + + } + +#pragma unroll + + for (int j = 0; j < 8; ++j) { + + n_idx[j] = warp_x * 64 + mx + j * 8; + + } + + + + __shared__ int32_t smem_A[2][128 * 17]; + + __shared__ int32_t smem_B[2][128 * 17]; + + + + int32_t accum[8][8] = {0}; + + + +#pragma unroll + + for (int step = 0; step < 2; ++step) { + + int load_idx = step * 256 + tid; + + int row = load_idx / 4; + + int col_int4 = load_idx & 3; + + + + int r = m_base + row; + + int token = token_ids[r] / topk; + + int64_t a_idx = (int64_t)token * K; + + int4 va = ((const int4*)(a + a_idx))[col_int4]; + + + + int sa = row * 17 + col_int4 * 4; + + smem_A[0][sa + 0] = va.x; + + smem_A[0][sa + 1] = va.y; + + smem_A[0][sa + 2] = va.z; + + smem_A[0][sa + 3] = va.w; + + + + int64_t b_idx = (int64_t)expert * N * K + (int64_t)(n_base + row) * K; + + int4 vb = ((const int4*)(b_col_major + b_idx))[col_int4]; + + + + int sb = row * 17 + col_int4 * 4; + + smem_B[0][sb + 0] = vb.x; + + smem_B[0][sb + 1] = vb.y; + + smem_B[0][sb + 2] = vb.z; + + smem_B[0][sb + 3] = vb.w; + + } + + __syncthreads(); + + + + for (int k_outer = 0; k_outer < K; k_outer += 64) { + + int comp_buf = (k_outer / 64) & 1; + + int load_buf = 1 - comp_buf; + + int next_k = k_outer + 64; + + + + if (next_k < K) { + +#pragma unroll + + for (int step = 0; step < 2; ++step) { + + int load_idx = step * 256 + tid; + + int row = load_idx / 4; + + int col_int4 = load_idx & 3; + + + + int r = m_base + row; + + int token = token_ids[r] / topk; + + int64_t a_idx = (int64_t)token * K + next_k; + + int4 va = ((const int4*)(a + a_idx))[col_int4]; + + + + int sa = row * 17 + col_int4 * 4; + + smem_A[load_buf][sa + 0] = va.x; + + smem_A[load_buf][sa + 1] = va.y; + + smem_A[load_buf][sa + 2] = va.z; + + smem_A[load_buf][sa + 3] = va.w; + + + + int64_t b_idx = (int64_t)expert * N * K + + + (int64_t)(n_base + row) * K + next_k; + + int4 vb = ((const int4*)(b_col_major + b_idx))[col_int4]; + + + + int sb = row * 17 + col_int4 * 4; + + smem_B[load_buf][sb + 0] = vb.x; + + smem_B[load_buf][sb + 1] = vb.y; + + smem_B[load_buf][sb + 2] = vb.z; + + smem_B[load_buf][sb + 3] = vb.w; + + } + + } + + + +#pragma unroll + + for (int k_step = 0; k_step < 16; ++k_step) { + + int32_t reg_A[8]; + + int32_t reg_B[8]; + + + +#pragma unroll + + for (int i = 0; i < 8; ++i) { + + reg_A[i] = smem_A[comp_buf][m_idx[i] * 17 + k_step]; + + } + +#pragma unroll + + for (int j = 0; j < 8; ++j) { + + reg_B[j] = smem_B[comp_buf][n_idx[j] * 17 + k_step]; + + } + + + +#pragma unroll + + for (int i = 0; i < 8; ++i) { + +#pragma unroll + + for (int j = 0; j < 8; ++j) { + + accum[i][j] = dp4a_compat(reg_A[i], reg_B[j], accum[i][j]); + + } + + } + + } + + __syncthreads(); + + } + + + + float scale_row[8]; + +#pragma unroll + + for (int i = 0; i < 8; ++i) { + + int r = m_base + m_idx[i]; + + int token = token_ids[r] / topk; + + scale_row[i] = scale_a[token] * moe_weights[r]; + + } + + + + float scale_col[8]; + +#pragma unroll + + for (int j = 0; j < 8; ++j) { + + int n = n_base + n_idx[j]; + + scale_col[j] = scale_b[(int64_t)expert * N + n]; + + } + + + +#pragma unroll + + for (int i = 0; i < 8; ++i) { + + int r = m_base + m_idx[i]; + +#pragma unroll + + for (int j = 0; j < 8; ++j) { + + int n = n_base + n_idx[j]; + + float v = (float)accum[i][j] * scale_row[i] * scale_col[j]; + + out[(int64_t)r * N + n] = __float2bfloat16(v); + + } + + } + +} + + + +static size_t device_allocation_size(const void* p) { + + mcDrvDeviceptr_t base = 0; + + size_t size = 0; + + (void)wcuMemGetAddressRange(&base, &size, (mcDrvDeviceptr_t)(uintptr_t)p); + + return size; + +} + + + +extern "C" void run_kernel( + + const int8_t* a, + + const int8_t* b_col_major, + + const float* scale_a, + + const float* scale_b, + + const float* moe_weights, + + const int32_t* token_ids, + + const int32_t* expert_ids, + + int64_t topk, + + __nv_bfloat16* out) + +{ + + size_t b_size = device_allocation_size(b_col_major); + + size_t out_size = device_allocation_size(out); + + + + int N = 7168; + + int K = 2048; + + if (b_size > 5000000000ULL) { + + N = 4096; + + K = 7168; + + } + + + + int EM = 4096; + + if (out_size > 128ULL * 1024ULL * 1024ULL) { + + EM = 32768; + + } else if (out_size == 0) { + + // Last-resort fallback if allocation-size probing is unavailable. + + int32_t host_tokens[4096]; + + cudaMemcpy(host_tokens, token_ids, sizeof(host_tokens), cudaMemcpyDeviceToHost); + + int max_token_id = 0; + + for (int i = 0; i < 4096; ++i) { + + if (host_tokens[i] > max_token_id) { + + max_token_id = host_tokens[i]; + + } + + } + + if (max_token_id >= 4096) { + + EM = 32768; + + } + + } + + + + dim3 block(256); + + dim3 grid(N / 128, EM / 128); + + w8a8_moe_gemm_kernel<<>>( + + a, b_col_major, scale_a, scale_b, moe_weights, + + token_ids, expert_ids, K, N, (int)topk, out); + +} \ No newline at end of file