如何计算块号

Question

我正在编写一个 CUDA 代码，我正在使用 GForce 9500 GT 显卡。

我正在尝试处理 20000000 个整数元素的数组，我使用的线程号是 256

warp 大小为 32。计算能力为 1.1

这是硬件http://www.geforce.com/hardware/desktop-gpus/geforce-9500-gt/specifications

现在，区块编号 = 20000000/256 = 78125 ?

这听起来不正确。如何计算块数？ 任何帮助将不胜感激。

我的 CUDA 核函数如下。这个想法是每个块将计算其总和，然后通过将每个块的总和相加来计算最终总和。

__global__ static void calculateSum(int * num, int * result, int DATA_SIZE)
{
    extern __shared__ int shared[];
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    shared[tid] = 0;
    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        shared[tid] += num[i];
    }

    __syncthreads();
    int offset = THREAD_NUM / 2;
    while (offset > 0) {
        if (tid < offset) {
            shared[tid] += shared[tid + offset];
        }
        offset >>= 1;
        __syncthreads();
    }

    if (tid == 0) {
        result[bid] = shared[0];

    }
}

我把这个函数称为

calculateSum <<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>> (gpuarray, result, size);

其中 THREAD_NUM = 256 gpu 数组的大小为 20000000。

这里我只是使用块号为 16 但不确定它是否正确？ 如何确保实现最大并行度？

这是我的 CUDA 占用计算器的输出。它说当块数为 8 时我将有 100% 的占用率。这意味着当块数 = 8 和线程数 = 256 时我将获得最大效率。对吗？

谢谢

Answer 1

如果每个线程处理一个元素，并且每个块有 256 个线程，您应该运行 20000000 个线程，结果正好是 78125 个块。这是一个完全有效的数字。

但是，有一个小问题。我手头没有 CC1.1 设备，但在 CC1.3 中：

Maximum sizes of each dimension of a grid:     65535 x 65535 x 1

因此，您应该为不同部分的数据运行多次内核，或者制作 2D 网格，然后将线程的 2D 地址简单地转换为数组元素的 1D 地址。

Answer 2

在您的情况下，线程总数 (20000000) 除以每个块的线程数 (256)，因此您可以使用该数字 (78125)。如果数字不均分，则常规整数除法会将其四舍五入，最终得到的线程数比需要的少。因此，在这种情况下，您需要使用如下函数将除法的结果四舍五入：

int DivUp(int a, int b) {
  return ((a % b) != 0) ? (a / b + 1) : (a / b);
}

由于此函数可能会为您提供比元素更多的线程，因此您还需要在内核中添加一个测试以中止对最后几个线程的计算：

int i(blockIdx.x * blockDim.x + threadIdx.x);
if (i >= n_items) {
  return;
}

但是，还有一个额外的障碍。您的硬件在一个网格中的每个维度中限制为最多 65535 个块，并且限制为两个维度（x 和 y）。因此，如果在使用 DivUp() 之后，您最终得到的计数高于该计数，那么您有两个选择。您可以拆分工作负载并多次运行内核，也可以使用两个维度。

要使用两个维度，您可以选择两个数字，每个数字都低于硬件限制，并且在相乘时成为您需要的实际块数。然后在内核顶部添加代码，将两个维度（x 和 y）组合成一个索引。

Answer 3

您发布的内核代码可以处理任何输入数据大小，与您选择启动的块数无关。选择应该简单地取决于性能。

根据经验，对于这种内核，您需要在单个多处理器上同时运行的块数乘以卡上多处理器的数量。第一个数字可以使用 CUDA 工具包中附带的 CUDA 占用电子表格获得，但上限为每个多处理器 8 个块，第二个数字将是 4对于您拥有的设备。这意味着不需要超过 32 个块来实现最大可能的并行度，但要准确回答需要访问编译器，而我目前没有。

您还可以使用基准测试来通过实验确定最佳块数，使用 4、8、12、16、20、24、28 或 32 个块之一（4 的倍数，因为这是您卡上多处理器的数量)。

Answer 4

您只在内核中使用网格的 x 维。因此，使用 cc 1.1 时，您只能使用 65535 个块。

20000000/256 = 78125 是正确的！

所以你肯定需要超过 1 个块。

内核：

//get unique block index
const unsigned int blockId = blockIdx.x //1D
    + blockIdx.y * gridDim.x //2D

//terminate unnecessary blocks
if(blockId >= 78124)
    return;

//... rest of kernel

最简单的方法是使用两个 y 块并在内核中检查块 ID。

dim3 gridDim = dim3(65535, 2); 

这将使超过 52945 个块无用，我不知道开销是多少，但先填充 x 然后填充 y 和 z 维度会创建很多未使用的块，尤其是在达到 z 维度时！

（Nvidia 应该明确地提供了一个实用函数，可以为内核内部的独特块使用获得最佳网格使用，就像这里的情况一样）

对于这个简单的例子，如何使用 x 和 y 并计算根。

grid(280, 280) = 78400 blocks //only 275 blocks overhead, less is not possible

这是计算能力 3.0 的一大优势。每个块上的 32 位范围通常使生活更轻松。 为什么它被限制为 65535 我一直不明白。

但我还是更喜欢向下兼容。

我还会测试@talonmies 的变化。