地址空间限定符

__gm__ 
  - global memory (DDR) 
__ubuf__ 
  - Vector Unified Buffer 
__ca__ 
  - Cube L0A Buffer 
__cb__ 
  - Cube L0B Buffer 
__cc__ 
  - Cube L0C Buffer 
__cbuf__ 
  - Cube L1 Buffer 
empty (Type *) 
  - private address (stack address)

地址空间限定符可以在变量声明中使用，用于指定对象分配的区域。如果对象的类型被地址空间名称限定，那么该对象将被分配在指定的地址空间中。同样地，可以通过地址空间来限定指针指向的类型，从而明确该指针所指向对象所在的地址空间。以下是具体的代码示例：

// declares a pointer p in the __gm__ address space that 
// points to an object(has int type) in the __gm__ address space 
__gm__ int *__gm__ p; 
 
__global__ [aicore] void foo(...) 
{ 
  // declares an array of 4 floats in the private address space. 
  // unlike OpenCL, CCE don't have explicit __private address space qualifier. 
  float x[4]; 
}

地址空间限定符不能用于非指针返回类型、非指针函数参数、函数类型，同一个类型上不允许使用多个地址空间限定符。以下是具体的代码示例：

// OK. 
[aicore] int f() {...} 
 
// Error. Address space qualifier cannot be used with a non-pointer return type. 
__ubuf__ [aicore] int f() { ... } 
 
// OK. Address space qualifier can be used with a pointer return type. 
__ubuf__ [aicore] 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;

// 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__ 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

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

 __global__ [aicore]  
void foo(__ca__ int * ptr) { // Error. Cannot have __ca__ 
                             // qualifier in kernel arguments 
  // OK 
  __ca__ int *ptr;  
}