全局变量

Host侧的全局变量与C++语言的全局变量特性相同，本节介绍的全局变量特指Device侧支持的全局变量。

Device侧的全局变量必须定义在命名空间作用域或者顶层作用域。由于异构结构的存在，用于区分Device侧和Host侧全局变量，所有Device侧的全局变量都需要关键字__gm__修饰。

__gm__ int g_device;                  //device side global variable 
int g_host;                           //host side global variable

全局变量的代码示例如下：

// built-in type global variable 
__gm__ bool bool_noinit; 
__gm__ int8_t int8_t_init = -8; 
__gm__ int32_t int32_t_arr_a[3] = {-321, -322}; 
__gm__ int64_t int64_t_arr_twodim[3][4] = {{-641, -642, -643, -644}}; 
enum DAY {MON=1,TUE,WED,THU,FRI,SAT,SUN}; 
// user defined type global variable 
__gm__ enum DAY day = FRI; 
struct SIMPLE{int32_t i32;uint8_t ui8;int16_t i16;}; 
struct COMPLEX{bool b;struct SIMPLE simple;}; 
__gm__ struct COMPLEX complex = {true, {-32, 8, -16 }}; 
union DATA{int32_t u_i32;uint64_t u_ui64;uint8_t  arr_ui8[16];}; 
__gm__ union DATA data = {.u_i32 = 32}; 
// pointers 
__gm__ int* __gm__ p;               // declares a pointer p in the gm address space  
                                    //that points to an object in the gm address space 
int * __gm__ p2;                    // declares a pointer p2 in the gm address space  
                                    //that points to an object in the private address space, 
                                    //usually on the stack 
// reference global variable 
__global__ [aicore] void kernel(__gm__ int* hbm_out) { 
  *hbm_out++ = int32_t_arr_a[1]; 
  *hbm_out++ = complex.simple.i32; 
}

由于异构程序中Device侧代码和runtime配合模式的特殊性，Device侧全局变量的使用需参见注意事项。