如其他语言一样，会提供一些内置的宏方便用户编写程序。预定义宏一节着重介绍一些用户在做异构编程时会经常用到的宏，以及宏的解释。
__NPU_ARCH__
__NPU_ARCH__是Device侧AI Core代码中的预处理宏，用于标识AI处理器的架构版本。该宏由四位数字组成，其中前三位数字用于标识AI Core的IP核(Intellectual Property Core)类型，第四位数字标识该AI Core同一个IP核的配置版本。通过该宏，开发者可以针对不同AI处理器，差异化进行代码适配和优化。AI处理器型号和__NPU_ARCH__对应关系如下表所示：
表 1 AI处理器型号和__NPU_ARCH__的对应关系
[object Object][object Object]
[object Object]以下为通过__NPU_ARCH__控制在不同AI处理器上算子输出值舍入模式的示例。
[object Object]ASCEND_IS_AIV、ASCEND_IS_AIC
ASCEND_IS_AIV和ASCEND_IS_AIC是通过C++宏实现的条件判断语句，用于在__aicore__修饰的函数中实现代码的条件编译。基于分离模式（AIC核和AIV核分离）开发融合算子时，算子逻辑中同时涉及AIV核和AIC核的处理逻辑，并需要进行核间同步，此时需要通过ASCEND_IS_AIV/ ASCEND_IS_AIC进行AIV和AIC核代码的隔离。
[object Object]
以MatmulNzCustom算子为例，该算子在分离模式下需要分别在AIV核和AIC核上实现不同的逻辑。具体而言，AIV核负责将矩阵数据搬入Unified Buffer，完成数据的重排（将矩阵数据转换为NZ格式），并将其写入Global Memory。而AIC核则直接从Global Memory读取已经重排好的NZ格式数据，并执行矩阵乘法（Matmul）计算。由于AIV核和AIC核的代码逻辑不同，需要通过ASCEND_IS_AIV和ASCEND_IS_AIC宏进行代码隔离，确保在编译时分别生成适用于AIV核和AIC核的代码。
示例伪码如下：
[object Object]ASCENDC_CUBE_ONLY
ASCENDC_CUBE_ONLY是通过C++宏实现的条件判断语句，用于在__aicore__修饰的函数中实现代码的条件编译。
基于分离模式开发非融合算子时，在只有矩阵计算的算子场景下，可以通过设置ASCENDC_CUBE_ONLY，使能纯Cube模式完成Matmul计算，减少消息通信的性能开销，提升算子性能。
[object Object]
以matmul_custom算子为例，高阶API Matmul默认使用MIX模式，即用户从AIV侧发起消息，通过消息通信框架中转消息后，在AIC侧执行Matmul计算。这套消息处理机制会带来额外的Scalar性能开销。相较于MIX模式，纯Cube模式可以直接跳过消息通信框架，完成Matmul计算，提升算子性能。
示例伪码如下：
[object Object]
函数执行空间限定符（Function Execution Space Qualifier）指示函数是在Host侧执行还是在Device侧执行，以及它是否可从Host侧或Device侧调用。
__global__
__global__执行空间限定符声明一个Kernel函数。Kernel函数有如下性质：在Device上执行；只能被Host侧函数调用；__global__只是表示这是Device侧函数的入口，并不表示具体的设备类型，具体的设备类型由__aicore__标记。具有如下使用约束：
- 一个__global__函数必须返回void类型，并且不能是class的成员函数。
- 主机侧调用__global__函数必须使用<<<>>>异构调用语法。
- __global__的调用是异步的，意味着函数返回，并不表示kernel函数在device侧已经执行完成，如果需要同步，需要使用Runtime同步接口显式同步，如aclrtSynchronizeStream接口。
__aicore__
__aicore__执行空间限定符声明一个函数，它具有如下属性：
- 在Device侧执行
- 只能被__global__函数，或者其他__aicore__函数调用
[object Object]__host__
__host__执行空间限定符声明一个函数，它具有如下属性：
- 只能在Host侧执行
- 只能被Host侧函数调用
- __global__ 和__host__不能一起使用
__host__限定符是可选项，无函数执行空间限定符定义的函数，默认是host函数。
[object Object]__aicpu__
AI CPU函数执行空间限定符__aicpu__用于指示函数是否为AI CPU Kernel函数，它具有如下属性：
- 在Device侧执行且只能被Host侧函数调用，因此必须与__global__同时声明。
- 一个__global__ __aicpu__函数不能是void返回类型，并且入参只能是一个指针。
- 一个__global__ __aicpu__函数不能在.asc文件中进行定义，只能声明，且需要使用extern。
- Host侧调用__global__ __aicpu__函数时必须使用<<<>>>异构调用语法，输入的函数入参在入参指针的基础上需要输入从指针中读取的数据大小。
- __global__的调用是异步的，意味着函数返回，并不表示kernel函数在Device侧已经执行完成，如果需要同步，需要使用Runtime同步接口显式同步，如aclrtSynchronizeStream接口
[object Object][object Object]__inline__
__inline__限定符声明一个函数，它具有如下属性：
- 标识Device侧函数强制内联，可以减少函数频繁调用产生的指令压栈、出栈的开销，但可能会导致算子二进制增加。
- 和C++函数修饰符inline的主要区别是Device侧__inline__是强制内联，C++的inline则是根据编译器优化选择性内联。
- AI Core对函数嵌套深度有限制，一般推荐嵌套深度不超过4层。使用强制内联可以减少调用层次。
__cube__
标识该核函数仅在Cube核执行。针对耦合模式的硬件架构，该修饰符不生效。
[object Object]__vector__
标识该核函数仅在Vector核执行。针对耦合模式的硬件架构，该修饰符不生效。
[object Object]__mix__(cube, vec)
标识该核函数同时在Cube核和Vector核上执行。(cube, vec)分别表示核函数启动的Cube核和Vector核的配比，支持的配比为(1, 0)，(0, 1)，(1, 1)， (1, 2)。针对耦合模式的硬件架构，该修饰符不生效。
__schedmode__(mode)
标识该核函数的执行调度模式。如下图所示：
mode = 0 : normal mode，尽可能选择空闲物理核下发执行核函数，若空闲物理核数无法满足当前核函数的需要，没有下发的部分等待核心空闲后执行。此时OP1和OP2算子会存在交叠执行（overlap）的情况。
mode = 1 : batch mode，在下发核函数时先进行判断，若空闲物理核数无法满足当前核函数的需要，则等待至空闲物理核数满足该核函数所需要的所有物理核时，同时下发执行，OP1和OP2的执行被切分（split）开，不会出现交叠执行的情况。
在多流并发场景，多算子并行执行时，若执行总核数超过最大物理核数，且多个算子逻辑使用SyncALL等核间同步接口时，建议设置mode为1，防止多个算子之间互相等待空闲核调度，导致死锁。默认值mode为0。
[object Object]
__simd_vf__
函数标记宏，用于标记SIMD VF入口函数，函数无返回值。使用asc_vf_call调用SIMD VF入口函数，启动VF子任务。
[object Object]__simd_vf__标记的SIMD VF有以下入参约束：
支持指针传参（Pass-by-Pointer），指针变量必须用__ubuf__地址空间限定符修饰。
不支持引用传参（Pass-by-Reference）。
不支持函数指针传参，函数对象。
__simd_vf__使用的示例如下：
[object Object]
__simd_callee__
函数标记宏，函数可以有返回值，允许被SIMD VF入口函数或其他非入口函数调用。
[object Object]
AI Core具备多级独立片上存储，各个地址空间独立编址，具备各自的访存指令，根据架构差异，有些存储空间具备统一地址空间（Generic Address Space），有些则没有。设备侧编程基于语法扩展允许地址空间作为合法的类型限定符，以提供针对不同地址空间的访问能力和地址空间合法性检查。
表 2 地址空间映射关系
[object Object][object Object]
[object Object]地址空间限定符可以在变量声明中使用，用于指定对象分配的区域。如果对象的类型被地址空间名称限定，那么该对象将被分配在指定的地址空间中。同样地，对于指针，指向的类型可以通过地址空间进行限定，以指示所指向的对象所在的地址空间。
地址空间限定符不能用于非指针返回类型，非指针函数参数，函数类型，同一个类型上不允许使用多个地址空间限定符。
[object Object]
private地址空间
private地址空间是大多数变量的默认地址空间，特别是局部变量。
[object Object]__gm__地址空间
__gm__地址空间限定符用来表示分配于设备侧全局内存的对象，全局内存对象可以声明为标量、用户自定义结构体的指针。
[object Object]__ubuf__地址空间
__ubuf__地址空间用来描述存储于AI Core核内UB存储空间的变量。
[object Object]__ca__, __cb__, __cc__, __cbuf__地址空间
上述几个地址空间主要用于特定的DMA指令访问，不具备标量直接访问能力。
[object Object]
[object Object][object Object]
[object Object][object Object][object Object]
[object Object]通常，建议用户使用内置变量对应的API获取所需值，不建议用户直接使用内置变量。因为内置变量反映的是单个硬件资源的配置信息，对于软件栈整合硬件资源、扩展硬件的功能，内置变量的值与实际语义可能不符。
例如，在Atlas 推理系列产品中，当启用KERNEL_TYPE_MIX_VECTOR_CORE时，算子会同时运行在AI Core和Vector Core上。此时，block_idx在这两种核心上都是从0开始计数，用户无法直接通过block_idx来切分数据和控制多核逻辑。而GetBlockIdx在Vector Core上对block_idx增加偏移量（AI Core的block_num），从而保证返回的值能够正确反映多核环境下的实际逻辑。
SIMD与SIMT混合编程场景中，SIMT VF的入口函数使用__simt_vf__进行标识，通过在SIMD的__aicore__函数中使用调用SIMT入口函数。被SIMT VF入口函数调用的函数使用__simt_callee__进行标识。
__simt_vf__ [object Object][object Object]
函数标记宏，用于标记SIMT VF入口函数，函数无返回值。使用asc_vf_call接口调用SIMT VF入口函数，启动VF子任务。
[object Object]__simt_vf__标记的SIMT VF函数支持的参数类型如下。__simt_vf__的使用示例请参考。
- 指针类型：__ubuf__ *、__gm__ *；
- 标量类型：bool、int8_t、uint8_t、int16_t、uint16_t、half、bfloat16、int32_t、uint32_t、float、int64_t、uint64_t。
__simt_callee__
函数标记宏，用于标记SIMT VF非入口函数，函数可以有返回值，允许被SIMT VF入口函数或其他非入口函数调用。
[object Object]
Ascend C为SIMT编程、SIMD与SIMT混合编程提供了布尔、整形、浮点型的标量数据类型和短向量数据类型，提供了用于表达线程块、线程网格三维信息的内置变量。 关于内置数据格式的详细说明请参见，内置变量请参见。