2

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

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

经线大小为 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 时我将获得最大效率。那是对的吗？

CUDA 占用计算 谢谢

4

4 回答 4

3

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

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

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

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

于 2012-05-02T12:17:26.383 回答
2

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

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

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

于 2012-05-02T17:50:35.143 回答
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）组合成一个索引。

于 2012-05-02T12:35:44.017 回答
1

您只在内核中使用网格的 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 的变化。

于 2012-05-03T00:37:44.873 回答