Atomic 性能测试总结

1. Atomic 测试背景

本测试针对 Ascend NPU 昇腾950（CANN 9.2.0 / bisheng 编译器） 上多核并发执行 AscendC::AtomicAdd 写 Global Memory（GM）的性能开销进行量化。

• 测试目标：测量多个 AI Core 并发对 GM 中 uint32_t 做原子加，并同步等待全部核完成（DSB + DCCI 轮询）的每轮平均时钟周期（tick）开销。
• 测试方法：每轮循环包含
  • [初始化] `SyncAll → 写初值 → SyncAll → clock() 起
  • [开始工作]→ AtomicAdd → dsb → dcci 轮询直至汇聚
  • [结束]→ clock() 止`，共 1000 轮取平均，再对 24 组配置各独立运行 3 次。
• 数据单位：ns/核/轮（取每核平均 tick 的均值）。
• 运行环境：正式未插桩逻辑，无中间计时、无逐轮回写；每次运行独立分配 GM，覆盖运行间设备状态/地址映射/调度波动。
• 结果：24 组 × 3 次 = 72 次均值

2. Atomic 的影响因子

影响因子	取值	作用机制
元素步长 / 字节间隔	0/0, 1/4, 16/64, 32/128, 64/256, 128/512	决定各核原子目标地址间距：stride=0 全核抢同一字（最大原子争用）；stride>0 各核写 x+coreId*stride 独立槽，降低直接争用但增大轮询扫描范围与缓存一致性开销
并发核数	4 / 8 / 16 / 32	决定原子操作争用强度与汇聚轮询的核数
重复轮数	1000	单次运行内的迭代数，平滑单轮抖动
重复运行次数	3	评估运行间波动范围（不是单轮极值）
同步与缓存策略	SyncAll + dsb(DSB_ALL) + dcci(ENTIRE_DATA_CACHE, CACHELINE_OUT)	固定开销，纳入每次测量；stride=0 轮询单字 x[0]==kCoreNum，stride>0 轮询遍历累加所有核槽位
构建/运行环境	CANN 9.2.0	保证测的是硬件真实开销，非观测开销

3. 各种不同测试方案

共 6 步长 × 4 核数 = 24 组配置，按"争用—间距"两维正交设计：

• 方案 A（stride=0，同地址atomic）：全核对同一 4 字节字做 AtomicAdd，最大原子串行化；轮询单字最便宜。4/8/16/32 核各一组。
• 方案 B（stride=4B，同缓存行）：相邻字，仍在同一 cache line，存在 false-sharing。
• 方案 C（stride= 64B，一行间隔）：恰好一个 cache line 间距。
• 方案 D（stride= 128B，两行间隔）：脱离单行 false-sharing。
• 方案 E（stride= 256B，四行间隔）：进一步拉开。
• 方案 F（stride=512B，八行间隔）：最大间距，低核数下几乎无共享，但 32 核下轮询扫描地址跨度最大。

每组配置：独立编译（bisheng -c + 链接 demo）→ 运行 → 校验 VALIDATION PASS → 收集每核 sum_ticks/avg_ticks → 汇总 mean_ticks。同一 source_sha256 在 3 次重复中保持一致，确保源码未变，差异仅来自运行环境。

4. 测试数据

4.1 三次运行均值（ns/核/轮，按 步长 × 核数）

步长(el)/字节(B)	4 核	8 核	16 核	32 核
0 / 0	587.87	793.76	1102.88	1916.24
1 / 4	639.22	819.27	1142.67	2096.50
16 / 64	742.09	933.82	1499.11	4193.31
32 / 128	528.62	746.89	1090.02	2243.78
64 / 256	437.92	691.07	851.49	3709.59
128 / 512	598.25	643.63	741.00	4568.40

4.2 运行间波动范围（最高均值 − 最低均值，ns）

步长(el)/字节(B)	4 核	8 核	16 核	32 核
0	221.37	10.75	291.39	66.80
4	83.19	61.93	144.58	47.48
64	93.56	96.36	605.18	30.42
128	9.33	29.09	23.70	69.42
256	232.97	41.07	27.15	33.39
512	129.53	0.66	21.97	140.21

4.3 单次 run_1 示例（stride= 4 核，per_core）

core	sum_ticks	avg_ticks
0	712729	712
1	726351	726
2	588618	588
3	591064	591

mean_ticks = 654.69（4 核均值），可见核间本身就有约 20% 的负载/调度偏差。

5. 测试代码示意

5.1 编译与运行命令（来自 commands.json）

# 编译：bisheng 按 dav-3510 架构编译 .asc 源码为 .o
/usr/local/Ascend/cann-9.2.0/bin/bisheng \
    -std=c++17 --npu-arch=dav-3510 -c --asc-aicore-lang \
    hello_world.asc -o hello_world.o

