8. C++语言扩展¶

8.1. 函数执行空间指示符 (Function Execution Space Specifiers)¶

函数执行空间指示符表示函数是在主机上执行还是在设备上执行，以及它是否可从主机或设备调用。

global¶

__global__ 执行空间指示符声明一个函数为入口核函数。这样的函数可以：

在计算设备上执行
仅可从主机端程序调用

一个 __global__ 函数必须具有 void 返回类型，并且不能是类的成员函数。

对 __global__ 函数的任何调用都必须按照执行配置的描述进行指定。

对 __global__ 函数的调用总是异步的，这意味着在设备完成执行之前它就会返回。

device¶

__device__ 执行空间指示符声明一个函数：

在计算设备上执行
仅可从设备端调用，也是从其他__global__或__device__核函数调用

__global__和__device__执行空间指示符不能用来同时修饰同一个函数。

host¶

__host__执行空间指示符声明一个函数：

在主机端上执行
仅可从主机端程序调用

当一个函数不加任何执行空间指示符时，就等同于加了__host__指示符，此时该函数只会被编译到主机端程序中。

__global__和__host__执行空间指示符不能用来同时修饰同一个函数。

__device__和__host__执行空间指示符可以用来同时修饰同一个函数，表示这个函数既可以在主机端被调用，也可以在设备端被调用。预定义宏__AGCU_ARCH__可以用来区分不同的设备端和主机端代码。

__host__ __device__ void func(void) {
#if __AGCU_ARCH__ >= 300
    // 兼容AGCU300的设备端代码
#elif __AGCU_ARCH__ >= 200
    // 兼容AGCU200的设备端代码
#elif !defined(__AGCU_ARCH__)
    // 主机端代码
#endif
}

cooperative¶

__cooperative__指示符必须和__global__一起使用，声明一个核函数：

所有线程块必须以协作方式同时启动
因而可以在整个线程网格内进行同步和通信

注意：被标记为__cooperative__的核函数，启动时线程块的数量不能超过当前设备上的计算单元集群数量。

8.2. 变量内存空间指示符（Variable Memory Space Specifiers）¶

变量内存空间指示符表示设备上变量的内存位置。

TOPS C++内存模型中定义了以下内存空间：通用内存空间、全局内存空间、共享内存空间、本地内存空间和常量内存空间。变量内存空间指示符可用于C/C++内置类型变量的定义，也可以用于结构体和对象的定义，以指定所分配的内存空间。TOPS C++扩展了C/C++语法的类型指示符，将内存空间名称作为有效的类型指示符。

在设备端代码中，一个不加__device__、__shared__和__constant__变量内存空间指示符的自动变量，对应到私有内存空间（Private Memory Space），且大概率会被编译器优化到寄存器中。当不使用变量内存空间指示符定义变量时：

将会分配在线程的私有内存空间中（容量限制参考《AGCU技术参数》中的每个线程的最大堆栈容量）
其生命期等同于所属线程的生命期
在不同的线程中对应不同的内存实例
只能被一个线程访问
每个线程中的私有变量可能会被分配到不同的内存地址上

定义指针时，如果没有显式声明内存空间指示符，则指针默认指向通用内存空间。内存空间A的指针只能赋值给指向和内存空间A相同内存空间的指针，或者父内存空间的指针。如果指针A、B的内存空间类型不相同，且B不是父内存空间类型，则将内存空间A的指针强制转换为内存空间B的指针是非法的。在TOPS C++中，除了常量内存空间，其它内存空间都是通用内存空间的子空间。

函数形参中包含指针时，不可以加内存空间指示符，默认指向通用内存空间。

device¶

变量内存空间指示符__device__表示对应的变量内存在设备内存中。

__device__可以和后续的指示符一起使用。当单独使用__device__指示符定义变量时：

将会分配在全局内存空间中
其生命期等同于设备上下文的生命期
在不同的设备上对应不同的内存实例
既可以在一个计算网格中的所有线程中访问，也可以通过运行时库的API从主机端访问（topsGetSymbolAddress / topsGetSymbolSize / topsMemcpyToSymbol / topsMemcpyFromSymbol）

