SIMD Built-In Keywords

Predefined Macros

Like other languages, some built-in macros are provided for users to write programs. This part describes some macros that are frequently used during heterogeneous programming.

• __NPU_ARCH__
  __NPU_ARCH__ is a preprocessing macro in the AI Core code on the device, which is used to identify the architecture version of the AI processor. This macro consists of four digits. The first three digits indicate the type of the intellectual property core (IP core) of the AI Core, and the fourth digit indicates the configuration version of the same IP core of the AI Core. With this macro, you can differentially adapt and optimize code for different AI processors. The following table lists the mapping between AI processor models and __NPU_ARCH__.

  Table 1 Mapping between AI processor models and __NPU_ARCH__

  AI processor Model	__NPU_ARCH__
  Atlas 350 Accelerator Card	3510
  Atlas A3 training product / Atlas A3 inference product	2201
  Atlas A2 training product / Atlas A2 inference product	2201
  Atlas 200I/500 A2 inference product	3002
  Atlas inference product	2002
  Atlas training product	1001

  The following is an example of using __NPU_ARCH__ to control the rounding mode of the operator output value on different AI processors:

  1 2 3 4 5 6 7 8 9 10 11 12

  __aicore__ static inline void CopyOut(uint64_t mulLen) { #if __NPU_ARCH__ == 2002 Cast(dstLocal, srcLocal, RoundMode::CAST_NONE, mulLen); // CAST_NONE indicates that the CAST_RINT mode is used when precision loss occurs during conversion. If precision loss is not involved, rounding is not performed. #elif __NPU_ARCH__ == 2201 Cast(dstLocal, srcLocal, RoundMode::CAST_RINT, mulLen); // CAST_RINT indicates the banker's rounding mode. #endif event_t eventVToMTE3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3)); SetFlag<HardEvent::V_MTE3>(eventVToMTE3); WaitFlag<HardEvent::V_MTE3>(eventVToMTE3); CommonCopyOut<float>(dstLocal, mulLen); // Copy a LocalTensor to a GlobalTensor. }

• ASCEND_IS_AIV, ASCEND_IS_AIC

  ASCEND_IS_AIV and ASCEND_IS_AIC are conditional statements implemented using C++ macros. They are used to implement conditional compilation of code in functions decorated with __aicore__. When developing a fused operator in separated mode (AIC and AIV cores are separated), the operator logic involves both AIV and AIC core processing logic, and inter-core synchronization is required. In this case, ASCEND_IS_AIV and ASCEND_IS_AIC are used to isolate the AIV and AIC core code.

  

  When Matmul high-level APIs are used, the AIV and AIC core code has been isolated using the REGIST_MATMUL_OBJ macro. You do not need to use this macro for processing.

  Take the MatmulNzCustom operator as an example. In separated mode, this operator needs to implement different logic on the AIV core and AIC core. Specifically, the AIV core is responsible for moving cube data into the Unified Buffer, rearranging the data (converting the cube data into the NZ format), and writing it to the global memory. The AIC core directly reads the rearranged NZ-format data from the global memory and performs cube multiplication (Matmul) computation. Because the code logic of the AIV core is different from that of the AIC core, the ASCEND_IS_AIV and ASCEND_IS_AIC macros are used to isolate the code to ensure that the code applicable to the AIV core and that applicable to the AIC core are generated separately during compilation.

  The example pseudocode is as follows:

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34

  template <typename AType, typename BType, typename CType, typename BiasType> __aicore__ inline void MatmulKernel<AType, BType, CType, BiasType>::Process(AscendC::TPipe *pipe) { // Use the vector unit of the AIV core to implement ND2NZ format conversion. In the following code, MatrixBtoNZ is the function for performing ND2NZ format conversion on cube B. if ASCEND_IS_AIV { pipe->InitBuffer(ubBuf, TOTAL_UB_SIZE); MatrixBtoNZ<typename B_TYPE::T>(tempGM, bGMNZ, tiling, isTransB, ubBuf, tiling.baseK, tiling.baseN); // ND2NZ function implemented on the vector SyncAll(); // Synchronization between the AIC and AIV cores AscendC::CrossCoreSetFlag<0x2, PIPE_MTE3>(0x4); return; } if ASCEND_IS_AIC { AscendC::CrossCoreWaitFlag(0x4); // Wait for the AIV core to complete ND2NZ format conversion. } ... ... // Set the left cube A, right cube B, and bias. matmulObj.SetTail(tailM, tailN); matmulObj.SetTensorA(aGlobal, false); matmulObj.SetTensorB(bGlobal, false); if (tiling.isBias) { matmulObj.SetBias(biasGlobal); } // Complete the cube multiplication operation. matmulObj.IterateAll(cGlobal); // End the cube multiplication operation. matmulObj.End(); }

• ASCENDC_CUBE_ONLY

  ASCENDC_CUBE_ONLY is a conditional statement implemented using C++ macros. It is used to implement conditional compilation of code in functions decorated with __aicore__.

  When developing non-fused operators in separated mode, you can set ASCENDC_CUBE_ONLY to enable the cube-only mode for Matmul computation in scenarios where only cube computation is involved. This reduces the performance overhead of message communication and improves operator performance.

  

  The ASCENDC_CUBE_ONLY macro must be set before #include "lib/matmul_intf.h".

  Take the matmul_custom operator as an example. By default, the Matmul high-level API uses the MIX mode. That is, a user initiates a message from the AIV, and Matmul computation is performed on the AIC after the message is forwarded by the message communication framework. This message processing mechanism causes extra scalar performance overhead. Compared with the MIX mode, the CUBE_ONLY mode can directly skip the message communication framework to complete Matmul computation, improving operator performance.

  The example pseudocode is as follows:

  1 2 3 4 5 6 7 8

  #define ASCENDC_CUBE_ONLY #include "lib/matmul_intf.h" using A_TYPE = AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, AType>; using B_TYPE = AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BType>; using C_TYPE = AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, CType>; using BIAS_TYPE = AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BiasType>; AscendC::Matmul<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, CFG_NORM> matmulObj;

Function Execution Space Qualifier

The function execution space qualifier indicates whether the function is executed on the host or device, and whether the function can be called from the host or device.

• __global__

  The execution space qualifier __global__ declares a kernel function. The kernel function is executed on the device, and called only by the function on the host. __global__ only indicates the entrypoint of the function on the device, rather than the specific device type. The specific device type is marked by __aicore__. The restrictions are as follows:

  • A __global__ function must return the void type, and cannot be a member function of class.
  • The host must call the __global__ function using the <<<>>> heterogeneous call syntax.
  • The __global__ call is asynchronous. The return of the function does not mean that the kernel function has been completely executed on the device. For explicit synchronization, use the runtime synchronization API, such as aclrtSynchronizeStream.
• __aicore__

  The execution space qualifier __aicore__ declares a function with the following attributes:

  • Executed on the device
  • Called only by the __global__ function or other __aicore__ functions

  // Only callable from device functions with same kind
  // of execution space
  __aicore__ void bar() {}
  
  // Define a kernel function execute on AI Core device
  __global__ __aicore__ void foo() {
    bar(); // OK.
  }

• __host__

  The execution space qualifier __host__ declares a function with the following attributes:

  • Executed only on the host
  • Called only by functions on the host
  • __global__ and __host__ cannot be used together.

  The __host__ qualifier is optional. If a function is not defined with any execution space qualifier, it is a host function by default.

  __aicore__ int f() {}
  
  // defines a host side function
  int foo() {}
  
  // defines a host side function
  __host__ int bar() {
    f();     // Error.
    foo();   // OK.
  }
  
  // Error.
  __global__ __host__ void kfunc() {}

• __aicpu__

  The function execution space qualifier __aicpu__ of the AI CPU indicates whether a function is an AI CPU kernel function. It has the following attributes:

  • The function is executed on the device and can be called only by functions on the host. Therefore, it must be declared with __global__.
  • A __global__ __aicpu__ function cannot return the void type, and the input parameter can only be a pointer.
  • A __global__ __aicpu__ function cannot be defined in the .asc file. It can only be declared and must be declared using extern.
  • When the __global__ __aicpu__ function is called on the host, the <<<>>> heterogeneous call syntax must be used. The input function parameter must include the size of the data read from the pointer based on the input parameter pointer.
  • The __global__ call is asynchronous. The return of the function does not mean that the kernel function has been completely executed on the device. For explicit synchronization, use the runtime synchronization API, such as aclrtSynchronizeStream.

  1 2 3 4 5 6 7

  // Define an AI CPU kernel function in the AI CPU device file. __aicpu__ void foo() {} // Error, single __aicpu__ identifier without _global__ __global__ void foo() {} // Error, single __global__ identifier without __aicpu__ __global__ __aicpu__ void foo() {} // Error, return type is void __global__ __aicpu__ int foo(void *a) {} // OK __global__ __aicpu__ int foo(int a) {} // Error, input param is not pointer __global__ __aicpu__ int foo(void *a, void *b) {} // Error, input param num is not one

  1 2 3

  // Declare an AI CPU kernel function in the .asc file. extern __global__ __aicpu__ uint32_t hello_world(void *args);// OK }

• __inline__

  The __inline__ qualifier declares a function with the following attributes:

  • Indicates that the function on the device is forcibly inlined. This reduces the overhead of instruction stack push and pop caused by frequent function calls, but may increase the operator binary size.
  • The main difference between __inline__ on the device and the C++ function modifier inline is that __inline__ indicates forcible inline on the device, while inline in C++ indicates selective inline based on compiler optimization.
  • AI Core has restrictions on the function nesting depth. It is recommended that the nesting depth be less than or equal to four layers. Using forcible inline can reduce the number of call layers.

• __cube__

  This qualifier indicates that the kernel function is executed only on the Cube Core. This qualifier does not take effect for the hardware architecture in coupled mode.

  1 2 3 4 5 6

  extern "C" __global__ __cube__ void mmad_custom(GM_ADDR a, GM_ADDR b, GM_ADDR c) { KernelMmad op; op.Init(a, b, c); op.Process(); }

• __vector__

  This qualifier indicates that the kernel function is executed only on the Vector Core. This qualifier does not take effect for the hardware architecture in coupled mode.

   __vector__ __global__ __aicore__ void add_custom(){}

• __mix__(cube, vec)

  This qualifier indicates that the kernel function is executed on both the Cube and Vector Cores. (cube, vec) indicates the ratio of Cube Cores to Vector Cores started by the kernel function. The supported ratios are (1, 0), (0, 1), (1, 1), and (1, 2). This qualifier does not take effect for the hardware architecture in coupled mode.

• __schedmode__(mode)

  This qualifier indicates the execution scheduling mode of the kernel function. The following figure shows the details:

  • mode = 0: normal mode. The system selects idle physical cores to deliver and execute the kernel function. If the number of idle physical cores cannot meet the requirements of the current kernel function, the remaining part of the kernel function is executed after the cores become idle. In this case, the execution of operators OP1 and OP2 may overlap.
  • mode = 1: batch mode. Before delivering the kernel function, the system checks whether the number of idle physical cores meets the requirements of the current kernel function. If not, the system waits until the number of idle physical cores meets the requirements of the kernel function, and then delivers and executes the kernel function at the same time. The execution of OP1 and OP2 is split, and no overlapping execution occurs.

  

  In multi-stream concurrency scenarios, if the total number of execution cores exceeds the maximum number of physical cores and multiple operators use inter-core synchronization APIs such as SyncALL, you are advised to set mode to 1 to prevent multiple operators from waiting for idle core scheduling from each other, which may cause a deadlock. The default value of mode is 0.

  1 2 3 4 5 6 7 8 9 10 11 12 13

  __schedmode__(1) __global__ __mix__(1, 2) void OP1() // OP1 uses the SyncAll API and may involve multi-stream concurrency. Therefore, the batch mode (mode 1) needs to be set. { AscendC::SyncAll(); .... } __schedmode__(1) __global__ __mix__(1, 2) void OP2() // OP2 uses the SyncAll API and may involve multi-stream concurrency. Therefore, the batch mode (mode 1) needs to be set. { AscendC::SyncAll(); .... } __schedmode__(0) __global__ __vector__ void OP3() {...} // OP3 does not use the SyncAll API. Therefore, the normal mode (mode 0) can be set, and the operator is executed according to the normal rules. or __global__ __vector__ void OP3() {...} // If __schedmode__ is not set, the normal mode is used by default.

Function Marking Macro

• __simd_vf__
  Function marking macro, which is used to mark the SIMD VF entry function. The function does not return a value. Use asc_vf_call to call the SIMD VF entry function and start the VF subtask.

  1	__simd_vf__ inline void KernelAdd(__ubuf__ float* x, __ubuf__ float* y, __ubuf__ float* z)

  The SIMD VFs marked by __simd_vf__ has the following input parameter restrictions:

  • Parameter pass by pointer (pass-by-pointer) is supported. Pointer variables must be modified by the address space qualifier __ubuf__.
  • Parameter pass by reference (pass-by-reference) is not supported.
  • Parameter pass by function pointer is not supported.

    The sample of using __simd_vf__ is as follows:

    1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17

    __simd_vf__ inline void simd_adds(__ubuf__ float *output, __ubuf__ float *input, uint32_t count, uint16_t one_repeat_size, uint16_t repeat_times) { AscendC::Reg::RegTensor<float> src_reg0; AscendC::Reg::RegTensor<float> dst_reg0; // asc_update_mask() will be supported later. // init MaskReg with the count of all numbers. AscendC::Reg::MaskReg mask_reg = AscendC::Reg::UpdateMask<float>(count); for (uint16_t i = 0; i < repeat_times; i++) { // asc_load, asc_adds and asc_store will be supported later. // load data from UB to RegTensor. AscendC::Reg::LoadAlign(src_reg0, input + i * one_repeat_size); AscendC::Reg::Adds(dst_reg0, src_reg0, 1.0f, mask_reg); // store data from RegTensor to UB. AscendC::Reg::StoreAlign(output + i * one_repeat_size, dst_reg0, mask_reg); } }

• __simd_callee__

  Function marking macro. The function can have a return value and can be called by the SIMD VF entry function or other non-entry functions.

  1	__simd_callee__ inline float add(float x, float y)

Address Space Qualifiers

The AI Core has multi-level independent on-chip storage. Each address space has independent addressing and its own memory access instructions. Depending on the architecture, some storage spaces have a generic address space, while others do not. On-device programming allows address spaces to be used as valid type qualifiers based on syntax extensions, providing access to different address spaces and check of address space validity.

Address space qualifiers can be used in variable declarations to specify the region in which an object is allocated. If the type of an object is qualified by an address space name, the object will be allocated in the specified address space. Similarly, for pointers, the pointing type can be qualified by an address space to indicate the address space in which the pointed-to object resides.

// declares a pointer p in the __gm__ address space that
// points to an object(has int type) in the __gm__ address space
__gm__ int *p;

__global__ __aicore__ void foo(...)
{
  // declares an array of 4 floats in the private address space.
  float x[4];
}

Address space qualifiers cannot be used for non-pointer return types, non-pointer function parameters, or function types. Multiple address space qualifiers cannot be used for the same type.

// OK.
__aicore__ int f() {...}

// Error. Address space qualifier cannot be used with a non-pointer return type.
__ubuf__ int f() { ... }

// OK. Address space qualifier can be used with a pointer return type.
__ubuf__ int *f() { ... }

// Error. Multiple address spaces specified for a type.
__ubuf__ __gm__ int i;

// OK. The first address space qualifies the object pointed to and the second
// qualifies the pointer.
__ubuf__ int * __gm__ ptr;

Important: The pointer size may vary with the address space. For example, sizeof(__gm__ int *) may not always be equal to sizeof(__ubuf__ int *). The compiler may store __ubuf__ pointers in 32-bit mode on some systems.

• Private address space

  The private address space is the default address space for most variables, especially local variables.

  // m is in a specific kernel parameter address space, 
  // it's physical location is implementation determined.
  __global__ __aicore__ void foo(int m) {
    // OK. i is an int variable allocated in private address space
    int i;
  }
  
  __aicore__ void bar(int k) { //OK. k is in private address space
    // OK. i is an int variable allocated in private address space
    int i; 
  }

• __gm__ address space

  The __gm__ address space qualifier is used to indicate objects allocated to the global memory on the device. Global memory objects can be declared as scalars or pointers to user-defined structures.

  __gm__ int *var; // var point to an array of int elements
  
  typedef struct {
      float a[3];
      int b[2];
  } foo_t;
  
  __gm__ foo_t *info; // info point to an array of foo_t elements

• __ubuf__ address space

  The __ubuf__ address space is used to describe variables stored in the UB storage space of the AI Core.

  __global__ __aicore__ void foo() {
    // ptr is in private address space, point to __ubuf__
    __ubuf__ int *ptr;
  }

• __ca__, __cb__, __cc__, and __cbuf__ address spaces

  The preceding address spaces are mainly used for access by specific DMA instructions and do not support direct access by scalars.

  class ObjTy{
    ObjTy(){...}
    void print(){...}
  
  private:
    int a;
    int b;
  };
  
  __global__ __aicore__ 
  void foo(__ca__ int * ptr) { // Error. Cannot have __ca__
                               // qualifier in kernel arguments
    // OK
    __ca__ int *ptr; 
  }

Built-in Constants

Built-in Variables

Generally, you are advised to use the APIs corresponding to the built-in variables to obtain the required values. You are advised not to directly use the built-in variables. This is because a built-in variable reflects the configuration information of a single hardware resource. For functions of integrating hardware resources and extending hardware by the software stack, the values of built-in variables may not match the actual semantics.

For example, in Atlas inference product , when KERNEL_TYPE_MIX_VECTOR_CORE is enabled, the operator runs on both the AI Core and Vector Core. In this case, block_idx starts counting from 0 on both cores, and users cannot directly use block_idx to split data and control multi-core logic. GetBlockIdx adds an offset (block_num of the AI Core) to block_idx on the Vector Core to ensure that the returned value correctly reflects the actual logic in a multi-core environment.

SIMD and SIMT Programming Scenario

In the SIMD and SIMT programming scenario, the entry function of the SIMT VF is identified by __simt_vf__, and asc_vf_call is used in the __aicore__ function of the SIMD to call the SIMT entry function. The function called by the SIMT VF entry function is identified by __simt_callee__.

• __simt_vf__
  Function marking macro, which is used to mark the SIMT VF entry function. The function does not return a value. Use asc_vf_call to call the SIMT VF entry function and start the VF subtask.

  1	__simt_vf__ inline void KernelAdd(__gm__ float* x, __gm__ float* y, __gm__ float* z)

  The SIMT VF function marked by __simt_vf__ supports the following parameter types. For the usage example of __simt_vf__, see Programming Sample.

  • Pointer types: __ubuf__ *, __gm__ *;
  • Scalar types: bool, int8_t, uint8_t, int16_t, uint16_t, half, bfloat16, int32_t, uint32_t, float, int64_t, uint64_t;
• __simt_callee__

  Function marking macro, which is used to mark SIMT VF non-entry functions. The function can have a return value and can be called by the SIMT VF entry function or other non-entry functions.

  1	__simt_callee__ inline float add(float x, float y)

Ascend C provides Boolean, integer, and floating-point scalar data types and short vector data types for SIMT programming and SIMD and SIMT programming, and provides built-in variables for expressing three-dimensional information of thread blocks and thread grids. For details about the built-in data formats, see Built-in Data Types. For details about the built-in variables, see Built-in Variables.