每个块的 CUDA 不同线程用于不同的功能答案

【问题标题】：CUDA different threads per block for different functions每个块的 CUDA 不同线程用于不同的功能
【发布时间】：2023-03-20 21:17:01
【问题描述】：

我正在编写一个 CUDA 程序，但遇到了一个问题。我有两个功能：

__global__ void cal_freq_pl(float *, char *, char *, int *, int *)
__global__ void cal_sum_vfreq_pl(float *, float *, char *, char *, int *)

我这样调用第一个函数： cal_freq_pl<<<M,512>>>( ... ); M是15左右的数字，所以我不担心。 512 是我的 GPU 上每个块的最大线程数。这可以正常工作，并为所有 M*512 值提供预期的输出。

但是当我以类似的方式调用第二个函数时： cal_sum_vfreq_pl<<<M,512>>>( ... ); 这是行不通的。在调试了那个函数的废话之后，我终于发现它以这些尺寸运行：cal_sum_vfreq_pl<<<M,384>>>( ... );，它比 512 小 128。它显示 512 没有错误，但结果不正确。

我目前只能访问 Compute1.0 arch，并且在 Windows 64 位机器上拥有 Nvidia Quadro FX4600 显卡。

我不知道为什么会发生这种行为，我确信第一个函数运行 512 个线程，而第二个函数只运行 384 个（或更少）。

有人可以提出一些可能的解决方案吗？

提前谢谢...

编辑：这是内核代码：

__global__ void cal_sum_vfreq_pl(float *freq, float *v_freq_vectors, char *wstrings, char *vstrings, int *k){
    int index = threadIdx.x;
    int m = blockIdx.x;
    int block_dim = blockDim.x;
    int kv = *k; int vv = kv-1; int wv = kv-2;
    int woffset = index*wv;
    int no_vstrings = pow_pl(4, vv);
    float temppp=0;
    char wI[20], Iw[20]; int Iwi, wIi;
    for(int i=0;i<wv;i++) Iw[i+1] = wI[i] = wstrings[woffset + i];
    for(int l=0;l<4;l++){
            Iw[0] = get_nucleotide_pl(l);
            wI[vv-1] = get_nucleotide_pl(l);
            Iwi = binary_search_pl(vstrings, Iw, vv);
            wIi = binary_search_pl(vstrings, wI, vv);
            temppp = temppp + v_freq_vectors[m*no_vstrings + Iwi] + v_freq_vectors[m*no_vstrings + wIi];
    }
    freq[index + m*block_dim] = 0.5*temppp;
}

【问题讨论】：

“它不起作用”是什么意思？是内核启动失败还是内核启动并在内部崩溃？也许只是内核内部的绑定检查失败。
@hubs：从 385 开始输出错误，这很奇怪。
如果它以 385 个线程开始并且结果将是错误的，那么您的内核源代码肯定会出现故障。但是不看你的内核代码就很难说。
@hubs：不，我的意思是，如果我输入 blockDim = 385 输出是错误的，并且整个输出都是错误的，不仅仅是从 385 开始，而是所有 M*512 值。我将添加内核代码...
作为一般规则，您应该在任何情况下检查线程是否要在您分配的内存中访问。尝试使用 cuda-memcheck 运行您的程序

标签： cuda

【解决方案1】：

看来您在第二个内核中分配了很多寄存器。由于硬件资源限制（例如每个块的寄存器数量），您不能总是达到每个块的最大线程数。

CUDA 提供了一个工具来帮助计算每个块的正确线程数。

http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_calculator.xls

您还可以在 CUDA 安装目录中找到此 .xls 文件。

【讨论】：

哦，我会调查的。一个简单的问题，char[20] 算作 1 个寄存器还是 20 个寄存器？
看来binary_search_pl( ... ); 是所有问题的根源。我不知何故需要用一些东西来替换它，这样线程/块就不会下降到 384。我仍然不确定寄存器变量在哪里使用，但它与这个 __device__ 函数有关。由于上述评论中的答案表明长字符数组会自动传输到本地内存，我也对其进行了测试，char[20] 不是问题，但__device__ 函数调用是（在这种情况下为binary_search_pl），它们使用了太多注册变量。
@user1961040，有编译选项可以为您输出寄存器使用情况。或者，我经常使用 Visual profile 来检查寄存器的使用情况。
我刚刚发现了一个新事物。因此，由于内核代码（和__device__ 函数调用），寄存器的数量减少了。但是我得到的错误输出不是由于内核代码的一些错误计算。代码甚至没有编译，但 VS2008 没有显示错误，即使使用 --ptxas-options=-v 选项。所以运行的代码是最后一个成功编译的版本。我所做的任何更改（线程数 > 384）都不会编译，因此会运行以前的版本。我经常遇到这个问题，所以想在这里为面临类似问题的人写它。