注意：变量内存空间指示符__device__和函数执行空间指示符__device__是不同的。

定义设备端的全局变量：

__device__ float global_var = 1.0f;

__device__ void kernel_func(void) {
    printf("%f\n", global_var);
}

声明设备端的全局变量：

extern __device__ float extern_global_var;

__device__ void kernel_func(void) {
    printf("%f\n", extern_global_var);
}

定义设备端的静态变量：

static __device__ float static_var = 1.0f;

__device__ void kernel_func(void) {
    static __device__ static_var2 = 1.0f;

    printf("%f %f\n", static_var, static_var2);
}

在主机端设置__device__变量：

__device__ float device_var;

void test(void) {
    float host_var = 1.0f;
    topsMemcpyToSymbol(&device_var, &host_var, sizeof(host_var));
}

constant¶

当使用变量内存空间指示符__constant__定义变量时：

将会分配在设备端的常量内存空间中（容量限制参考《AGCU技术参数》中的每个核函数的最大常量数据段容量）
其生命期等同于设备上下文的生命期
在不同的设备上对应不同的内存实例
既可以在一个计算网格中的所有线程中访问，也可以通过运行时库的API从主机端访问（topsGetSymbolAddress / topsGetSymbolSize / topsMemcpyToSymbol / topsMemcpyFromSymbol）

定义设备端的全局常量：

__constant__ float const_var = 1.0f;

__device__ void kernel_func(void) {
    printf("%f\n", const_var);
}

指示符__constant__可以和指示符__device__一起使用。

__device__ __constant__ float device_const_var = 1.0f;

__device__ void kernel_func(void) {
    printf("%f\n", device_const_var);
}

在主机端设置__constant__变量：

__constant__ float const_var;

void host_func(void) {
    float host_var = 1.0f;
    topsMemcpyToSymbol(&const_var, &host_var, sizeof(host_var));
}

shared¶

当使用变量内存空间指示符__shared__定义变量时：

将会分配在线程块的共享内存空间中（容量限制参考《AGCU技术参数》中的每个线程块的最大共享内存容量）
其生命期等同于所属线程块的生命期
在不同的线程块中对应不同的内存实例
只能被同一个线程块中的所有线程访问
每个线程块中的共享变量可能会被分配到不同的内存地址上

在核函数内定义确定大小的共享变量或数组，则在编译器可以自动计算出共享内存的大小，会在启动核函数的时候自动分配：

__device__ void kernel_func(void) {
    __shared__ float shared_var;
    if (threadIdx.x == 0) {
        shared_var = 1.0f;
    }
}

当开发者无法在编译期确定共享数组的大小时，可以定义一个extern的共享变量：

extern __shared__ float shared_array[];

__global__ void kernel_func(void) {
    int* array0 = (int*)shared_array;
    float* array1 = (float*)&array0[127];
}

int main() {
    // 启动一个包含动态共享内存的核函数
    int shared_mem_size = 128 * 1024;
    kernel_func<<<1, N, shared_mem_size>>>();
}

extern __shared__是动态共享变量，必须按照不定长数组的方式定义，且在一个作用域中只能定义一个。如果一个核函数使用了动态共享内存，则在启动时必须指定共享内存的大小。

local¶

当使用变量内存空间指示符__local__定义变量时：

将会分配在线程的本地内存空间中（容量限制参考《AGCU技术参数》中的每个线程的最大本地内存容量）
其生命期等同于所属线程的生命期
在不同的线程中对应不同的内存实例
只能被一个线程访问
每个线程中的本地变量可能会被分配到不同的内存地址上
本地内存主要是用来存放参与计算的向量和张量数据

定义一个核函数内部的本地变量：

__device__ void kernel_func(void) {
    __local__ __vector float shared_var[128];
}

restrict¶

