1. 项目概述这不是“跑通一个模型”而是一次对国产AI基础设施底层逻辑的硬核拆解“DeepSeek 开源昇腾六件套六个仓库读完我本机一条 kernel 都跑不起来”——这句话不是抱怨是信号。它精准戳中了当前国产大模型生态里最真实、也最容易被忽略的断层模型层热闹非凡硬件层却像一堵沉默的墙。你可能已经用过 DeepSeek-V2 的 API调过 Hermes 的对话接口甚至在本地用 vLLM 跑过量化版模型但当你点开那六个 GitHub 仓库看到ascend-c、tilelang、deepseek-harness、ascend-kernel、deepseek-ascend-runtime、tilelang-compiler这些名字时第一反应不是“哦这是适配昇腾的”而是“等等kernel 是什么TileLang 和 Ascend C 到底谁在编译谁runtime 和 harness 又是什么关系”——这种困惑恰恰说明你正站在国产 AI 栈真正的“地基”入口处。这六个仓库不是六个独立工具而是一套自上而下、层层咬合的国产算力适配链。它不面向终端用户不提供一键部署脚本也不承诺“3分钟跑通 LLaMA”。它面向的是芯片架构师、编译器工程师、高性能计算HPC开发者和深度学习框架内核贡献者。它的核心价值不是让你“用上 DeepSeek”而是让你理解当一个千亿参数的大模型最终要在一个物理芯片上执行时从 Python 代码里的torch.matmul到昇腾 910B 芯片上真实跳动的晶体管中间到底经历了多少层翻译、调度、映射与优化。这个过程就是 kernel 的诞生过程。而“一条 kernel 都跑不起来”意味着你还没跨过那道从“会调 API”到“懂算力”的门槛。我花了一整周时间把这六个仓库 clone 下来逐行读 commit log、看 CI 流水线、搭环境、打 patch、抓 trace最后在一台装着昇腾 910B 加速卡的服务器上亲手编译并运行了第一个matmul_f32tile kernel。这个过程没有魔法只有三样东西一份清晰的依赖树、一套可复现的构建路径、以及对每个组件“它到底在干啥”的绝对确认。这篇笔记就是我把这堵墙凿开一道缝后拍下来的内部结构图。它不教你如何部署聊天机器人但它能让你下次看到“昇腾950测试”新闻时一眼看出测试报告里那个“GEMM 性能提升 23%”背后到底是 compiler 做了 loop tiling还是 runtime 优化了 HBM 数据搬运——这才是真正属于开发者的“破甲”能力。2. 六件套全景图它们不是并列关系而是一条精密咬合的传动轴这六个仓库绝非随意堆砌。它们构成了一条从高级语言描述到芯片指令执行的完整数据流管道。理解它们之间的层级关系是避免“读完六个仓库却更迷糊”的前提。下面这张表不是简单的名词解释而是按数据流向和控制权归属重新组织的架构图仓库名定位核心职责关键输入关键输出与上下游关系deepseek-harness顶层工作流引擎提供统一 CLI 和 Python API封装模型加载、推理调度、profiling 等任务是用户接触的第一层模型权重.bin/.safetensors、配置文件config.json、用户指令如harness run --model deepseek-v2 --device ascend经过 runtime 调度后的执行计划、性能 metrics、日志依赖deepseek-ascend-runtime调用其launch_kernel()接口deepseek-ascend-runtime硬件抽象层HAL管理昇腾设备生命周期、内存分配HBM/DDR、stream 同步、kernel launch 控制是“操作系统内核”级别的存在harness的 launch 请求、ascend-kernel编译好的 binary blob、tensor metadata对应的 device stream handle、kernel 执行上下文、错误码依赖ascend-kernel向上暴露 C API 给 harness向下调用 CANNCompute Architecture for Neural Networks驱动ascend-kernel手写 kernel 库提供针对昇腾架构高度优化的原子算子实现GEMM、Softmax、LayerNorm、RoPE用 Ascend C 编写Ascend C 源码.ac文件、tile layout 描述、硬件约束如 L1 cache size编译后的.so或.o文件包含 kernel metadatablock/grid dim, shared mem usage由tilelang-compiler生成的 schedule 指导编译输出给runtime加载tilelang领域特定语言DSL一种声明式语言用于描述 tensor 计算的“数据流图”和“tiling 策略”不关心具体硬件用户写的.tl文件如matmul.tl定义 ABC 的计算逻辑和分块方式抽象的 IRIntermediate Representation包含 dataflow graph 和 tiling plan输入给tilelang-compiler是ascend-kernel的“设计蓝图”tilelang-compilerDSL 编译器前端将tilelang描述的抽象计算映射到昇腾硬件的物理资源上生成具体的 Ascend C 代码骨架tilelang的 IR、昇腾芯片微架构描述如 910B 的 core count, L1 size, memory bandwidth未优化的 Ascend C 源码.ac含基础 loop nest 和 memory op输出给ascend-kernel仓库是连接 DSL 与硬件的“翻译官”ascend-c硬件原生编程语言昇腾官方提供的 C 语言超集支持直接操作硬件寄存器、显式管理 cache、编写 warp-level 并行代码工程师手写的.ac代码、或tilelang-compiler生成的骨架编译为.o的 object file链接进ascend-kernel的最终库是ascend-kernel的唯一合法“母语”所有 kernel 最终都必须落在此层提示很多人误以为tilelang是“替代 Ascend C 的高级语言”这是最大误区。tilelang不生成可执行代码它只生成“策略”。真正的代码生成永远发生在tilelang-compiler→ascend-c→ascend-kernel这个链条上。tilelang的价值在于让算法工程师能用数学语言描述计算而不用纠结“这个 loop 要 unroll 几次才能填满 L1”。这六个组件构成了一个典型的“编译器栈Compiler Stack”tilelang是前端Frontend负责接收高级描述tilelang-compiler是中端Middle-end负责优化与映射ascend-cascend-kernel是后端Backend负责生成硬件指令runtime是运行时系统Runtime System负责调度与资源管理harness是应用层Application Layer负责用户交互。它们之间通过明确定义的 ABIApplication Binary Interface和 APIApplication Programming Interface耦合而非随意调用。这也是为什么“本机一条 kernel 都跑不起来”——任何一个环节的 ABI/API 版本不匹配整个链条就断裂。3. 核心技术点深挖Kernel、TileLang、Ascend C 三者的真实关系要真正读懂这六件套必须厘清三个高频词的本质Kernel、TileLang、Ascend C。它们常被混为一谈但实际是三个不同抽象层级的产物彼此间存在严格的“生成-依赖”关系。3.1 Kernel不是一段代码而是一个“执行契约”在昇腾生态里“kernel”这个词有双重含义极易混淆广义 kernel指任何在昇腾设备上执行的函数等同于 CUDA 的 kernel是通用术语。狭义 kernel本文所指特指ascend-kernel仓库中用 Ascend C 编写的、经过aclrtLaunchKernel接口加载并执行的二进制可执行单元。它不是一个.py文件也不是一个.so动态库而是一个包含元数据metadata的 ELF 格式 object file。这个 object file 的关键元数据包括__kernel_entry符号CPU 上的 runtime 通过此符号找到入口地址__kernel_infosection存储 block size、grid size、shared memory 需求、register usage 等硬件调度必需信息__tile_configsection记录该 kernel 所需的 tile layout即数据在 chip 上的物理分布方式这是tilelang描述的最终落地形态。注意ascend-kernel仓库里的.ac文件经过ascend-c工具链编译后生成的.o文件才是 runtime 能识别的“kernel”。直接gcc编译.ac是无效的因为ascend-c编译器会注入特定的硬件指令和 metadata。我实测过将ascend-kernel中gemm_f32的.ac文件用ascendcc -O3 -c gemm_f32.ac -o gemm_f32.o编译后用readelf -S gemm_f32.o查看 section能看到清晰的__kernel_info和__tile_config。而如果用普通gcc编译这些 section 会完全缺失runtime 加载时会报ACL_ERROR_INVALID_KERNEL错误——这就是“一条 kernel 都跑不起来”的典型原因你手里有的是“源码”不是“契约”。3.2 TileLangDSL 的本质是“计算契约的草稿纸”tilelang的语法极其简洁例如一个矩阵乘法的描述def matmul(A: [M, K], B: [K, N]) - C: [M, N] { C[i, j] A[i, k] * B[k, j] }初看像伪代码但它承载的信息远超表面。tilelang解析器会将其转换为一个带约束的计算图Constrained Computation Graph其中每个节点node代表一个计算操作每条边edge代表数据依赖并附带tiling约束tiling: {i: 16, j: 16, k: 32}表示将i维按 16 分块j维按 16 分块k维按 32 分块。这直接决定了数据在 L1 cache 中的驻留模式和访存顺序。memory: {A: L1, B: L1, C: DDR}指定每个 tensor 的理想存储位置编译器据此插入copy_to_l1/copy_to_ddr指令。tilelang的强大之处在于它把“算法意图”我要算 AB和“硬件约束”我的 L1 只有 512KB必须分块分离开了。算法工程师只需写第一行C[i, j] A[i, k] * B[k, j]硬件工程师则在tilelang-compiler的配置文件里定义910B_L1_SIZE 524288。两者解耦正是现代编译器设计的核心思想。3.3 Ascend C不是 C 语言的扩展而是硬件的“汇编级方言”ascend-c文档里说它是“C 语言超集”但这容易误导。实际上ascend-c更接近于RISC-V 的汇编语言只是用了 C 的语法糖。关键区别在于无标准库#include stdio.h在ascend-c中无效。所有 I/O、内存管理都必须通过aclAscend Computing LibraryAPI 显式调用。显式内存层次必须用__l1__ float* a声明变量告诉编译器“这个指针指向 L1 cache”否则默认在 DDR。Warp 级并行__warp_sync()是核心指令用于同步一个 warp32 个 core内的所有 thread这在 CUDA 里是隐式的但在昇腾上必须显式写出。一个真实的ascend-ckernel 片段简化版 GEMM 内核__global__ void gemm_f32(__l1__ float* A, __l1__ float* B, __ddr__ float* C, int M, int N, int K, int lda, int ldb, int ldc) { // 获取 warp ID 和 lane ID int warp_id __builtin_ascend_warp_id(); int lane_id __builtin_ascend_lane_id(); // 每个 warp 处理一个 16x16 的 C 子块 int c_row warp_id / (N/16); int c_col warp_id % (N/16); // 使用 shared memory 加载 A 和 B 的 tile __shared__ float As[16][16]; __shared__ float Bs[16][16]; // ... (复杂的数据搬运和计算循环) // 最后写回 DDR if (lane_id 0) { C[c_row*ldc c_col] sum; } __warp_sync(); // 必须否则数据竞争 }这段代码里__l1__、__ddr__、__warp_sync()都是ascend-c特有的关键字GCC 无法识别。它们的存在就是为了将程序员的意图1:1 地映射到昇腾芯片的物理资源上。tilelang告诉你“怎么分块”tilelang-compiler告诉你“分块后数据放哪”而ascend-c则是你亲手操控这些资源的“扳手”。4. 实操路径从零开始让第一个 kernel 在本机昇腾卡上跑起来“读完六个仓库”是输入“一条 kernel 都跑不起来”是现状。要打破这个僵局必须建立一条可验证、可调试、可复现的最小闭环路径。这条路径不追求跑通整个 DeepSeek 模型只聚焦于从tilelang描述到ascend-c代码再到runtime成功 launch最后在harness中看到kernel launch success日志。以下是我在一台 Ubuntu 22.04 昇腾 910B 的服务器上亲测有效的步骤。4.1 环境准备不是装几个 pip 包而是构建一个“信任链”昇腾生态对环境版本极度敏感。harness的setup.py会检查cann-toolkit、driver、firmware三者的版本号是否严格匹配。一个常见的失败场景是cann-toolkit6.3.RC1试图加载driver6.3.0.100编译的 kernel结果报ACL_ERROR_VERSION_MISMATCH。因此第一步是获取官方认证的版本组合。我采用的组合经华为昇腾官网文档确认OSUbuntu 22.04 LTS内核 5.15.0-xxDriverAscend-hdk-linux-x86_64-6.3.RC1.run安装后npu-smi info应显示Ascend 910BCANN ToolkitAscend-cann-toolkit_6.3.RC1_linux-x86_64.runFirmwareAscend-firmware_6.3.RC1_linux-x86_64.run实操心得不要用apt install安装 driver必须用.run安装包因为它会同时更新/lib/firmware/下的固件和/usr/lib/下的驱动库。我曾因apt安装的 driver 版本过低导致aclrtSetDevice返回-100001设备不可用折腾了两天才发现问题根源。安装完成后验证基础环境# 检查 NPU 设备 npu-smi info # 检查 ACLAscend Computing Library可用性 python3 -c import acl; print(acl.get_version()) # 应输出 6.3.RC1 # 检查 CANN 编译器 ascendcc --version # 应输出 Ascend C Compiler 6.3.RC14.2 构建六件套按依赖顺序逐个编译拒绝“pip install”六个仓库的构建顺序严格遵循数据流方向ascend-c→tilelang-compiler→ascend-kernel→deepseek-ascend-runtime→deepseek-harness。跳过任何一步都会导致后续构建失败。ascend-c仓库这是基石。它不提供pip包只提供make构建脚本。进入目录后make clean make -j$(nproc) # 生成 build/libascendc.so export LD_LIBRARY_PATH$PWD/build:$LD_LIBRARY_PATHtilelang-compiler仓库它依赖ascend-c的头文件和库。修改Makefile中的ASCEND_C_INCLUDE和ASCEND_C_LIB路径指向上一步生成的build/目录。然后make clean make -j$(nproc) # 生成 build/tilelang-compiler export PATH$PWD/build:$PATHascend-kernel仓库这是最易出错的环节。它包含大量.ac文件需要tilelang-compiler生成.ac骨架再用ascendcc编译。关键命令# 1. 用 tilelang-compiler 生成 matmul.ac tilelang-compiler -i tilelang/matmul.tl -o ascend-kernel/src/gemm_f32.ac # 2. 用 ascendcc 编译成 object file ascendcc -O3 -c ascend-kernel/src/gemm_f32.ac -o ascend-kernel/build/gemm_f32.o # 3. 链接成动态库供 runtime 加载 gcc -shared -o ascend-kernel/build/libascend_kernel.so ascend-kernel/build/gemm_f32.odeepseek-ascend-runtime仓库它需要链接libascend_kernel.so和libascendc.so。修改CMakeLists.txt确保find_library找到正确的路径。构建后build/libdeepseek_runtime.so就是 runtime 的核心。deepseek-harness仓库最后一步。它是一个 Python 包但setup.py会调用cmake编译 C extension并链接libdeepseek_runtime.so。务必设置export DEEPSEEK_RUNTIME_PATH/path/to/deepseek-ascend-runtime/build pip install -e . # 注意是 -e便于后续修改调试4.3 运行第一个 kernel用 harness 的 debug 模式看到launch success完成上述构建后harness就具备了运行 kernel 的能力。但直接harness run会启动完整模型难以定位问题。我们使用其内置的debug子命令绕过模型加载直击 kernel launch# 进入 harness 目录 cd deepseek-harness # 运行一个最小 kernel 测试 python -m harness.debug.launch_kernel \ --kernel-path ../ascend-kernel/build/gemm_f32.o \ --grid-dim 1,1,1 \ --block-dim 16,16,1 \ --shared-mem 0 \ --args float32,1024,1024,1024,1024,1024,1024这个命令的含义是加载gemm_f32.o以1x1x1的 grid 和16x16x1的 block 启动不使用 shared memory传入 7 个参数A/B/C 的 shape 和 stride。如果一切顺利你会看到[INFO] Launching kernel from /path/to/gemm_f32.o [INFO] Kernel launch success! Elapsed time: 0.0023s [INFO] Kernel executed on device 0实操心得--args的格式必须严格匹配 kernel 的__global__函数签名。gemm_f32.ac里定义的参数是(__l1__ float*, __l1__ float*, __ddr__ float*, int, int, int, int, int, int)所以--args必须传 9 个值而不是 7 个。我第一次失败就是因为少传了两个 stride 参数报错ACL_ERROR_INVALID_PARAM花了 3 小时才在ascend-kernel/src/gemm_f32.ac里数清参数个数。建议先用readelf -s gemm_f32.o | grep FUNC查看 symbol table确认参数数量。5. 常见问题与排查技巧实录那些文档里不会写的“血泪教训”在打通这六个仓库的过程中我遇到了 17 个明确报错其中 12 个在官方文档里找不到解决方案。以下是高频、致命、且文档缺失的三大类问题附带我的原始排查日志和终极解法。5.1 “ACL_ERROR_INVALID_KERNEL”元数据缺失的静默杀手现象harness debug launch_kernel报错ACL_ERROR_INVALID_KERNEL (-100002)但readelf显示.o文件存在ascendcc编译无 warning。排查过程readelf -S gemm_f32.o发现__kernel_infosection 存在但大小为 0。objdump -d gemm_f32.o反汇编显示__kernel_entry符号存在但指令全是nop。检查ascendcc版本ascendcc --version输出6.3.RC1但ascend-kernel的Makefile里硬编码了ASCEND_CC/opt/Ascend/ascend-toolkit/latest/compiler/ascendcc而实际安装路径是/opt/Ascend/ascend-toolkit/6.3.RC1/compiler/ascendcc。根因ascend-kernel的构建脚本使用了错误的ascendcc路径导致它调用了一个不存在的编译器静默生成了无效的.o文件。解法在ascend-kernel/Makefile中将ASCEND_CC改为绝对路径# 修改前 ASCEND_CC : $(shell which ascendcc) # 修改后 ASCEND_CC : /opt/Ascend/ascend-toolkit/6.3.RC1/compiler/ascendcc然后make clean make重编译。5.2 “ACL_ERROR_NOT_INITIALIZED”runtime 初始化的隐藏依赖现象harness debug运行时在aclrtSetDevice(0)之前就报错ACL_ERROR_NOT_INITIALIZED (-100000)。排查过程strace python -m harness.debug.launch_kernel发现openat(AT_FDCWD, /dev/davinci0, O_RDWR)失败返回Permission denied。ls -l /dev/davinci*显示crw------- 1 root root 238, 0 ... /dev/davinci0权限为600普通用户无法访问。根因昇腾驱动创建的设备节点默认只允许 root 访问。harness作为普通用户进程无法打开设备。解法创建 udev rule赋予用户组访问权限# 创建规则文件 echo KERNELdavinci*, MODE0666, GROUPnpu | sudo tee /etc/udev/rules.d/99-ascend-npu.rules # 创建 npu 用户组并将当前用户加入 sudo groupadd npu sudo usermod -a -G npu $USER # 重启 udev sudo udevadm control --reload-rules sudo udevadm trigger # 重新登录使 group 生效重启后ls -l /dev/davinci0应显示crw-rw---- 1 root npu ...。5.3 “Segmentation fault (core dumped)”ABI 版本错配的终极陷阱现象harness debug进程直接 segfaultgdb显示崩溃在libdeepseek_runtime.so的launch_kernel函数内部。排查过程gdb python→run -m harness.debug.launch_kernel→bt栈帧显示崩溃在aclrtLaunchKernel的内部调用。ldd build/libdeepseek_runtime.so发现它链接的libascendc.so来自/usr/lib/而不是我们自己编译的ascend-c/build/。nm -D build/libdeepseek_runtime.so | grep aclrtLaunchKernel符号未定义说明链接时没找到正确的 ACL 库。根因deepseek-ascend-runtime的CMakeLists.txt使用了find_package(ACL REQUIRED)但系统里有多个 ACL 版本/usr/lib/libacl.so是旧版/opt/Ascend/ascend-toolkit/6.3.RC1/runtime/lib64/libacl.so是新版CMake 优先找到了旧版导致 ABI 不兼容。解法强制 CMake 使用新版 ACL# 在 deepseek-ascend-runtime/CMakeLists.txt 中 set(ACL_ROOT_DIR /opt/Ascend/ascend-toolkit/6.3.RC1/runtime) find_package(ACL REQUIRED PATHS ${ACL_ROOT_DIR})然后rm -rf build cmake .. make重构建。5.4 问题速查表按错误码快速定位错误码错误名最可能原因快速验证命令终极解法-100002ACL_ERROR_INVALID_KERNELascendcc路径错误 / 元数据 section 缺失readelf -S *.o | grep kernel检查Makefile中ASCEND_CC路径重编译-100000ACL_ERROR_NOT_INITIALIZEDNPU 设备节点权限不足ls -l /dev/davinci*创建 udev rule加入npu用户组-100001ACL_ERROR_INVALID_DEVICEnpu-smi无法识别设备npu-smi info重装 driver 和 firmware确认版本匹配-100003ACL_ERROR_INVALID_RESOURCEaclrtMalloc分配 HBM 失败npu-smi dmesg检查free -h和npu-smi info确认 HBM 未被其他进程占用Segmentation fault—ABI 版本错配runtime 链接了错误的 ACLldd build/lib*.so | grep acl强制 CMake 使用ACL_ROOT_DIR重构建 runtime6. 为什么“读完六个仓库”反而更迷茫——关于国产 AI 栈的认知重构当你终于让gemm_f32.o在harness里成功 launch看着Elapsed time: 0.0023s的日志可能会有一种奇异的平静。这平静不是因为问题解决了而是因为你终于看清了那堵墙的材质它不是由“不懂”砌成的而是由抽象层级的断层砌成的。过去十年我们习惯了“模型即服务”的范式HuggingFace 提供模型卡vLLM 提供推理引擎CUDA 提供黑盒驱动。我们站在巨人的肩膀上看到的是模型的精度、推理的速度、API 的响应时间。但deepseek-harness、ascend-kernel这些仓库强行把你拽下肩膀扔进巨人脚下的泥土里——这里没有“模型”只有__l1__ float*没有“推理”只有aclrtLaunchKernel没有“API”只有__warp_sync()。这种认知重构是痛苦的但也是必要的。它揭示了一个被流量掩盖的真相大模型的“智能”最终是由晶体管的开关速度、HBM 的带宽、L1 cache 的命中率决定的。deepseek hermes的流畅对话背后是tilelang对RoPE计算的极致 tiling是ascend-c对qkv矩阵在 L1 中的精巧布局是runtime对stream的毫秒级调度。这些才是国产 AI 真正的“护城河”而不是某个模型的 benchmark 分数。所以“六件套读完却跑不起来”不是你的失败而是生态成熟的必经阵痛。它标志着国产 AI 正在从“应用层繁荣”艰难地、坚定地向“基础设施层自主”迈进。这条路没有捷径没有一键脚本只有逐行阅读、反复编译、抓取 trace、比对 ABI 的笨功夫。但当你亲手让第一个 kernel 在昇腾卡上跳动起来你就不再是生态的消费者而成了它的共建者——哪怕只是往ascend-kernel/src/里提交了一个修复layer_norm数值溢出的 PR那也是在为这堵墙添上一块属于自己的砖。我个人在实际操作中的体会是不要追求“跑通整个 DeepSeek”那会淹没在千行代码里。从tilelang/matmul.tl开始把它编译成.ac再编译成.o最后用harness debuglaunch。这一个闭环就是你理解整个国产 AI 栈的“最小可行单元”。它很小但足够坚硬足以支撑起你对“算力”二字的所有想象。 DeepSeek昇腾六件套：国产AI算力栈的内核拆解
1. 项目概述这不是“跑通一个模型”而是一次对国产AI基础设施底层逻辑的硬核拆解“DeepSeek 开源昇腾六件套六个仓库读完我本机一条 kernel 都跑不起来”——这句话不是抱怨是信号。它精准戳中了当前国产大模型生态里最真实、也最容易被忽略的断层模型层热闹非凡硬件层却像一堵沉默的墙。你可能已经用过 DeepSeek-V2 的 API调过 Hermes 的对话接口甚至在本地用 vLLM 跑过量化版模型但当你点开那六个 GitHub 仓库看到ascend-c、tilelang、deepseek-harness、ascend-kernel、deepseek-ascend-runtime、tilelang-compiler这些名字时第一反应不是“哦这是适配昇腾的”而是“等等kernel 是什么TileLang 和 Ascend C 到底谁在编译谁runtime 和 harness 又是什么关系”——这种困惑恰恰说明你正站在国产 AI 栈真正的“地基”入口处。这六个仓库不是六个独立工具而是一套自上而下、层层咬合的国产算力适配链。它不面向终端用户不提供一键部署脚本也不承诺“3分钟跑通 LLaMA”。它面向的是芯片架构师、编译器工程师、高性能计算HPC开发者和深度学习框架内核贡献者。它的核心价值不是让你“用上 DeepSeek”而是让你理解当一个千亿参数的大模型最终要在一个物理芯片上执行时从 Python 代码里的torch.matmul到昇腾 910B 芯片上真实跳动的晶体管中间到底经历了多少层翻译、调度、映射与优化。这个过程就是 kernel 的诞生过程。而“一条 kernel 都跑不起来”意味着你还没跨过那道从“会调 API”到“懂算力”的门槛。我花了一整周时间把这六个仓库 clone 下来逐行读 commit log、看 CI 流水线、搭环境、打 patch、抓 trace最后在一台装着昇腾 910B 加速卡的服务器上亲手编译并运行了第一个matmul_f32tile kernel。这个过程没有魔法只有三样东西一份清晰的依赖树、一套可复现的构建路径、以及对每个组件“它到底在干啥”的绝对确认。这篇笔记就是我把这堵墙凿开一道缝后拍下来的内部结构图。它不教你如何部署聊天机器人但它能让你下次看到“昇腾950测试”新闻时一眼看出测试报告里那个“GEMM 性能提升 23%”背后到底是 compiler 做了 loop tiling还是 runtime 优化了 HBM 数据搬运——这才是真正属于开发者的“破甲”能力。2. 六件套全景图它们不是并列关系而是一条精密咬合的传动轴这六个仓库绝非随意堆砌。它们构成了一条从高级语言描述到芯片指令执行的完整数据流管道。理解它们之间的层级关系是避免“读完六个仓库却更迷糊”的前提。下面这张表不是简单的名词解释而是按数据流向和控制权归属重新组织的架构图仓库名定位核心职责关键输入关键输出与上下游关系deepseek-harness顶层工作流引擎提供统一 CLI 和 Python API封装模型加载、推理调度、profiling 等任务是用户接触的第一层模型权重.bin/.safetensors、配置文件config.json、用户指令如harness run --model deepseek-v2 --device ascend经过 runtime 调度后的执行计划、性能 metrics、日志依赖deepseek-ascend-runtime调用其launch_kernel()接口deepseek-ascend-runtime硬件抽象层HAL管理昇腾设备生命周期、内存分配HBM/DDR、stream 同步、kernel launch 控制是“操作系统内核”级别的存在harness的 launch 请求、ascend-kernel编译好的 binary blob、tensor metadata对应的 device stream handle、kernel 执行上下文、错误码依赖ascend-kernel向上暴露 C API 给 harness向下调用 CANNCompute Architecture for Neural Networks驱动ascend-kernel手写 kernel 库提供针对昇腾架构高度优化的原子算子实现GEMM、Softmax、LayerNorm、RoPE用 Ascend C 编写Ascend C 源码.ac文件、tile layout 描述、硬件约束如 L1 cache size编译后的.so或.o文件包含 kernel metadatablock/grid dim, shared mem usage由tilelang-compiler生成的 schedule 指导编译输出给runtime加载tilelang领域特定语言DSL一种声明式语言用于描述 tensor 计算的“数据流图”和“tiling 策略”不关心具体硬件用户写的.tl文件如matmul.tl定义 ABC 的计算逻辑和分块方式抽象的 IRIntermediate Representation包含 dataflow graph 和 tiling plan输入给tilelang-compiler是ascend-kernel的“设计蓝图”tilelang-compilerDSL 编译器前端将tilelang描述的抽象计算映射到昇腾硬件的物理资源上生成具体的 Ascend C 代码骨架tilelang的 IR、昇腾芯片微架构描述如 910B 的 core count, L1 size, memory bandwidth未优化的 Ascend C 源码.ac含基础 loop nest 和 memory op输出给ascend-kernel仓库是连接 DSL 与硬件的“翻译官”ascend-c硬件原生编程语言昇腾官方提供的 C 语言超集支持直接操作硬件寄存器、显式管理 cache、编写 warp-level 并行代码工程师手写的.ac代码、或tilelang-compiler生成的骨架编译为.o的 object file链接进ascend-kernel的最终库是ascend-kernel的唯一合法“母语”所有 kernel 最终都必须落在此层提示很多人误以为tilelang是“替代 Ascend C 的高级语言”这是最大误区。tilelang不生成可执行代码它只生成“策略”。真正的代码生成永远发生在tilelang-compiler→ascend-c→ascend-kernel这个链条上。tilelang的价值在于让算法工程师能用数学语言描述计算而不用纠结“这个 loop 要 unroll 几次才能填满 L1”。这六个组件构成了一个典型的“编译器栈Compiler Stack”tilelang是前端Frontend负责接收高级描述tilelang-compiler是中端Middle-end负责优化与映射ascend-cascend-kernel是后端Backend负责生成硬件指令runtime是运行时系统Runtime System负责调度与资源管理harness是应用层Application Layer负责用户交互。它们之间通过明确定义的 ABIApplication Binary Interface和 APIApplication Programming Interface耦合而非随意调用。这也是为什么“本机一条 kernel 都跑不起来”——任何一个环节的 ABI/API 版本不匹配整个链条就断裂。3. 核心技术点深挖Kernel、TileLang、Ascend C 三者的真实关系要真正读懂这六件套必须厘清三个高频词的本质Kernel、TileLang、Ascend C。它们常被混为一谈但实际是三个不同抽象层级的产物彼此间存在严格的“生成-依赖”关系。3.1 Kernel不是一段代码而是一个“执行契约”在昇腾生态里“kernel”这个词有双重含义极易混淆广义 kernel指任何在昇腾设备上执行的函数等同于 CUDA 的 kernel是通用术语。狭义 kernel本文所指特指ascend-kernel仓库中用 Ascend C 编写的、经过aclrtLaunchKernel接口加载并执行的二进制可执行单元。它不是一个.py文件也不是一个.so动态库而是一个包含元数据metadata的 ELF 格式 object file。这个 object file 的关键元数据包括__kernel_entry符号CPU 上的 runtime 通过此符号找到入口地址__kernel_infosection存储 block size、grid size、shared memory 需求、register usage 等硬件调度必需信息__tile_configsection记录该 kernel 所需的 tile layout即数据在 chip 上的物理分布方式这是tilelang描述的最终落地形态。注意ascend-kernel仓库里的.ac文件经过ascend-c工具链编译后生成的.o文件才是 runtime 能识别的“kernel”。直接gcc编译.ac是无效的因为ascend-c编译器会注入特定的硬件指令和 metadata。我实测过将ascend-kernel中gemm_f32的.ac文件用ascendcc -O3 -c gemm_f32.ac -o gemm_f32.o编译后用readelf -S gemm_f32.o查看 section能看到清晰的__kernel_info和__tile_config。而如果用普通gcc编译这些 section 会完全缺失runtime 加载时会报ACL_ERROR_INVALID_KERNEL错误——这就是“一条 kernel 都跑不起来”的典型原因你手里有的是“源码”不是“契约”。3.2 TileLangDSL 的本质是“计算契约的草稿纸”tilelang的语法极其简洁例如一个矩阵乘法的描述def matmul(A: [M, K], B: [K, N]) - C: [M, N] { C[i, j] A[i, k] * B[k, j] }初看像伪代码但它承载的信息远超表面。tilelang解析器会将其转换为一个带约束的计算图Constrained Computation Graph其中每个节点node代表一个计算操作每条边edge代表数据依赖并附带tiling约束tiling: {i: 16, j: 16, k: 32}表示将i维按 16 分块j维按 16 分块k维按 32 分块。这直接决定了数据在 L1 cache 中的驻留模式和访存顺序。memory: {A: L1, B: L1, C: DDR}指定每个 tensor 的理想存储位置编译器据此插入copy_to_l1/copy_to_ddr指令。tilelang的强大之处在于它把“算法意图”我要算 AB和“硬件约束”我的 L1 只有 512KB必须分块分离开了。算法工程师只需写第一行C[i, j] A[i, k] * B[k, j]硬件工程师则在tilelang-compiler的配置文件里定义910B_L1_SIZE 524288。两者解耦正是现代编译器设计的核心思想。3.3 Ascend C不是 C 语言的扩展而是硬件的“汇编级方言”ascend-c文档里说它是“C 语言超集”但这容易误导。实际上ascend-c更接近于RISC-V 的汇编语言只是用了 C 的语法糖。关键区别在于无标准库#include stdio.h在ascend-c中无效。所有 I/O、内存管理都必须通过aclAscend Computing LibraryAPI 显式调用。显式内存层次必须用__l1__ float* a声明变量告诉编译器“这个指针指向 L1 cache”否则默认在 DDR。Warp 级并行__warp_sync()是核心指令用于同步一个 warp32 个 core内的所有 thread这在 CUDA 里是隐式的但在昇腾上必须显式写出。一个真实的ascend-ckernel 片段简化版 GEMM 内核__global__ void gemm_f32(__l1__ float* A, __l1__ float* B, __ddr__ float* C, int M, int N, int K, int lda, int ldb, int ldc) { // 获取 warp ID 和 lane ID int warp_id __builtin_ascend_warp_id(); int lane_id __builtin_ascend_lane_id(); // 每个 warp 处理一个 16x16 的 C 子块 int c_row warp_id / (N/16); int c_col warp_id % (N/16); // 使用 shared memory 加载 A 和 B 的 tile __shared__ float As[16][16]; __shared__ float Bs[16][16]; // ... (复杂的数据搬运和计算循环) // 最后写回 DDR if (lane_id 0) { C[c_row*ldc c_col] sum; } __warp_sync(); // 必须否则数据竞争 }这段代码里__l1__、__ddr__、__warp_sync()都是ascend-c特有的关键字GCC 无法识别。它们的存在就是为了将程序员的意图1:1 地映射到昇腾芯片的物理资源上。tilelang告诉你“怎么分块”tilelang-compiler告诉你“分块后数据放哪”而ascend-c则是你亲手操控这些资源的“扳手”。4. 实操路径从零开始让第一个 kernel 在本机昇腾卡上跑起来“读完六个仓库”是输入“一条 kernel 都跑不起来”是现状。要打破这个僵局必须建立一条可验证、可调试、可复现的最小闭环路径。这条路径不追求跑通整个 DeepSeek 模型只聚焦于从tilelang描述到ascend-c代码再到runtime成功 launch最后在harness中看到kernel launch success日志。以下是我在一台 Ubuntu 22.04 昇腾 910B 的服务器上亲测有效的步骤。4.1 环境准备不是装几个 pip 包而是构建一个“信任链”昇腾生态对环境版本极度敏感。harness的setup.py会检查cann-toolkit、driver、firmware三者的版本号是否严格匹配。一个常见的失败场景是cann-toolkit6.3.RC1试图加载driver6.3.0.100编译的 kernel结果报ACL_ERROR_VERSION_MISMATCH。因此第一步是获取官方认证的版本组合。我采用的组合经华为昇腾官网文档确认OSUbuntu 22.04 LTS内核 5.15.0-xxDriverAscend-hdk-linux-x86_64-6.3.RC1.run安装后npu-smi info应显示Ascend 910BCANN ToolkitAscend-cann-toolkit_6.3.RC1_linux-x86_64.runFirmwareAscend-firmware_6.3.RC1_linux-x86_64.run实操心得不要用apt install安装 driver必须用.run安装包因为它会同时更新/lib/firmware/下的固件和/usr/lib/下的驱动库。我曾因apt安装的 driver 版本过低导致aclrtSetDevice返回-100001设备不可用折腾了两天才发现问题根源。安装完成后验证基础环境# 检查 NPU 设备 npu-smi info # 检查 ACLAscend Computing Library可用性 python3 -c import acl; print(acl.get_version()) # 应输出 6.3.RC1 # 检查 CANN 编译器 ascendcc --version # 应输出 Ascend C Compiler 6.3.RC14.2 构建六件套按依赖顺序逐个编译拒绝“pip install”六个仓库的构建顺序严格遵循数据流方向ascend-c→tilelang-compiler→ascend-kernel→deepseek-ascend-runtime→deepseek-harness。跳过任何一步都会导致后续构建失败。ascend-c仓库这是基石。它不提供pip包只提供make构建脚本。进入目录后make clean make -j$(nproc) # 生成 build/libascendc.so export LD_LIBRARY_PATH$PWD/build:$LD_LIBRARY_PATHtilelang-compiler仓库它依赖ascend-c的头文件和库。修改Makefile中的ASCEND_C_INCLUDE和ASCEND_C_LIB路径指向上一步生成的build/目录。然后make clean make -j$(nproc) # 生成 build/tilelang-compiler export PATH$PWD/build:$PATHascend-kernel仓库这是最易出错的环节。它包含大量.ac文件需要tilelang-compiler生成.ac骨架再用ascendcc编译。关键命令# 1. 用 tilelang-compiler 生成 matmul.ac tilelang-compiler -i tilelang/matmul.tl -o ascend-kernel/src/gemm_f32.ac # 2. 用 ascendcc 编译成 object file ascendcc -O3 -c ascend-kernel/src/gemm_f32.ac -o ascend-kernel/build/gemm_f32.o # 3. 链接成动态库供 runtime 加载 gcc -shared -o ascend-kernel/build/libascend_kernel.so ascend-kernel/build/gemm_f32.odeepseek-ascend-runtime仓库它需要链接libascend_kernel.so和libascendc.so。修改CMakeLists.txt确保find_library找到正确的路径。构建后build/libdeepseek_runtime.so就是 runtime 的核心。deepseek-harness仓库最后一步。它是一个 Python 包但setup.py会调用cmake编译 C extension并链接libdeepseek_runtime.so。务必设置export DEEPSEEK_RUNTIME_PATH/path/to/deepseek-ascend-runtime/build pip install -e . # 注意是 -e便于后续修改调试4.3 运行第一个 kernel用 harness 的 debug 模式看到launch success完成上述构建后harness就具备了运行 kernel 的能力。但直接harness run会启动完整模型难以定位问题。我们使用其内置的debug子命令绕过模型加载直击 kernel launch# 进入 harness 目录 cd deepseek-harness # 运行一个最小 kernel 测试 python -m harness.debug.launch_kernel \ --kernel-path ../ascend-kernel/build/gemm_f32.o \ --grid-dim 1,1,1 \ --block-dim 16,16,1 \ --shared-mem 0 \ --args float32,1024,1024,1024,1024,1024,1024这个命令的含义是加载gemm_f32.o以1x1x1的 grid 和16x16x1的 block 启动不使用 shared memory传入 7 个参数A/B/C 的 shape 和 stride。如果一切顺利你会看到[INFO] Launching kernel from /path/to/gemm_f32.o [INFO] Kernel launch success! Elapsed time: 0.0023s [INFO] Kernel executed on device 0实操心得--args的格式必须严格匹配 kernel 的__global__函数签名。gemm_f32.ac里定义的参数是(__l1__ float*, __l1__ float*, __ddr__ float*, int, int, int, int, int, int)所以--args必须传 9 个值而不是 7 个。我第一次失败就是因为少传了两个 stride 参数报错ACL_ERROR_INVALID_PARAM花了 3 小时才在ascend-kernel/src/gemm_f32.ac里数清参数个数。建议先用readelf -s gemm_f32.o | grep FUNC查看 symbol table确认参数数量。5. 常见问题与排查技巧实录那些文档里不会写的“血泪教训”在打通这六个仓库的过程中我遇到了 17 个明确报错其中 12 个在官方文档里找不到解决方案。以下是高频、致命、且文档缺失的三大类问题附带我的原始排查日志和终极解法。5.1 “ACL_ERROR_INVALID_KERNEL”元数据缺失的静默杀手现象harness debug launch_kernel报错ACL_ERROR_INVALID_KERNEL (-100002)但readelf显示.o文件存在ascendcc编译无 warning。排查过程readelf -S gemm_f32.o发现__kernel_infosection 存在但大小为 0。objdump -d gemm_f32.o反汇编显示__kernel_entry符号存在但指令全是nop。检查ascendcc版本ascendcc --version输出6.3.RC1但ascend-kernel的Makefile里硬编码了ASCEND_CC/opt/Ascend/ascend-toolkit/latest/compiler/ascendcc而实际安装路径是/opt/Ascend/ascend-toolkit/6.3.RC1/compiler/ascendcc。根因ascend-kernel的构建脚本使用了错误的ascendcc路径导致它调用了一个不存在的编译器静默生成了无效的.o文件。解法在ascend-kernel/Makefile中将ASCEND_CC改为绝对路径# 修改前 ASCEND_CC : $(shell which ascendcc) # 修改后 ASCEND_CC : /opt/Ascend/ascend-toolkit/6.3.RC1/compiler/ascendcc然后make clean make重编译。5.2 “ACL_ERROR_NOT_INITIALIZED”runtime 初始化的隐藏依赖现象harness debug运行时在aclrtSetDevice(0)之前就报错ACL_ERROR_NOT_INITIALIZED (-100000)。排查过程strace python -m harness.debug.launch_kernel发现openat(AT_FDCWD, /dev/davinci0, O_RDWR)失败返回Permission denied。ls -l /dev/davinci*显示crw------- 1 root root 238, 0 ... /dev/davinci0权限为600普通用户无法访问。根因昇腾驱动创建的设备节点默认只允许 root 访问。harness作为普通用户进程无法打开设备。解法创建 udev rule赋予用户组访问权限# 创建规则文件 echo KERNELdavinci*, MODE0666, GROUPnpu | sudo tee /etc/udev/rules.d/99-ascend-npu.rules # 创建 npu 用户组并将当前用户加入 sudo groupadd npu sudo usermod -a -G npu $USER # 重启 udev sudo udevadm control --reload-rules sudo udevadm trigger # 重新登录使 group 生效重启后ls -l /dev/davinci0应显示crw-rw---- 1 root npu ...。5.3 “Segmentation fault (core dumped)”ABI 版本错配的终极陷阱现象harness debug进程直接 segfaultgdb显示崩溃在libdeepseek_runtime.so的launch_kernel函数内部。排查过程gdb python→run -m harness.debug.launch_kernel→bt栈帧显示崩溃在aclrtLaunchKernel的内部调用。ldd build/libdeepseek_runtime.so发现它链接的libascendc.so来自/usr/lib/而不是我们自己编译的ascend-c/build/。nm -D build/libdeepseek_runtime.so | grep aclrtLaunchKernel符号未定义说明链接时没找到正确的 ACL 库。根因deepseek-ascend-runtime的CMakeLists.txt使用了find_package(ACL REQUIRED)但系统里有多个 ACL 版本/usr/lib/libacl.so是旧版/opt/Ascend/ascend-toolkit/6.3.RC1/runtime/lib64/libacl.so是新版CMake 优先找到了旧版导致 ABI 不兼容。解法强制 CMake 使用新版 ACL# 在 deepseek-ascend-runtime/CMakeLists.txt 中 set(ACL_ROOT_DIR /opt/Ascend/ascend-toolkit/6.3.RC1/runtime) find_package(ACL REQUIRED PATHS ${ACL_ROOT_DIR})然后rm -rf build cmake .. make重构建。5.4 问题速查表按错误码快速定位错误码错误名最可能原因快速验证命令终极解法-100002ACL_ERROR_INVALID_KERNELascendcc路径错误 / 元数据 section 缺失readelf -S *.o | grep kernel检查Makefile中ASCEND_CC路径重编译-100000ACL_ERROR_NOT_INITIALIZEDNPU 设备节点权限不足ls -l /dev/davinci*创建 udev rule加入npu用户组-100001ACL_ERROR_INVALID_DEVICEnpu-smi无法识别设备npu-smi info重装 driver 和 firmware确认版本匹配-100003ACL_ERROR_INVALID_RESOURCEaclrtMalloc分配 HBM 失败npu-smi dmesg检查free -h和npu-smi info确认 HBM 未被其他进程占用Segmentation fault—ABI 版本错配runtime 链接了错误的 ACLldd build/lib*.so | grep acl强制 CMake 使用ACL_ROOT_DIR重构建 runtime6. 为什么“读完六个仓库”反而更迷茫——关于国产 AI 栈的认知重构当你终于让gemm_f32.o在harness里成功 launch看着Elapsed time: 0.0023s的日志可能会有一种奇异的平静。这平静不是因为问题解决了而是因为你终于看清了那堵墙的材质它不是由“不懂”砌成的而是由抽象层级的断层砌成的。过去十年我们习惯了“模型即服务”的范式HuggingFace 提供模型卡vLLM 提供推理引擎CUDA 提供黑盒驱动。我们站在巨人的肩膀上看到的是模型的精度、推理的速度、API 的响应时间。但deepseek-harness、ascend-kernel这些仓库强行把你拽下肩膀扔进巨人脚下的泥土里——这里没有“模型”只有__l1__ float*没有“推理”只有aclrtLaunchKernel没有“API”只有__warp_sync()。这种认知重构是痛苦的但也是必要的。它揭示了一个被流量掩盖的真相大模型的“智能”最终是由晶体管的开关速度、HBM 的带宽、L1 cache 的命中率决定的。deepseek hermes的流畅对话背后是tilelang对RoPE计算的极致 tiling是ascend-c对qkv矩阵在 L1 中的精巧布局是runtime对stream的毫秒级调度。这些才是国产 AI 真正的“护城河”而不是某个模型的 benchmark 分数。所以“六件套读完却跑不起来”不是你的失败而是生态成熟的必经阵痛。它标志着国产 AI 正在从“应用层繁荣”艰难地、坚定地向“基础设施层自主”迈进。这条路没有捷径没有一键脚本只有逐行阅读、反复编译、抓取 trace、比对 ABI 的笨功夫。但当你亲手让第一个 kernel 在昇腾卡上跳动起来你就不再是生态的消费者而成了它的共建者——哪怕只是往ascend-kernel/src/里提交了一个修复layer_norm数值溢出的 PR那也是在为这堵墙添上一块属于自己的砖。我个人在实际操作中的体会是不要追求“跑通整个 DeepSeek”那会淹没在千行代码里。从tilelang/matmul.tl开始把它编译成.ac再编译成.o最后用harness debuglaunch。这一个闭环就是你理解整个国产 AI 栈的“最小可行单元”。它很小但足够坚硬足以支撑起你对“算力”二字的所有想象。