# 链接：生成可执行 demo
/usr/local/Ascend/cann-9.2.0/bin/bisheng \
    --npu-arch=dav-3510 hello_world.o -o demo

# 运行
./demo

5.2 kernel 测试代码（stride=0 同字争用）

核心循环：仅 core 0 复位共享字，全核对该字做 AtomicAdd，轮询 x[0]==kCoreNum 汇聚。

#include <cstdio>
#include <cstdint>
#include "utils/debug/asc_printf.h"
#include "utils/debug/asc_time.h"
#include "kernel_operator.h"
#include "acl/acl.h"

constexpr uint32_t kCoreNum    = 4;   // 可调：4 / 8 / 16 / 32
constexpr uint32_t kStride     = 0;   // 可调：元素步长，字节间隔 = 4 * kStride
constexpr uint32_t kIterations = 1000;

__global__ __vector__ void hello_world(__gm__ uint32_t* x)
{
    const uint32_t coreId = AscendC::GetBlockIdx();
    uint64_t sumTime = 0;
    for (uint32_t j = 0; j < kIterations; ++j) {
        AscendC::SyncAll();
        if (coreId == 0) {                      // 仅 core 0 复位共享字
            AscendC::WriteGmBypassDCache(x, 0u);
        }
        AscendC::SyncAll();
        uint64_t start = clock();
        AscendC::AtomicAdd(x + coreId * kStride, 1u);  // 全核抢同一字
        dsb(mem_dsb_t::DSB_ALL);
        while (true) {                          // 汇聚轮询：单字
            dcci(static_cast<__gm__ void*>(x),
                 cache_line_t::ENTIRE_DATA_CACHE,
                 dcci_dst_t::CACHELINE_OUT);
            if (x[0] == kCoreNum) { break; }
            asm volatile("nop");
        }
        uint64_t end = clock();
        sumTime += end - start;
    }
    AscendC::SyncAll();
    printf("RESULT core=%u stride=%u cores=%u iterations=%u sum_ticks=%llu avg_ticks=%llu\n",
           coreId, kStride, kCoreNum, kIterations,
           (unsigned long long)sumTime,
           (unsigned long long)(sumTime / kIterations));
}

5.3 kernel 测试代码（stride>0 分散槽位）

核心差异：各核复位各自槽位 x + coreId*kStride，轮询改为累加全部核槽位 n==kCoreNum。下例为 stride=64 / 32 核（最大跨度扫描）。

constexpr uint32_t kCoreNum = 32;
constexpr uint32_t kStride  = 64;   // 字节间隔 = 256B
constexpr uint32_t kIterations = 1000;

__global__ __vector__ void hello_world(__gm__ uint32_t* x)
{
    const uint32_t coreId = AscendC::GetBlockIdx();
    uint64_t sumTime = 0;
    for (uint32_t j = 0; j < kIterations; ++j) {
        AscendC::SyncAll();
        AscendC::WriteGmBypassDCache(x + coreId * kStride, 0u);  // 各核复位各自槽位
        AscendC::SyncAll();
        uint64_t start = clock();
        AscendC::AtomicAdd(x + coreId * kStride, 1u);            // 各核写独立槽
        dsb(mem_dsb_t::DSB_ALL);
        while (true) {                          // 汇聚轮询：累加全部核槽位
            dcci(static_cast<__gm__ void*>(x),
                 cache_line_t::ENTIRE_DATA_CACHE,
                 dcci_dst_t::CACHELINE_OUT);
            uint32_t n = 0;
            for (uint32_t i = 0; i < kCoreNum; ++i) {
                n += x[i * kStride];            // 跨度 = kCoreNum * 字节间隔
            }
            if (n == kCoreNum) { break; }
            asm volatile("nop");
        }
        uint64_t end = clock();
        sumTime += end - start;
    }
    AscendC::SyncAll();
    printf("RESULT core=%u stride=%u cores=%u iterations=%u sum_ticks=%llu avg_ticks=%llu\n",
           coreId, kStride, kCoreNum, kIterations,
           (unsigned long long)sumTime,
           (unsigned long long)(sumTime / kIterations));
}

两版本仅 kCoreNum/kStride 常量与汇聚分支不同（stride=0 单字判等 vs stride>0 累加遍历），其余同步/计时骨架一致，保证不同方案测得开销可比。

5.4 Host 侧：分配、启动、校验

#define ACL_CHECK(expr) do { \
    aclError err = (expr); \
    if (err != ACL_SUCCESS) { \
        std::fprintf(stderr, "ACL_ERROR %s code=%d\n", #expr, (int)err); \
        return 1; \
    } \
} while (0)