__restrict__关键字可以用来定义受限指针（restrict pointer）。受限指针一般和const修饰符一起使用，用来标记函数的形参指针在函数体内部不会存在别名。比如以下代码中，指针a指向的值不能缓存在寄存器中，每次都必须从指针指向的内存中读入，否则计算结果是错误的。

__device__ void kernel_func(const float* a, const float* b, float* c) {
    *a++;
    *b++;
    *c = *a + *b;
}

__device__ void kernel_func2(void) {
    float a = 1.0f;
    kernel_func(&a, &a, &a);
}

以下代码使用__restrict__修饰了形参指针，告诉编译器指针a在函数kernel_func中不会存在别名：

__device__ void kernel_func(const float* __restrict__ a, const float* __restrict__ b, float* c) {
    *a++;
    *b++;
    *c = *a + *b;
}

__device__ void kernel_func2(void) {
    float a = 1.0f;
    float b = 1.0f;
    float c = 0.0f;
    kernel_func(&a, &b, &c);
}

注意：在加上__restrict__指示符后，编译器并不会检查调动时的实参是否指向了相同的内存空间，需要开发者自己保证这一点。

注意：开发者应该尽量将函数形参指针标记为const __restrict__，这样可以给编译器和硬件更多的优化空间，降低访问内存的频率，并且可以降低寄存器分配压力。

转换规则¶

允许非通用地址空间的指针到通用地址空间指针的隐式转换

__local__ char* ptr0；
char* ptr1 = ptr0;

通用地址空间指针向非通用地址空间的指针转换必须要显式强制转换

char* ptr0;
__local__ char* ptr1 = (__local__ char*) ptr0;

非通用地址空间之间不能互相转换

__local__ char* ptr0;
// 以下代码编译器会报错
__shared__ char* ptr1 = (__shared__ char*) ptr0;

8.3. 内置向量类型和操作 (Built-in Vector Types & Operators)¶

内置向量类型 (Built-in Vector Types)¶

TOPS C++内置了向量类型，以方便开发者进行向量运算。在定义向量类型时，类型指示符和内存空间指示符位于定义中的关键字__vector（或__vector2 __vector4）之前。以下代码中提供了带__vector（或__vector2 __vector4）关键字语法的大多数合法形式。为了清晰起见，图中省略了一些变体：类型指示符（如const）和内存类指示符（如static）可以在声明中以任何顺序出现，只要两者都不紧跟在关键字__vector（或__vector2 __vector4）之后即可。

__device__ void test(void) {
    // 定义一个包含128个int的向量
    __vector4 int v1;
    // 定义一个包含128个bool的向量
    __vector4 bool int v2;

    // 定义一个包含256个short的向量
    __vector4 short v3;
    // 定义一个包含256个bool的向量
    __vector4 bool short v4;
-}

关键字	GCU210	GCU300
__vector	128 B	128 B
__vector2	256 B	256 B
__vector4	512 B	512 B

注意：当前版本的TOPS C++还不支持在GCU210上使用向量类型以及相关运算，只可以在GCU300上使用。

