昇腾C++代码是显式管理数据搬运和流水线的，而CUDA代码更像是为每个线程独立编写逻辑，硬件自动完成并行。对比同样做 C = A + B 的向量加法，差异如下：

1. 昇腾 SIMD 风格 (基于 Ascend C)

在昇腾上，你是在为一个AI Core编程，关注的是如何把数据切成小块，高效地喂给SIMD流水线。（Single Instruction, Multiple Data，单指令多数据流）

• 显式流水线：代码被划分为 CopyIn（搬入）、Compute（计算）、CopyOut（搬出）三个阶段，使用 TQue 队列和 TPipe 内存管理来实现流水线并行。
• 核心指令：调用 Add(zLocal, xLocal, yLocal, tileLength) 完成向量相加，tileLength 代表这一批处理的数据量。
• 数据搬运：数据要搬进 LocalTensor 才能计算，计算完再主动搬回显存。
• 特点：这种风格对程序员要求更高，但能榨干硬件在特定任务上的能效。

以下是基于Ascend C的完整核函数（Kernel）写法，核心是管理数据流水线，每个物理核心上只能运行一个thread.

// 文件名: add_custom.cpp
#include "kernel_operator.h"

// 定义核函数
extern "C" __global__ __aicore__ void add_custom(__gm__ float* a, __gm__ float* b, __gm__ float* c)
{
    // 1. 获取当前核的索引和总核数，用于数据切分
    uint32_t blockIdx = AscendC::GetBlockIdx();
    uint32_t blockNum = AscendC::GetBlockNum();

    // 2. 初始化流水线管理对象 (TPipe)
    AscendC::TPipe pipe;
    // 声明内存块，用于存放数据 (UB: Unified Buffer, 片上存储)
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueA, inQueueB; // 输入队列
    AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueC;          // 输出队列
    pipe.InitBuffer(inQueueA, 256 * sizeof(float)); // 分配单次处理256个float
    pipe.InitBuffer(inQueueB, 256 * sizeof(float));
    pipe.InitBuffer(outQueueC, 256 * sizeof(float));

    // 3. 循环Tiling，分块处理整段数据
    uint32_t totalLength = N; // 假设全局N已知
    uint32_t perBlockLength = totalLength / blockNum;
    uint32_t startIdx = blockIdx * perBlockLength;
    uint32_t endIdx = (blockIdx + 1) * perBlockLength;

    for (uint32_t i = startIdx; i < endIdx; i += 256) {
        // 4. 搬运数据: 从全局内存(GM)搬入片上内存(UB)
        AscendC::LocalTensor<float> aLocal = inQueueA.AllocTensor<float>();
        AscendC::LocalTensor<float> bLocal = inQueueB.AllocTensor<float>();
        AscendC::DataCopy(aLocal, a + i, 256); 
        AscendC::DataCopy(bLocal, b + i, 256);

        // 5. 【核心SIMD计算】一条指令同时计算256个数
        AscendC::LocalTensor<float> cLocal = outQueueC.AllocTensor<float>();
        AscendC::Add(cLocal, aLocal, bLocal, 256); // 硬件执行向量加法

        // 6. 搬运结果: 从片上内存搬回全局内存
        AscendC::DataCopy(c + i, cLocal, 256);

        // 7. 释放资源
        inQueueA.FreeTensor(aLocal);
        inQueueB.FreeTensor(bLocal);
        outQueueC.FreeTensor(cLocal);
    }
}

2. NVIDIA  SIMT 风格 (基于 CUDA)

你是在为海量并发线程编程，每个线程只处理一个元素，硬件负责将它们组织成Warp（32线程一组）协同执行。SIMT(Single Instruction, Multiple Threads，单指令多线程）

• 线程即核心：一个线程只负责计算 C[tid] = A[tid] + B[tid]。通过内置变量 threadIdx 和 blockIdx 找到自己是哪个“工人”。
• 计算方式：硬件将大量线程组成 Warp，以32个线程为一组，执行相同的指令处理不同的数据，这就是SIMT的核心。
• 数据搬运：调用 cudaMemcpy 在启动内核前搬运整块数据，核函数直接读写全局内存。
• 特点：编程模型更通用灵活，对开发者更友好，但在部分场景下能效不如专门优化的SIMD。

这是标准的CUDA核函数，核心是利用线程ID处理对应元素。

// 文件名: add_kernel.cu
#include <cuda_runtime.h>

// 核函数定义，__global__表示在设备端执行
__global__ void add_kernel(const float* a, const float* b, float* c, int N)
{
    // 1. 计算当前线程负责处理的全局索引, 全局线程ID计算
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    //           ↑           ↑            ↑
    //        Block索引    Block大小     Thread索引
    // 2. 步长 (用于处理超出线程总数的情况)
    int stride = blockDim.x * gridDim.x; 
    //            ↑                ↑
    //  Block大小(多少个thread)  Grid大小(多少个Block)

    // 3. 循环处理，每个线程处理多个元素 (Grid Stride Loop)
    for (int i = idx; i < N; i += stride) {
        // 4. 【核心SIMT计算】每个线程执行一次标量加法
        //    硬件以Warp(32线程)为单位，执行相同指令处理不同数据
        c[i] = a[i] + b[i]; 
    }
}

// 宿主端调用代码 (Host侧)
int main() {
    int N = 1 << 20; // 1M个元素
    size_t bytes = N * sizeof(float);
    // ... (分配内存cudaMalloc、拷贝数据cudaMemcpy省略)

    // 启动内核: 配置网格和线程块
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    add_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);
    
    // ... (拷回结果cudaMemcpy、释放内存)
    return 0;
}

CUDA编程模型中，threadIdx.x、blockIdx.x、blockDim.x、gridDim.x是CUDA运行时环境在启动内核时自动提供的内置变量，编译器自动设置这些变量。

Grid、Block、Thread、Warp对应的概念，这是CUDA编程的基础。

3.昇腾与英伟达对比

简单来说，Ascend C代码在解决“如何高效搬运数据”，而CUDA代码在解决“如何分配任务给线程”。前者类似编排工厂流水线，后者类似调动工人团队。