CUDA与AscendC对比（一）

对于《NVIDIA CUDA编程实操一：环境检查与入门》的扩展。

这一段cuda例子，global代表什么？我看AscendC里面也有这个？是否AscendC和CUDA很像呢？ 另外helloFromGPU<<<1, 5>>>()这种调用方式也像。这里的1,5代表什么呢？但threadIdx.x是从0开始编号的

__global__ 是一个函数类型限定符，它告诉编译器这个函数是一个核函数（Kernel）。

• 由主机（Host）调用，在设备（Device）上执行：__global__ 函数被编译并发送到GPU（或NPU）上，由成千上万的并行计算单元去执行。

• 返回void类型：__global__ 函数不能有返回值。

Ascend C在设计上大量借鉴了CUDA的编程思想，所以看起来会非常眼熟。

CUDA让你直接管理成千上万个细粒度的“线程”，思考“我的第i个线程要处理哪个数据”；而Ascend C让你管理粗粒度的“任务块”，思考“我的第i个核要处理哪一块数据”。

helloFromGPU<<<1, 5>>>() 会创建1个Block，这个Block里有5个线程。这是CUDA（以及Ascend C）启动核函数的特殊语法，括号里的两个数字就是执行配置（Execution Configuration）。

• 第一个参数 1 (gridDim)：代表整个任务被划分成1个线程块（Block）。注意不是划分为几个grid，但可以说grid有几个block组成。

• 第二个参数 5 (blockDim)：代表这个线程块里包含5个线程（Thread）。

threadIdx.x则是CUDA为每个线程内置提供的一个3维向量结构体（threadIdx）的x分量。从0开始编号的。因此，你看到的输出会是 Hello World from GPU thread 0! 到 Hello World from GPU thread 4!，但顺序是不确定的，因为它们是并行执行的。

这里面涉及到一些GPU规格，请你从对代码开发者有帮助的角度来解释下

1. Compute capability: 6.0 - 功能集版本：GPU架构的"版本号"，决定了支持哪些硬件特性

   • 新版本支持的特性（如动态并行、统一内存、Tensor Core等）在老版本上不可用

   • 6.0对应的Pascal架构，支持统一内存、16位浮点运算，但不支持Tensor Core（从7.0开始）

   • 编译时用 -arch=sm_60 针对这个架构优化

1. Warp size: 32 threads - 调度单元：GPU执行指令的最小单位，32个线程同时执行相同指令

// ❌ 差：会导致线程束发散
if (threadIdx.x % 2 == 0) {
    // 奇数线程在这段分支里空闲
} else {
    // 偶数线程在这段分支里空闲
}
​
// ✅ 好：整个线程束走同一路径
if (threadIdx.x < 16) {
    // 前16个线程执行
} else {
    // 后16个线程执行（仍然是整个线程束）
}

• 建议线程块大小设为32的倍数（128、256最常见）

• 条件分支要尽量按线程束对齐，避免性能损失

1. Max threads per block: 1024 - 线程块容量上限

// ❌ 这个启动配置会失败！
dim3 block(1024, 2);  // 总共2048 > 1024
​
// ✅ 只要乘积不超过1024即可
dim3 block(32, 32);   // 1024 threads
dim3 block(32, 16);   // 512 threads

• 处理大数组时要合理选择2D/3D线程块形状

• 这个限制决定了单个Block能处理的连续数据量

1. Shared memory per block: 48 KB - 片上高速缓存：每个Block独享的高速缓存（类似CPU的L1，但软件可控）

__global__ void matrixMul(float *A, float *B, float *C) {
    __shared__ float tileA[32][32];  // 占用4KB
    __shared__ float tileB[32][32];  // 占用4KB
    // 总共8KB，合理使用
    
    // 利用shared memory可以避免重复从全局内存读取
    // 速度提升
}

• 常见用途：矩阵分块计算、数据重用、线程间通信

• 注意：48KB是上限，声明过多会导致kernel启动失败

• 可以通过cudaFuncSetAttribute()动态配置部分SM的L1缓存为shared memory

1. Max threads per multiprocessor: 2048 - 流处理器容量：每个SM（Streaming Multiprocessor，计算单元）能同时"活跃"的线程总数

// 假设你的配置：
helloFromGPU<<<1000, 256>>>();  // 每个Block 256线程
​
// P100有56个SM，每个SM同时最多2048线程
// 所以每个SM最多同时运行 2048/256 = 8 个Block
// 理论上可同时运行 56 * 8 = 448 个Block
// 但实际硬件会尽力调度1000个Block

• 延迟隐藏：当线程等待内存时，SM会切换到其他线程束

• 线程太少会浪费计算能力

• 线程太多（> 4-5倍SM容量）会提升寄存器压力

1. Multiprocessor count: 56 - 并行计算单元数量，GPU内部有多少个独立的计算核心组（SM）

总结

核心原则：根据你的实际数据量和工作负载，选择能最大化SM占用率（active warps / max warps）的参数。目标是让每个SM至少有 4-8个活跃线程束 来隐藏内存延迟。

AscendC编程似乎没有涉及的这么复杂

Ascend C 的设计初衷就是在保留并行计算强大能力的同时，尽可能地隐藏硬件细节，降低开发者的心智负担。

1. 从“管理线程”到“设计流水线”

2. 自动化的并发与同步

在CUDA中，你需要显式地管理线程间的同步（例如使用__syncthreads()），用起来很繁琐。

• Ascend C的做法：引入了队列 (Queue) 和内存管理模块 (Pipe)。CopyIn、Compute、CopyOut三个“任务”通过队列来传递数据，Pipe统一管理内存。当一个任务往队列里放入数据后，下一个任务就能自动获取并处理。这种流水线并行的同步机制由框架自动完成，你只需按部就班地调用 EnQue（入队）、DeQue（出队）API即可，无需操心底层的锁或同步指令。

1. 围绕“块”而非“线程”的编程

CUDA里有 threadIdx，你需要精确到每个线程。而Ascend C的数据操作单元是 LocalTensor，它代表一块连续的内存数据。

• API设计：Ascend C的API（如DataCopy、Add）大多是对整块数据的操作。它允许你以half*这样的指针，操作一整段数据，而不是像CUDA那样写一个循环让每个线程处理一个元素。

• 核函数启动：CUDA用 <<<gridDim, blockDim>>> 启动一个二维的线程网格。Ascend C则简化为了 <<<blockDim, ...>>>，这个 blockDim 直接指定要在多少个AI Core上运行这个核函数。你可以在核函数里通过 GetBlockIdx() 知道自己是第几个核，然后根据自己的“核编号”去处理对应的那一段数据。

它将你从CUDA那种必须精确控制每一个线程的“微观管理”中解放出来，让你可以用更贴近“数据流动”和“任务流水线”的思维去编程。