向量类型	向量长度
`__vector int`	包含32个int
`__vector unsigned int`	包含32个unsigned int
`__vector short`	包含64个short
`__vector unsigned short`	包含64个unsigned short
`__vector char`	包含128个char
`__vector unsigned char`	包含128个unsigned char
`__vector float`	包含32个float
`__vector __fp16`	包含64个__fp16
`__vector __bf16`	包含64个__bf16
`__vector2 int`	包含64个int
`__vector2 unsigned int`	包含64个unsigned int
`__vector2 short`	包含128个short
`__vector2 unsigned short`	包含128个unsigned short
`__vector2 char`	包含256个char
`__vector2 unsigned char`	包含256个unsigned char
`__vector2 float`	包含64个float
`__vector2 __fp16`	包含128个__fp16
`__vector2 __bf16`	包含128个__bf16
`__vector4 int`	包含128个int
`__vector4 unsigned int`	包含128个unsigned int
`__vector4 short`	包含256个short
`__vector4 unsigned short`	包含256个unsigned short
`__vector4 char`	包含512个char
`__vector4 unsigned char`	包含512个unsigned char
`__vector4 long`	包含128个long
`__vector4 unsigned long`	包含128个unsigned long
`__vector4 float`	包含128个float
`__vector4 __fp16`	包含256个__fp16
`__vector4 __bf16`	包含256个__bf16
`__vector bool int`	包含32个bool
`__vector bool unsigned int`	包含32个bool
`__vector bool short`	包含64个bool
`__vector bool unsigned short`	包含64个bool
`__vector bool char`	包含128个bool
`__vector bool unsigned char`	包含128个bool
`__vector2 bool int`	包含128个bool
`__vector2 bool unsigned int`	包含64个bool
`__vector2 bool short`	包含128个bool
`__vector2 bool unsigned short`	包含128个bool
`__vector2 bool char`	包含256个bool
`__vector2 bool unsigned char`	包含256个bool
`__vector4 bool int`	包含128个bool
`__vector4 bool unsigned int`	包含128个bool
`__vector4 bool short`	包含256个bool
`__vector4 bool unsigned short`	包含256个bool
`__vector4 bool char`	包含512个bool
`__vector4 bool unsigned char`	包含512个bool
`__vector4 bool long`	包含128个bool
`__vector4 bool unsigned long`	包含128个bool

向量运算符 (Built-in vector Operators)¶

双目运算符 (Binary Operators)¶

运算符	功能描述	支持的数据类型
`x+y`	加法运算	支持所有相同类型的标量和向量操作
`x-y`	减法运算	支持所有相同类型的标量和向量操作
`x*y`	减法运算	支持所有相同类型的标量和向量操作
`x<<y`	逻辑左移运算	支持所有无符号整型标量和向量操作
`x>>y`	逻辑右移运算	支持所有无符号整型标量和向量操作

单目运算符 (Unary Operators)¶

运算符	功能描述	支持的数据类型
`[x]`	获取向量中某个元素，或者修改向量中某个元素的值。（注：后续会添加越界检查）	支持所有向量类型
`-x`	负号运算	支持有符号数据类型向量
`++x`	自增运算	支持所有向量类型
`--x`	自减运算	支持所有向量类型
`sizeof(x)`	获取类型大小（单位：字节）	支持所有向量类型

比较运算符 (Comparison Operators)¶

对于比较运算，输入为标量时返回布尔值，输入为向量时返回对应的掩码整型向量(maskbit vector)。

运算符	功能描述	支持的数据类型
`x==y`	等于运算	支持所有标量和向量类型
`x!=y`	不等于运算	支持所有标量和向量类型
`x<y`	小于运算	支持所有标量和向量类型
`x<=y`	小于等于运算	支持所有标量和向量类型
`x>y`	大于运算	支持所有标量和向量类型
`x>=y`	大于等于运算	支持所有标量和向量类型

位运算符 (Bitwise Operators)¶

运算符	功能描述	支持的数据类型
`x&y`	求与运算	无符号整型向量、掩码整型向量
`x\|y`	求或运算	无符号整型向量、掩码整型向量
`x^y`	异或运算	无符号整型向量、掩码整型向量
`~x`	按位取反运算	无符号整型向量
`x&~y`	将y求非后再与x求与	掩码整型向量
`x\|~y`	将y求非后再与x求或	掩码整型向量
`~(x&y)`	与非运算	掩码整型向量
`~(x\|y)`	或非运算	掩码整型向量
`~(x^y)`	异或然后再求反	掩码整型向量

备注1：上表中掩码整型向量包括如下数据类型：__vector bool int, __vector bool short, __vector bool char, __vector2 bool int, __vector2 bool short, __vector2 bool char, __vector4 bool int, __vector4 bool short, __vector4 bool char。

