我是那个在CANN训练营里，刚刚用 DataCopy 征服了数据搬运的开发者。当我的单核算子性能提升65%后，我一度以为这就是技术的尽头。直到训练营的导师在代码评审会上抛出一个问题：“你的算子现在跑满一个AI Core的百分之几？”

我查了下性能日志，答案令人羞愧：不到5%。原来，我一个核函数实例，只利用了昇腾AI处理器庞大算力资源的冰山一角。导师指着 kernel_name<<<TILE_NUM, nullptr>>>(...) 中的 TILE_NUM 说：“你想让几百个AI Core核心为你同时工作吗？钥匙，就在这里。”

这个 TILE_NUM，就是核函数接口中那个看似不起眼的 uint32_t blockDim。在 [2025年昇腾CANN训练营第二季] 的“码力全开特辑”中，我花了整整一周时间与它“搏斗”，终于从迷惑到通透。今天，就让我们一同解开多核并行的启动奥秘。

>> 突破单核思维，需要体系化的指导：点击加入训练营，驾驭多核编程

第一章：一个人的工地与一支工程队

在深入技术细节前，让我们先建立一个心智模型。

单核模式：一个人的工地
我之前写的所有算子，无论是向量加还是复杂的运算，都只启动了一个核函数实例。这就像把一整栋大楼的建造任务，交给一个全能的工人。他需要自己测量、自己搬砖、自己砌墙……即使他效率再高，对于大型工程来说，也是杯水车薪。

多核模式：一支工程队
而 blockDim 参数，就是你作为项目总指挥，派出的工程队小队数量。

• blockDim = 8：派出8个小队。
• blockDim = 256：派出256个小队。

每个小队都拿着完全相同的施工图纸（也就是你的核函数代码），但他们会根据自己小队的编号（blockIdx），自动去完成属于自己那部分任务。

第二章：揭开面纱——blockDim 的身份与使命

在我们熟悉的核函数调用中：

vector_add_custom<<<TILE_NUM, nullptr, stream>>>(totalLength, TILE_NUM, ...);

这个 TILE_NUM 被传递给了核函数的 uint32_t blockDim 形参。它的官方名字叫 “任务网格中X维度的尺寸” ，通俗讲，就是你希望启动的核函数实例的总数。

1. blockDim 如何驱动多核？

关键在于核函数内部的两个“秘密武器”：GET_BLOCK_IDX() 和 GET_BLOCK_NUM()。

• GET_BLOCK_NUM()：返回的就是 blockDim 的值，告诉每个实例：“我们总共有多少兄弟在并肩作战。”
• GET_BLOCK_IDX()：返回当前实例的索引（从0开始），告诉它：“你是第几号？”

这样，每个核函数实例就知道了自己的唯一身份和团队规模，从而可以计算出自己应该处理哪一部分数据。

2. 一个“反直觉”的深刻理解

这里有一个至关重要的、曾让我困惑许久的点：你只编译了一份核函数二进制代码，但通过指定 blockDim，运行时系统会创建 blockDim 个独立的实例，每个实例都执行这份相同的代码，但 GET_BLOCK_IDX() 的返回值各不相同。

这就好比你把一份相同的施工图纸复印了100份，发给100个工程队。图纸一样，但每个队根据自己编号的不同，去建造不同楼层的公寓。

第三章：实战！将向量加法从单核升级到多核

让我们以向量加法为例，完成这次关键的升级。

单核版本的局限：
单核版本需要处理全部 totalLength 个数据。如果数据量巨大，一个核就会成为瓶颈。

多核改造四步法：

第一步：修改核函数——从“我全干”到“我干我的那份”

__global__ __aicore__ void vector_add_multi_core(
    uint32_t totalLength,   // 总数据量
    uint32_t tileNum,       // 这就是blockDim！任务块总数
    uint8_t* x1,
    uint8_t* x2,
    uint8_t* y)
{
    // 1. 获取身份信息
    int32_t blockIdx = GET_BLOCK_IDX();  // “我是第几号工人？”
    int32_t blockDim = GET_BLOCK_NUM();  // “我们总共有多少工人？”

    // 2. 任务划分：计算每个核应该处理的数据量和起始位置
    uint32_t dataPerBlock = totalLength / blockDim; // 每个核的基础工作量
    uint32_t remainder = totalLength % blockDim;    // 无法整除的余数

    // 3. 负载均衡：处理余数，确保所有数据都被处理
    // 策略：前 `remainder` 个核，每人多干1个
    uint32_t myLength = dataPerBlock + (blockIdx < remainder ? 1 : 0);
    uint32_t myOffset = blockIdx * dataPerBlock + (blockIdx < remainder ? blockIdx : remainder);

    // 4. 基于 myLength 和 myOffset 进行后续的数据搬运和计算
    // ... (这与单核流程完全一样，但现在是处理局部数据)
    __gm__ uint8_t* globalX1 = x1 + myOffset;
    __gm__ uint8_t* globalX2 = x2 + myOffset;
    __gm__ uint8_t* globalY = y + myOffset;

    constexpr int32_t TILE_LENGTH = 256;
    uint8_t localX1[TILE_LENGTH];
    uint8_t localX2[TILE_LENGTH];

    // 注意：这里需要循环处理 myLength 个数据，可能超过TILE_LENGTH，此处为简化示例
    if (myLength <= TILE_LENGTH) {
        DataCopy<LocalTensor, GM_ADDR>(localX1, globalX1, myLength / BLOCK_SIZE, 0, 0);
        DataCopy<LocalTensor, GM_ADDR>(localX2, globalX2, myLength / BLOCK_SIZE, 0, 0);

        // 计算
        for (int i = 0; i < myLength; ++i) {
            localY[i] = localX1[i] + localX2[i];
        }

        DataCopy<GM_ADDR, LocalTensor>(globalY, localY, myLength / BLOCK_SIZE, 0, 0);
    } else {
        // 处理比TILE_LENGTH大的情况，需要分多次搬运和计算
    }
}

