CUDA与AscendC对比（三）

__global__ 是一个函数类型限定符，它告诉编译器这个函数是一个核函数（Kernel） 既然如此，叫kernel不是更直观？

__kernel__ 被AMD和OpenCL用了，NVIDIA选择了不同的名字。

1. 2006年CUDA刚出来时，"kernel"这个词在并行计算领域已经被OpenCL（2009年才发布）之前的Brook、Sh等研究项目广泛使用。

2. 语义准确性：__global__ 强调的是函数的可见性和调用范围——它既可以被Host（CPU端）调用，又运行在Device（GPU端）上，是一种"全局"的函数。而"kernel"在操作系统里已经有含义（内核态），可能会引起混淆。

CUDA内部其实也有__device__和__host__

__global__ void kernel_1() {...}   // 被CPU调用，在GPU上执行
__device__ void kernel_2() {...}   // 被GPU调用，在GPU上执行
__host__ void kernel_3() {...}     // 被CPU调用，在CPU上执行（普通函数）
__host__ __device__ void kernel_4() {...}  // 既可在CPU也可在GPU

这里__global__的特殊性就体现出来了：它是唯一一个被Host启动的关键字。

最后说一句：__global__函数是唯一一个可以被CPU代码调用、但实际在GPU上执行的函数。应该没错吧。

核函数（Kernel）只是在设备上运行的函数，为什么叫核函数呢？这个核很容易让人联想到Linux的内核态，但与之又毫不相干

CUDA 的“核”指的是计算核心 (Core)，而不是操作系统内核 (Kernel)。这个命名源于 GPU 硬件本身的结构。在 GPU 架构中，执行计算任务的基本单元叫做“核心”（Core）。英伟达 GPU 里有成百上千个 CUDA 核心，AMD GPU 里有流处理器，华为昇腾 NPU 里有 AI Core。因此，一个在“核心”上执行的函数，就被自然地称为“核函数”（Kernel Function）。这个叫法更精准地区分了普通 CPU 函数和并行设备函数，并且在学术界和工业界都作为标准术语被广泛使用。它虽然借用了英文词 Kernel，但指向的是硬件核心 (Core)，而不是操作系统内核。

英伟达 GPU 里有成百上千个 CUDA 核心。这个核心是前面说到的warp吗？还是thread?

• CUDA 核心 (Core)： 这是一个物理硬件。你可以把它理解为GPU里的一个计算器或算术逻辑单元。它在一个时钟周期内只能做一个简单的运算（比如一次浮点数加法）。

• Thread (线程)： 这是一个逻辑概念。它是你写代码时，一个任务的最小执行单元。

• Warp (线程束)： 这是一个调度单位，是一个逻辑概念。它是32个Threads的集合。你可以把它理解为一个小任务组。

Warp (32个线程) 是 SM 接收和调度任务的基本单位，而 CUDA核心 只是 SM 内部用来实际做计算的工具。所以，当你听到“GPU有几千个核心”时，你应该理解为“GPU有几千个可以做简单计算的硬件单元”，而不是“同时独立运行几千个不同任务的处理器”。

__global__ __aicore__ 这里的 aicore 是否多余，是不是因为NPU上还有AI CPU单元，而GPU没有？

NVIDIA GPU的核心计算单元类型相对单一。虽然它内部有负责调度的SM（流多处理器）和负责计算的CUDA核心，但它们本质上是同一套逻辑单元。因此，当一个函数被标记为 __global__ 时，编译器和硬件非常清楚：这就是要在GPU的通用计算核心上跑，没有第二种可能。

昇腾AI处理器内部存在多种专门化的计算单元。

• AI Core。它内部集成了强大的Vector（向量）单元和Cube（立方体）单元，专门用来加速深度学习中大量的矩阵、向量运算，是执行AI算子的主力。

• AI CPU：昇腾NPU内部除了AI Core，还集成了AI CPU。它主要负责执行一些逻辑复杂的、不太适合AI Core处理的任务，比如数据预处理、后处理，或者作为某些特殊算子的备选执行单元。

正是因为存在这两种功能迥异的“核”，编译器就必须明确知道：“嘿，我这段代码到底要扔给AI Core算，还是扔给AI CPU算？”