int main()
{
    ACL_CHECK(aclInit(nullptr));
    ACL_CHECK(aclrtSetDevice(0));
    aclrtStream stream = nullptr;
    ACL_CHECK(aclrtCreateStream(&stream));

    const size_t count = kStride == 0 ? 1 : kCoreNum * kStride;  // stride=0 仅 1 个字
    const size_t bytes = count * sizeof(uint32_t);
    std::vector<uint32_t> x(count, 0);
    uint32_t* xDevice = nullptr;
    ACL_CHECK(aclrtMalloc((void**)&xDevice, bytes, ACL_MEM_MALLOC_HUGE_FIRST));
    ACL_CHECK(aclrtMemcpy(xDevice, bytes, x.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE));

    hello_world<<<kCoreNum, 0, stream>>>(xDevice);               // 启动 kCoreNum 个核
    ACL_CHECK(aclrtSynchronizeStream(stream));
    ACL_CHECK(aclrtMemcpy(x.data(), bytes, xDevice, bytes, ACL_MEMCPY_DEVICE_TO_HOST));

    // 校验：stride=0 共享字应为 kCoreNum；stride>0 槽位间隔处应为 1
    bool valid = true;
    for (size_t i = 0; i < count; ++i) {
        uint32_t expected = kStride == 0 ? kCoreNum : (i % kStride == 0 ? 1u : 0u);
        if (x[i] != expected) { std::fprintf(stderr, "MISMATCH index=%zu\n", i); valid = false; }
    }
    std::printf("VALIDATION %s stride=%u stride_bytes=%u cores=%u bytes=%zu\n",
                valid ? "PASS" : "FAIL", kStride, kStride * 4, kCoreNum, bytes);

    ACL_CHECK(aclrtFree(xDevice));
    ACL_CHECK(aclrtDestroyStream(stream));
    ACL_CHECK(aclrtResetDevice(0));
    ACL_CHECK(aclFinalize());
    return valid ? 0 : 2;
}

5.5 运行输出示例（output.log，stride=0 / 4 核）

VALIDATION PASS stride=0 stride_bytes=0 cores=4 bytes=4
[AIV Block 0/4] RESULT core=0 stride=0 cores=4 iterations=1000 sum_ticks=712729 avg_ticks=712
[AIV Block 1/4] RESULT core=1 stride=0 cores=4 iterations=1000 sum_ticks=726351 avg_ticks=726
[AIV Block 2/4] RESULT core=2 stride=0 cores=4 iterations=1000 sum_ticks=588618 avg_ticks=588
[AIV Block 3/4] RESULT core=3 stride=0 cores=4 iterations=1000 sum_ticks=591064 avg_ticks=591

6. 总结

1. 核数是主导因子：开销随并发核数单调上升。低核数（4→8→16）线性温和增长；32 核出现跃变——多数配置 16→32 核开销翻 2~6 倍（如 stride 128：741→4568，约 6.2×；stride 16：1499→4193，约 2.8×），表明 32 核并发原子 + 全 cache DCI 汇聚成为瓶颈。

2. 步长效应非单调，与核数耦合：

   • 低核数（4/8/16）下大步长更优：stride 64/128 把各核槽位拉出共享缓存行，false-sharing 消失，开销最低（4 核 stride 64 仅 437.92，16 核 stride 128 仅 741.00）。
   • 32 核下小步长反超：stride 0（同字）以 1916.24 拿下 32 核最低值，因为轮询只读单字最便宜；而 stride 128/64 在 32 核下成为全局最差（4568 / 3709），原因是 dcci 全局失效 + 遍历 32 个大跨度槽位累加的轮询代价远超原子争用本身。
   • stride 16 在所有核数下都偏差（16 核 1499、32 核 4193），疑似与 cache line / bank 对齐产生持续干扰，是最差步长之一。

3. 运行间稳定性：大部分配置波动 <100 ns，可复现性好；最稳定为 stride 128 / 8 核（0.66 ns），最不稳定为 stride 16 / 16 核（605.18 ns）与 stride 0 / 16 核（291.39 ns），提示中等核数 × 中等步长区存在对调度/地址映射敏感的"共振点"，工程上应避开。

4. 工程建议：

   • 多核原子汇聚场景，优先把目标地址分散到 ≥256B（stride≥64）以消除 false-sharing，但仅在 ≤16 核时成立。
   • 32 核满载时，要么用同字原子（stride 0）让硬件串行化、轮询最简；要么改用层级/分批汇聚，避免大跨度 dcci 轮询。
   • 避开 stride 16（64B，单 cache line 间距）这一最易触发抖动的配置。
   • 报告均值应附带运行间波动范围；对 16/32 核关键配置建议增加重复次数（≥5）以收敛方差。