第二步：修改Host调用——启动千军万马

int main() {
    // ... 准备数据、申请设备内存等步骤不变

    // 关键修改：选择并设置 blockDim
    uint32_t totalLength = 1024 * 1024; // 1M 个数据
    uint32_t blockDim = 128; // 我们启动128个核实例！

    // 调用核函数！
    vector_add_multi_core<<<blockDim, nullptr>>>(totalLength, blockDim, deviceX1, deviceX2, deviceY);

    // ... 后续同步、拷贝结果等步骤不变
}

第四章：灵魂拷问——blockDim 到底该设为多少？

这是我当时最大的疑问。是不是越大越好？我试过设为1、16、128、1024，结果发现性能并非线性增长，在某个值之后甚至会下降。

训练营的课程给出了科学的指导原则：

1. 与数据规模匹配：blockDim 最好是总数据量的约数，这样负载最均衡。同时，每个核处理的数据量（myLength）不能太小，否则任务划分的开销会占比过高。
2. 与硬件资源匹配：一个昇腾AI处理器有固定的物理核心数量。如果 blockDim 远大于物理核心数，那么多出来的实例就需要排队等待，由硬件进行时分复用，这会带来额外的调度开销。
3. 与资源限制匹配：每个核实例都会消耗LM、寄存器等资源。blockDim 设置过大，可能导致总资源需求超出硬件能力，编译或运行失败。
4. 黄金法则：在实践中，通常需要一个性能调优的过程。 使用 msprof 等工具进行 profiling，尝试不同的 blockDim 值（例如从64开始，以2的幂次递增到512），观察算子的耗时，找到一个性能的“甜蜜点”。

第五章：常见“翻车”现场与调试心得

1. 翻车一：数据覆盖与丢失
   我的第一个多核版本，结果总是少了一部分。原因是任务划分逻辑有误，myOffset 计算错误，导致两个核处理了同一块数据，而另一块数据无人问津。
   调试技巧：在核函数里打印 blockIdx, myLength, myOffset，确保每个核的任务区间是连续且覆盖完整的 [0, totalLength-1]。

2. 翻车二：负载不均
   当 totalLength 不能被 blockDim 整除时，如果简单地对所有核分配 dataPerBlock，那么最后一个核的任务会少一些。虽然没问题，但没能让所有核同时干完活，存在微小的性能浪费。我上面代码中“前 remainder 个核多干1个”的策略，就是一种经典的负载均衡手段。

3. 翻车三：原子操作的陷阱
   当多个核需要更新同一个全局变量时（例如在全局统计一个值），就会发生数据竞争。这时必须使用原子操作（如 atomic_add），否则结果不可预测。这是我后面在实现复杂算子时遇到的又一个深坑。

结语：从“工匠”到“指挥官”的思维跃迁

理解并掌握了 blockDim 的用法，是我在CANN训练营学习中一个里程碑式的时刻。它意味着我的编程思维，从关注“一个核如何把事情做对”的工匠思维，跃迁到了“如何指挥成百上千个核高效协作”的指挥官思维。

我不再只关心单核内部的流水线和寄存器优化，更要站在全局视角，思考任务的分解、负载的均衡、资源的分配。这才是异构并行计算的真正魅力所在。

现在，我的向量加法算子在处理百万级数据时，性能相比单核版本提升了近百倍。但这远不是终点。在训练营接下来的课程中，我们将探索如何在每个核内部，让数据搬运和计算再次“并行”起来——这就是核内流水线并行。多核并行与核内并行相结合，才能将昇腾处理器的算力压榨到极致。征程，才刚刚开始。

---

想要学会如何指挥AI Core的千军万马吗？>> 立即报名2025年CANN训练营第二季，成为并行编程的指挥官