备注2：上表中无符号整型向量包括如下数据类型：__vector unsigned char, __vector unsigned short, __vector unsigned int, __vector2 unsigned char, __vector2 unsigned short, __vector2 unsigned int, __vector4 unsigned char, __vector4 unsigned short, __vector4 unsigned int。

8.4. 数学函数 (Mathmatical Functions)¶

函数	功能描述	支持的数据类型
`abs(x)`	求绝对值运算	所有向量数据类型
`mod(x)`	求模运算	__vector4整型向量
`div(x,y)`	除法运算	__vector4向量
`sign(x)`	获取符号信息	__vector4浮点向量
`round(x)`	取整操作	__vector4浮点向量
`ln(x)`	对数函数	__vector4浮点向量、__vector浮点向量
`log1p(x)`	返回log(x+1)	__vector4浮点向量
`power(x)`	求幂函数	__vector4浮点向量
`sqrt(x)`	平方根函数	__vector4浮点向量
`rec(x)`	求倒数函数	__vector浮点向量
`rsqrt(x)`	平方根倒数函数	__vector4浮点向量、__vector浮点向量
`exp(x)`	指数函数	__vector4浮点向量、__vector浮点向量
`expm1(x)`	返回exp(x)-1	__vector4浮点向量
`isfinite(x)`	判断有限值	__vector4浮点向量
`floor(x)`	向下取整函数	__vector4浮点向量
`ceil(x)`	向上取整函数	__vector4浮点向量
`sin(x)`	正弦函数	__vector4浮点向量、__vector浮点向量
`cos(x)`	余弦函数	__vector4浮点向量
`tan(x)`	正切函数	__vector4浮点向量、__vector浮点向量
`sinh(x)`	双曲正弦函数	__vector4浮点向量
`cosh(x)`	双曲余弦函数	__vector4浮点向量
`tanh(x)`	双曲正切函数	__vector4浮点向量、__vector浮点向量
`asin(x)`	反正弦函数	__vector4浮点向量
`acos(x)`	反余弦函数	__vector4浮点向量
`atan(x)`	反正切函数	__vector4浮点向量
`atan2(x,y)`	返回以弧度表示的 x/y 的反正切	__vector4浮点向量
`softplus(x)`	softplus激活函数	__vector4浮点向量、__vector浮点向量
`sigmoid(x)`	sigmoid激活函数	__vector4浮点向量、__vector浮点向量
`gelu(x)`	gelu激活函数	__vector4浮点向量、__vector浮点向量

备注：

上表中的__vector4整型向量包括如下数据类型：__vector4 int, __vector4 short, __vector4 char, __vector4 unsigned int, __vector4 unsigned short, __vector4 unsigned char。

上表中的__vector4浮点向量包括如下数据类型：__vector4 __bf16, __vector4 __fp16, __vector4 float。

上表中的__vector4向量包括如下数据类型：__vector4 int, __vector4 short, __vector4 char, __vector4 unsigned int, __vector4 unsigned short, __vector4 unsigned char, __vector4 __bf16, __vector4 __fp16, __vector4 float。

上表中的__vector浮点向量包括如下数据类型：__vector __bf16, __vector __fp16, __vector float。

8.5. 内置变量 (Built-in Variables)¶

内建变量指定了线程网格和线程块的维度以及线程块和线程的索引。它们仅在在设备上执行的函数中有效。

gridDim¶

该变量的类型为dim3，包含了线程网格的维度。

blockIdx¶

该变量的类型为uint3，包含了线程块在线程网格中的索引。

blockDim¶

该变量的类型为dim3，包含了线程块的维度。

threadIdx¶

该变量的类型为uint3，包含了线程在线程块中的索引。

8.6. 内置宏定义 (Built-in Macro Definitions)¶

GCU_ARCH¶

GCU_ARCH是编译器内置宏，表示当前编译过程所针对的硬件架构版本，开发者可以直接在代码中使用。

宏定义	硬件架构
GCU_ARCH = 210	GCU210
GCU_ARCH = 300	GCU300

AGCU_ARCH¶

