NPU Architecture Version 200x

This section describes the hardware constraints and recommended solutions. The corresponding product model is the Atlas inference product.

Global variable usage constraints

The NPU architecture version 200x does not support generic addressing (for address spaces such as UB, stack, and GM). Therefore, the address space at the language level must match. Different address space information cannot be converted due to syntax non-compliance. Global variables are located on the GM, and passed parameters are located on the stack. An error is reported when global variables are used to pass parameters in a function. Currently, the compiler adapts to constexpr as follows only when the optimization level is O0: Global variables are moved from the DDR to the stack. Therefore, global variables can be defined and used only with constexpr in the O0 scenario. In other scenarios, global variables are not supported.

Supported scenarios

In the O0 scenario, the reference of small-scale data (integer and floating-point types) of constexpr global variables is supported, and elements can be obtained from arrays.

constexpr int a=1;
int *pa=&a; 
// Array
constexpr uint8_t padList[4] = {0, 1, 3, 5};
__aicore__ uint64_t Compute(uint32_t padNumber){
  uint64_t regFMatrix =3;
    for (uint32_t i = 0; i < padNumber; i++) {
        regFMatrix |= uint64_t(padList[i]);    } 
 return regFMatrix;
}

Scenarios not supported

An error will be reported if global variables other than constexpr are used.

const float aa = 3.141592653589; // const is not supported. Change it to constexpr.
template<typename T>
__global__ __aicore__ void hello_world3()
{   
    AscendC::printf("global var is %f\n", aa);
};

Parent topic: Hardware Constraints