5. 线程模型¶

5.1. 线程层次结构 (Thread Hierarchy)¶

在TOPS编程模型中，计算资源的最小颗粒是一个线程(Thread)，每个线程有自己独立的寄存器和堆栈。同时，每个线程可以属于不同层次的多个线程组(Thread Group)，每个线程组中的线程可以互相之间通过共享内存共享数据，也可以互相之间进行同步。编译器内部内置了两个线程组：线程块(Thread Block)和线程网格(Thread Grid)。其中，线程网格表示启动一次核函数产生的所有线程，而线程块则只包含了其中部分线程。

如下图所示，在一个线程网格中包含了4个线程块，每个线程块中包含了8个线程(以箭头表示)，总计32个线程。

每次启动核函数产生的线程，都被组织成线程网格、线程块、线程三层结构。有两种不同的启动核函数的方式，分别提供不同的共享和同步范围。

启动方式	线程块内共享或同步	线程网格内共享或同步
topsLaunchKernel	支持	不支持
topsCooperativeLaunchKernel	支持	支持

如下表所示，不同的启动方式还有不同的线程数量限制：

启动方式	线程块数量限制	线程数量限制
topsLaunchKernel	<=1024	<=计算单元集群中的计算单元数量
topsCooperativeLaunchKernel	<=计算单元集群数量	<=计算单元集群中的计算单元数量

每个线程拥有自己的线程编号(thread index)，线程编号提供 x y z 三个成员变量。在使用时，既可以只使用x维度表示一维数据，也可以使用 x y 维度表示二维数据，或者使用 x y z 维度表示三维数据。线程编号是根据线程组的分层进行编号的，下表中定义了两层、四个内置变量供开发者使用：

变量名称	含义
gridDim	一个线程网格中包含的线程块数量
blockIdx	当前线程块在线程网格中的编号
blockDim	一个线程块中包含的线程数量
threadIdx	当前线程在线程块中的编号

上述四个内置变量都包含 .x .y .z 三个成员变量，分别表示在 x y z 三个维度上的数值。如果只使用一个维度，则其他维度为一，一个线程网格中包含(gridDim.x * blockDim.x)个线程；如果使用两个维度，则一个线程网格中包含(gridDim.x * gridDim.y * blockDim.x * blockDim.y)个线程；如果使用三个维度，则一个线程网格中包含(gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z)个线程。

每个线程块中的线程数量，以及每个线程网格中的线程块数量，会根据硬件版本的不同和核函数启动方式的不同而有所变化。具体数量限制可以参考《AGCU计算能力》章节。

5.2. 软硬件映射 (Software to Hardware Mapping)¶

线程模型是针对具体硬件资源的抽象，规定了多个线程之间可以在哪些范围内同步和共享资源。下表中定义了不同层次的线程模型概念是如何映射到GCU芯片硬件资源上的。

线程模型概念	GCU硬件资源
线程	计算单元
线程块	计算单元集群
线程网格	芯片
本地内存	一级内存
共享内存	二级内存
全局内存	设备内存或主机端内存

本地内存(Local Memory)属于单个线程所有，无法在多个线程间共享数据，其访问延迟较低。
共享内存(Shared Memory)由单个线程块中的所有线程共享，可以在这些线程之间共享数据，其访问延迟中等。
全局内存(Global Memory)由所有线程共享，包括不同核函数或不同线程网格所产生的线程之间都可以共享，其访问延迟最大。