AGCU_ARCH是编译器内置宏，表示当前编译过程所针对的抽象硬件架构版本，开发者可以直接在代码中使用。

宏定义	抽象硬件架构	实际硬件架构
AGCU_ARCH = 200	AGCU200	GCU210和GCU300
AGCU_ARCH = 300	AGCU300	GCU300和下一代

8.7. 数据变换引擎 (Data Transformation Engine)¶

数据变换引擎简称DTE，是用于数据传输和转换（切片、反切片、转置等）的硬件引擎，提供满足计算引擎要求的数据。

DTE的特点包括：

独立于SIP进行工作
数据流向不需要寄存器的分配，数据流转的效率更高
节约了SIP的算力，数据传输和转换不需要SIP计算地址和初始化值

DTE上下文 (DTE Context)¶

DTE上下文用于维护DTE资源，对DTE上下文的基本操作包括：定义、初始化、释放。TOPS C++目前支持以下两种DTE上下文：

// 定义一个线程块共享的DTE
__shared__ tops_dte_ctx_t ctx;
// 定义一个线程私有的DTE
tops_dte_ctx_t ctx;

修饰符类型	可见性	占用资源	数据传输
`__shared__`	表示一个线程块内的所有线程共享的DTE上下文	使用线程块的DTE硬件资源	主要用作全局内存和共享内存之间的数据传输
无（不加任何修饰符）	表示线程私有的DTE上下文	使用线程的DTE硬件资源	主要用作全局内存和本地内存、共享内存和本地内存之间的数据传输

// 初始化DTE上下文
tops_init_dte(&ctx);

// 释放DTE上下文
tops_destroy_dte(&ctx);

DTE的数据参数 (DTE Parameters)¶

DTE 操作相关接口都使用mdspan类型作为配置数据参数。通过mdspan数据结构来给设备地址附加额外的信息（如维度、形状、所属内存空间、总大小等），其构造函数的参数包括：内存空间（可选）、内存地址、形状维度。

声明mdspan的两种方式：

// 方式 1：使用平铺展开形状维度的方式初始化形状
tops::mdspan src2((char*)(address), N, H, W, C);

// 方式 2：使用数组变量的方式初始化形状
int shape[] = {N, H, W, C};
tops::mdspan src4((char*)(address), shape);

DTE的操作 (DTE Operations)¶

DTE数据搬运操作接口一般是配置和启动合并的方式，通常至少包含两个参数：目标对象、源对象，目标对象和源对象类型统一使用mdspan数据类型。

简要起见，本小节介绍中，将目标对象简称为dst，源对象简称为src。

DTE支持两种启动方式：同步启动和异步启动。

同步启动：使用同步DTE数据搬运接口（如，memcpy()）启动并等待

tops_dte_ctx_t ctx;
ctx.init();
tops::memcpy(ctx, dst, src);
ctx.destroy();

异步启动：使用异步DTE数据搬运接口（如，memcpy_async()）启动，会返回一个 tops::event，可以使用 tops::wait 来同步等待 tops::event

tops_dte_ctx_t ctx;
tops_init_dte(ctx);
tops::event ev = tops::memcpy_async(ctx, dst, src);
tops::wait(ev);
tops_destroy_dte(ctx);

DTE数据搬运接口包括下列函数，同时也为每个函数提供了 xxx_async() 后缀的异步版本。

下表中是GCU210、GCU300都支持的数据搬运接口：

