---
title: NPU架构版本200x
description: "本节介绍硬件约束以及解决方案建议。对应的产品型号为： Atlas 推理系列产品 。"
url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/programug/Ascendcopdevg/atlas_ascendc_10_00062.html
sourcePath: /source/zh/canncommercial/900/programug/Ascendcopdevg/atlas_ascendc_10_00062.html
indexId: 2fc90f93399c3ec83f1329301bbac98360b0959480d6ef8b79d46c31fbe4075481
---
# 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); };