167

如何组织线程以由 GPU 执行？

4

2 回答 2

302

硬件

例如，如果一个 GPU 设备有 4 个多处理单元，每个单元可以运行 768 个线程：那么在给定时刻，真正并行运行的线程不会超过 4*768 个（如果您计划更多线程，它们将等待轮到他们了）。

软件

线程以块的形式组织。一个块由一个多处理单元执行。可以使用 1Dimension(x)、2Dimensions (x,y) 或 3Dim 索引 (x,y,z) 来识别（索引）块的线程，但对于我们的示例，无论如何 x y z <= 768（其他限制适用到 x、y、z，请参阅指南和您的设备功能）。

显然，如果您需要超过 4*768 个线程，则需要超过 4 个块。块也可以索引为 1D、2D 或 3D。有一个等待进入 GPU 的块队列（因为在我们的示例中，GPU 有 4 个多处理器并且只有 4 个块同时执行）。

现在一个简单的案例：处理一个 512x512 的图像

假设我们想要一个线程处理一个像素 (i,j)。

我们可以使用每个 64 个线程的块。然后我们需要 512*512/64 = 4096 个块（所以要有 512x512 个线程 = 4096*64）

通常在 blockDim = 8 x 8（每个块 64 个线程）的 2D 块中组织（使索引图像更容易）线程。我更喜欢称它为threadsPerBlock。

dim3 threadsPerBlock(8, 8);  // 64 threads

和 2D gridDim = 64 x 64 块（需要 4096 个块）。我更喜欢称它为 numBlocks。

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
              imageHeight/threadsPerBlock.y);

内核是这样启动的：

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );

最后：会有类似“4096 个块的队列”的东西，其中一个块正在等待分配 GPU 的多处理器之一以执行其 64 个线程。

在内核中，线程要处理的像素 (i,j) 是这样计算的：

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;
于 2010-03-06T11:16:49.243 回答
10

假设一个 9800GT GPU：

  • 它有 14 个多处理器 (SM)
  • 每个 SM 有 8 个线程处理器（AKA 流处理器、SP 或内核）
  • 每个块最多允许 512 个线程
  • warpsize 是 32（这意味着每个 14x8=112 线程处理器可以调度多达 32 个线程）

https://www.tutorialspoint.com/cuda/cuda_threads.htm

一个块不能有比 512 更多的活动线程，因此__syncthreads只能同步有限数量的线程。即如果您使用 600 个线程执行以下操作：

func1();
__syncthreads();
func2();
__syncthreads();

那么内核必须运行两次，执行顺序为：

  1. func1 为前 512 个线程执行
  2. func2 为前 512 个线程执行
  3. func1 为剩余线程执行
  4. func2 为剩余线程执行

笔记：

要点是__syncthreads块范围的操作，它不会同步所有线程。

我不确定__syncthreads可以同步的线程的确切数量，因为您可以创建一个具有超过 512 个线程的块并让 warp 处理调度。据我了解，更准确的说法是：至少在前 512 个线程中执行 func1。

在我编辑这个答案之前（早在 2010 年），我测量了 14x8x32 线程是使用__syncthreads.

如果有人再次对此进行测试以获得更准确的信息，我将不胜感激。

于 2010-06-14T06:25:52.623 回答