接口	描述
`tops::memcpy(ctx, dst, src)`	以src总大小拷贝src到dst，开发者需确保dst 的大小足够
`tops::memset(ctx, dst, const_value)`	将dst所指定的内存内容设置为const_value
`tops::slice(ctx, dst, src, offset)`	按照dst指定的形状和offset指定的偏置从src中拷贝数据到dst所指定的位置，开发者需确保dst的大小和偏置不超过src的相应值，否则会发生自动填充
`tops::deslice(ctx, dst, src, offset)`	把src所指定的数据拷贝并覆盖dst中由offset所指定的位置
`tops::transpose(ctx, dst, src, layout)`	按照layout对src的数据进行拷贝并转置到dst所指定的位置
`tops::slice_transpose(ctx, dst, src, offset, layout)`	slice和transpose的组合，先拷贝src中dst所指定的切片接着按照layout对切片进行转置并把结果保存在dst所指定的位置
`tops::transpose_deslice(ctx, dst, src, offset, layout)`	transpose和deslice的组合，先把src所指定的数据按照layout进行转置，接着将数据拷贝并覆盖到dst中offset所指定的位置
`tops::pad(ctx, dst, src, pad_low, pad_high, pad_mid, pad_value)`	pad操作，把src指定的数据，按照dst所指定的形状和大小，用pad_value的值设置到src的首部（pad_low有效），尾部（pad_high有效），或者中间（pad_mid有效），并把结果移动到dst所指的位置
`tops::mirror_tb(ctx, dst, src)`	按照第一维（X轴，shape的最后一个元素）将src指定的数据翻转，并把结果移动到dst所指的位置
`tops::mirror_lr(ctx, dst, src)`	按照第二维（Y轴，shape的倒数第二个元素）将src指定的数据翻转，并把结果移动到dst所指的位置
`tops::broadcast(ctx, dst, src)`	按照src到dst维度的变化，将src的数据做相应的扩张，并把结果移动到dst所指的位置

下表中是GCU300以及之后新增的数据搬运接口（GCU210不支持）：

接口	描述
`tops::slice_pad(ctx, dst, src, src_offsets, slice_shape, pad_low, pad_high, pad_mid, pad_value)`	按照slice_shape指定的形状和src_offsets指定的偏置从src中拷贝数据到dst所指定的位置，dst的大小和偏置超过 src的相应值时，发生自动填充，用pad_value的值设置到src的首部（pad_low有效），尾部（pad_high有效），或者中间（pad_mid）有效，并把结果移动到dst所指的位置
`tops::slice_deslice(ctx, dst, src, src_offsets, slice_shape, dst_offsets)`	按照slice_shape指定的形状和src_offsets指定的偏置从src中拷贝数据，并覆盖dst中由dst_offsets所指定的位置
`tops::slice_broadcast(ctx, dst, src, src_offsets, slice_shape)`	按照slice_shape指定的形状和src_offsets指定的偏置从src中拷贝数据，并将数据做dst维度相应的扩张，并把结果移动到dst所指定的位置
`tops::slice_expand(ctx, dst, src, src_offsets, slice_shape, phase, ratio)`	按照slice_shape指定的形状和src_offsets指定的偏置从src中拷贝数据，并把结果移动到dst按照ratio和phase选取数据的指定位置（举例，如果dst bpe = 2 * src bpe，那么ratio必须设置为0，phase可以取值 0或1，取值0意味着指定dst每个数据的低一半，作为拷贝目的，高一半补0）
`tops::shrink_deslice(ctx, dst, src, phase, ratio, dst_offsets)`	把src所指定的数据按照ratio和phase选取src数据（举例，如果src bpe = 2 * dst bpe，那么ratio必须设置为0，phase可以取值 0或1，取值0意味着选取src每个数据的低一半），拷贝并覆盖dst中由dst_offsets所指定的位置
`tops::memset_deslice(ctx, dst, src, dst_offsets, value)`	把src所指定的数据初始化为value值后，拷贝并覆盖dst中由offset所指定的位置
`tops::mirror_tb_pad(ctx, dst, src, pad_low, pad_high, pad_mid, value)`	按照第一维（X轴，shape的最后一个元素）将src指定的数据翻转，并把结果移动到dst所指的位置，dst的大小和偏置超过 src的相应值时，发生自动填充，用value的值设置到数据的首部（pad_low有效），尾部（pad_high有效），或者中间（pad_mid）有效，并把结果移动到dst所指的位置
`tops::mirror_lr_pad(ctx, dst, src, pad_low, pad_high, pad_mid, value)`	按照第二维（Y轴，shape的倒数第二个元素）将src指定的数据翻转，并把结果移动到dst所指的位置，dst的大小和偏置超过 src的相应值时，发生自动填充，用value的值设置到数据的首部（pad_low有效），尾部（pad_high有效），或者中间（pad_mid）有效，并把结果移动到dst所指的位置
`tops::mirror_tb_deslice(ctx, dst, src, dst_offsets)`	按照第一维（X轴，shape的最后一个元素）将src指定的数据翻转，并把结果拷贝并覆盖dst中由dst_offsets所指定的位置
`tops::mirror_lr_deslice(ctx, dst, src, dst_offsets)`	按照第二维（Y轴，shape的倒数第二个元素）将src指定的数据翻转，并把结果拷贝并覆盖dst中由dst_offsets所指定的位置
`tops::sub_sample(ctx, dst, src, dim_stride)`	采样操作，把src的dim_stride维度做下采样，并把结果拷贝到dst中

