大家好，依然是那个在CANN训练营里摸爬滚打的新人。成功运行了“向量加法”后，我兴奋了没多久，就被一个更根本的问题缠住了：为什么核函数非要加上 __global__ 这个奇怪的修饰符？那个看起来人畜无害的 kernel_name，又藏着多少我们不知道的秘密？

在 [2025年昇腾CANN训练营第二季] 的课程中，老师反复强调：“理解核函数，是理解Ascend C的钥匙。” 起初我以为这只是一句口号，直到我因为一个函数名的拼写错误调试了一下午后，我才明白，这两个看似简单的语法，其实是连接Host与Device世界的“规则契约”。今天，我们就来把这把钥匙彻底磨亮。

>> 系统学习离不开好课程，训练营的“0基础入门系列”正是这样的存在：点击加入CANN训练营

第一章：从一次“诡异”的编译错误讲起

事情是这样的，我在模仿样例写第二个算子时，自信满满地写下了如下代码：

// 我的“精心”之作
extern "C" void my_awesome_kernel(uint8_t* x, uint8_t* y) {
    // ... 一些计算
}

在 main.cpp 里，我理所当然地调用它：

my_awesome_kernel<<<1, nullptr>>>(deviceX, deviceY);

结果，编译顺利通过，运行时却直接卡死，设备日志报错：找不到对应的内核函数。

我懵了。代码逻辑检查了无数遍，明明没错啊！最后还是训练营的答疑老师在代码里一眼看出了问题：“你的核函数，少了 __global__ 和 __aicore__ 修饰符。”

这一刻我才明白，在Ascend C的世界里，不是所有函数都能被叫做“核函数”的。它必须通过特定的“身份认证”，而这个认证，就始于 __global__。

第二章：__global__——核函数的“身份证”

你可以把 __global__ 理解为一个特殊的通行证。当编译器看到这个关键字时，它会立刻意识到：“哦，这个函数不是跑在CPU上的，是要被编译成在AI Core上执行的代码。”

它的核心作用有三个：

1. 身份标识：明确告知编译器和运行时系统，这是一个设备侧核函数。
2. 编译指引：编译器会启动一套完全不同的编译流程，针对昇腾AI Core的指令集架构进行编译优化，生成的不是普通的x86或ARM代码。
3. 调用约定：它规定了这类函数的特殊调用方式，即使用我们见过的 kernel_name<<<...>>>(...) 这种语法。

那么 __aicore__ 呢？
你可以把它看作是 __global__ 的专属搭档。它进一步限定了这个函数所使用的硬件资源（如寄存器、本地内存等）必须符合AI Core的架构规范。在绝大多数情况下，它们都是成对出现的。训练营老师打了个比方：__global__ 说“我要上高速”，而 __aicore__ 则指定了“我必须开一辆符合高速公路标准的车”。

所以，一个正确、完整的核函数声明应该长这样：

extern "C" __global__ __aicore__ void vector_add_custom(...) {
    // 核函数体
}

少一个，你的函数就无法被正确识别和启动。

第三章：kernel_name——连接两个世界的“密钥”

解决了编译问题，我们来谈谈名字 kernel_name。这不仅仅是一个函数名，它是Host代码与Device代码之间的契约。

1. 命名本身就是契约

在Host侧（main.cpp）我们通过 vector_add_custom<<<...>>> 来调用。在Device侧（.cpp文件）我们必须提供一个完全同名的函数实现。这个匹配是大小写敏感且强绑定的。我之前的错误就是把 awesome 拼成了 awsome，导致链接器找不到符号。

2. extern "C" 的必要性

你有没有注意到核函数声明前的 extern "C"？这可不是摆设。C++编译器为了支持函数重载，会进行名称修饰，把 vector_add_custom 变成一个像 _Z16vector_add_custom... 这样的奇怪名字。这会导致Host侧按原名称找不到它。

extern "C" 的作用就是禁止C++的名称修饰，确保函数在二进制层面保持我们定义的 vector_add_custom 这个名字。这样，运行时系统才能像对暗号一样，准确地找到并启动它。

3. 名字的“可见性”与部署

当我们的应用程序运行时，核函数的代码并不是以源代码的形式存在的。它已经被提前编译成一个独立的内核二进制文件。当Host程序执行 <<<...>>> 调用时，运行时系统会根据你提供的 kernel_name，在一个已注册的内核列表中去查找对应的二进制代码，然后加载到AI Core上执行。

所以，kernel_name 就是那个在部署好的“内核仓库”里，精准抓取所需功能包的唯一密钥。

第四章：一个“错误”示例的剖析

让我们来看一个集大成的错误示例，看看你是否能一眼找出所有问题：

// 核函数文件 (kernel.cpp)
extern "C" __global__ void my_kernel(uint8_t* data) {
    // ... 计算
}

// Host文件 (main.cpp)
int main() {
    // ...
    My_Kernel<<<1, nullptr>>>(deviceData); // 注意这里的大小写！
    // ...
}

这里至少有三个致命错误：

1. 核函数缺少 __aicore__ 修饰符，行为未定义。
2. Host侧调用与核函数定义名不一致：my_kernel vs My_Kernel。C/C++是大小写敏感语言！
3. （潜在风险）如果核函数在C++文件中定义而未用 extern "C"，会因为名称修饰导致链接失败。

正确的写法，是像对待法律条文一样，确保每一个字符都准确无误。

结语：规矩背后是深邃的思考

以前写C++，函数就是函数。但在Ascend C的异构世界里，我第一次如此深刻地感受到，一个函数的定义和调用，竟能牵扯到如此多的底层机制。__global__ 和 kernel_name 这两条简单的语法规则，背后是异构计算架构为了协调两个完全不同硬件（CPU与AI Core）而设立的通信协议。

理解它们，就像是掌握了和AI Core对话的基本语法。从此，你不再是模糊地复制粘贴代码，而是清楚地知道，你写的每一行指令，将如何跨越架构的鸿沟，在专用的处理器上绽放出算力的光芒。

在训练营的后续课程中，我们马上就要接触到更刺激的多核并行，那时我们会看到 <<<TILE_NUM, nullptr>>> 里的 TILE_NUM 如何化身千军万马。我已经准备好了。

---

理解底层原理，才能更好地驾驭上层代码。系统学习Ascend C，从训练营开始 >> 立即报名2025年CANN训练营第二季