昇腾与英伟达编程架构对比

索木目-黄淮

252人浏览 · 2026-07-15 09:51:50

索木目-黄淮 · 2026-07-15 09:51:50 发布

昇腾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运行时环境在启动内核时自动提供的内置变量，编译器自动设置这些变量。

变量	值	示例
`threadIdx.x`	当前线程在Block中的位置	0, 1, 2, ..., 255
`blockIdx.x`	当前Block在Grid中的位置	0, 1, 2, ..., 1023
`blockDim.x`	每个Block的线程数	256（你设置的）
`gridDim.x`	Grid中的Block总数	1024（你设置的）

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

层次	软件概念	硬件概念	通信范围	同步成本
Grid	所有Block的集合	整个GPU	全局同步（极慢）	需要核函数结束
Block	线程组	SM（流多处理器）	共享内存 + 同步	`__syncthreads()`（中等）
Thread	单个线程	CUDA Core	寄存器/局部内存	无（独立执行）
Warp	32个连续线程（隐形）	Warp调度单元	洗牌指令（`__shfl`）	零延迟（同一Warp内）

3.昇腾与英伟达对比

维度	昇腾 (Ascend C)	NVIDIA (CUDA)
并行模式	数据并行 + 流水并行：代码描述的是“搬运一块->计算一块->搬出一块”的流水线。	线程并行：代码描述的是“一个线程负责一个或几个元素”的逻辑。
核心计算	`AscendC::Add(... , 256)`：一条指令，操作256个数据。	`c[i] = a[i] + b[i]`：一条指令，只操作1个数据，由硬件并行执行。
数据管理	显式管理：必须使用`DataCopy`搬入/搬出`LocalTensor`，控制片上存储。	隐式管理：直接读写全局内存指针，由L1/L2缓存辅助。
内存模型	分层手动控制：显存(GM) -> 片上(UB) -> 计算单元。	统一内存模型：虽然也有共享内存，但基础计算直接面向全局内存。
单次处理量	固定块大小（如256），由代码定义。	单个元素，由`threadIdx`定位。