CUDA：性能更好，占用率更低答案

【问题标题】：CUDA: Better performances with lower occupancyCUDA：性能更好，占用率更低
【发布时间】：2022-01-21 06:26:27
【问题描述】：

我是一名学习 CUDA 的学生，我正在尝试编写一个用于计数排序的 CUDA 算法：

__global__ void kernelCountingSort(int *array, int dim_array, int *counts) {
    // define index
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int count = 0;
    // check that the thread is not out of the vector boundary
    if (i >= dim_array) return;

    for (int j = 0; j < dim_array; ++j) {
        if (array[j] < array[i])
            count++;
        else if (array[j] == array[i] && j < i)
            count++;
    }
    counts[count] = array[i];

我试图通过增加块大小来分析我的算法性能，这是具有相应块大小的时间图：

使用 64 作为块大小时，我有 100% 的占用率，但是我获得了最佳性能，因此最小执行时间，块大小为 32。我在问是否有可能以更少的占用率获得更好的表演。

我将 colab 与 Tesla T4 搭配使用，规格如下：

【问题讨论】：

我正在使用一个包含 5000 个元素的整数数组来测试算法。
This 可能感兴趣。
请注意，该算法效率不高，因为它具有二次复杂度（即O(n*n)）。实际上，它遍历数组的所有项只是为了找到在输出数组中插入该项的位置。在 GPU 上工作的排序算法要快得多，例如 Bitonic 排序 (O(n (log n)^2))。最先进的算法通常使用基数排序（O(n)——这是 CUB 高度优化的 CUDA 库所做的）。
请注意，访问 counts[count] = array[i] 具有竞争条件。两个线程可以得出相同的计数，对吗？不应该是counts[array[i]] = count吗？
我检查过，两个线程不可能通过我在 if 条件中创建的控件得出相同的计数。

标签： performance optimization cuda gpu nvidia

【解决方案1】：

我在问是否有可能以更少的占用率获得更好的性能。

是的，这是可能的，并且广受好评的 papers 已经写了关于该主题的文章。

无法使用an incomplete snippet of code 解释这在您的特定情况下是否有意义，并且没有关于 GPU 或执行环境的信息。

【讨论】：