NPU架构版本200x

本节介绍硬件约束以及解决方案建议。对应的产品型号为：Atlas 推理系列产品。

• 全局变量使用约束

  NPU架构版本200x不支持Generic Addressing通用寻址（针对UB、stack、GM等地址空间），因此语言层面地址空间必须匹配。不同的地址空间信息不允许转换，不符合语法。全局变量位于GM上，传参位于stack上，函数内传参使用全局变量时会报错。目前编译器仅在优化级别为O0的场景下，对constexpr做了适配处理：将全局变量先从DDR上放到stack上。所以全局变量仅支持在O0下使用constexpr进行定义和使用，在其他场景均不支持。

  • 支持场景

    在O0下支持constexpr全局变量小规模类型（整型、浮点型）数据的引用，支持数组取元素。

    1 2 3 4 5 6 7 8 9 10

    constexpr int a=1； int *pa=&a; // 数组 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; }

  • 不支持场景

    对于使用constexpr之外的全局参数将报错。

    1 2 3 4 5 6

    const float aa = 3.141592653589; // 不支持const 应修改为constexpr template<typename T> __global__ __aicore__ void hello_world3() { AscendC::printf("global var is %f\n", aa); };