8.8. 同步机制 (Synchronization Mechanism)¶

topscc支持两种同步机制：线程组的同步和计数同步。

线程组同步¶

阻塞指定范围内所有线程，直到所有线程都到达后，解除阻塞。

接口	描述
`void __syncthreads()`	线程块内所有线程做一次同步。
`void __syncblocks()`	线程网格内所有线程做一次同步。

注意：使用__syncblocks()的核函数必须声明为__cooperative__

计数同步¶

该机制通过屏障（Barrier）对象实现，可以阻塞一组已知数量的线程，直到要求数量的线程都到达屏障。屏障对象的生命周期由一个或多个（count）阶段组成，每个阶段定义一个等待线程阻塞的阶段同步点。线程可以到达（arrive）屏障，并等待（wait）以阻塞在相位同步点上。

屏障阶段由以下步骤组成：

每次调用到达arrive()或arrive_and_wait()时，预期计数都会减少。
当预期计数达到零时，在阶段同步点上阻塞的所有线程都将解除阻塞。

接口	描述
`tte::barrier<BarrierScope>(int count)`	屏障构造函数，初始化屏障的阶段计数 enum BarrierScope { SIP = 0, CLUSTER = 1 };
`tte::barrier::arrive()`	到达屏障并减少计数
`void tte::barrier::wait()`	阻塞在同步点直到阶段计数为零
`void tte::barrier::arrive_and_wait()`	到达屏障并减少计数，阻塞在同步点直到阶段计数为零

8. C++语言扩展¶

8.1. 函数执行空间指示符 (Function Execution Space Specifiers)¶

__global__¶

__device__¶

__host__¶

__cooperative__¶

8.2. 变量内存空间指示符（Variable Memory Space Specifiers）¶

__device__¶

__constant__¶

__shared__¶

__local__¶

__restrict__¶

转换规则¶

8.3. 内置向量类型和操作 (Built-in Vector Types & Operators)¶

内置向量类型 (Built-in Vector Types)¶

向量运算符 (Built-in vector Operators)¶

双目运算符 (Binary Operators)¶

单目运算符 (Unary Operators)¶

比较运算符 (Comparison Operators)¶

位运算符 (Bitwise Operators)¶

8.4. 数学函数 (Mathmatical Functions)¶

8.5. 内置变量 (Built-in Variables)¶

gridDim¶

blockIdx¶

blockDim¶

threadIdx¶

8.6. 内置宏定义 (Built-in Macro Definitions)¶

GCU_ARCH¶

AGCU_ARCH¶

8.7. 数据变换引擎 (Data Transformation Engine)¶

DTE上下文 (DTE Context)¶

DTE的数据参数 (DTE Parameters)¶

DTE的操作 (DTE Operations)¶

8.8. 同步机制 (Synchronization Mechanism)¶

线程组同步¶

计数同步¶

global¶

device¶

host¶

cooperative¶

device¶

constant¶

shared¶

local¶

restrict¶