c - 查找 __global__ 函数允许的最大块/线程数

Question

这似乎是一个简单的问题，但我无法在任何地方找到答案。我有一个全局函数，我可以这样调用:

func<<<nbBlocks,nbThreadByBlock, nbBytesOfSharedMmy>>>(args);

如果我理解正确的话，我永远不能对 nbThreadByBlock 使用超过 1024 个，但是我如何动态知道我的函数 func 允许的 nbThreadByBlock 最大值是多少> 对于我的 GPU ？

如果我的 func 函数使用更多局部变量，每个 block 的最大线程数就会减少，我的想法是否正确？

关于我可以使用的 block 总数，有上限吗？我在想，如果我放置的 block 多于可能的数量，它们将按顺序处理，是真的吗？

谢谢!

Answer 1

下面是一段代码

cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, 0); //assuming current device ID is 0

将设备的属性收集到deviceProp中。如您所见here ，成功调用 cudaGetDeviceProperties 后，您将能够访问具有所需设备属性的 deviceProp 成员。例如，deviceProp.maxThreadsPerMultiProcessor 表示每个多处理器的最大线程数，deviceProp.maxThreadsPerBlock 表示每个 block 的最大线程数，等等。

每个 block 的适当线程数以及调用函数的 block 总数主要取决于您的设备属性和程序。您调用的每个 block 都会占用 SM 的一部分。多少取决于您的 block 请求的资源:线程、寄存器和共享内存。
考虑这个例子。假设您的设备 SM 最多可以有 2048 个线程、48 KB 共享内存和 64 KB 寄存器。如果您的 block 需要 512 个线程，并且同时使用 SM 可用的所有共享内存和寄存器，则 SM 中不可能有另一个具有相同特征的 block 。因此，由于无法使用 2048 减去 512 个潜在 SM 线程，您可以将最大实现占用率降低到 25%。现在，如果您设计 block 的方式是将 block 中的线程增加到 1024 个，则可以消耗相同数量的寄存器和共享内存，则占用率将增加一倍，达到 50%。

通常不建议使用大量 block 。 GPU 将新 block 调度到可用的 SM。如果所有 SM 都被占用，则它将对该 block 进行排队，直到 SM 有足够的空闲资源用于该 block 。调度新 block 会给 GPU 带来开销(尽管很小)。最好在找到最佳 block 大小后，计算(或分析)该 block 在 SM 上的占用情况，然后调用与占用所有 GPU SM 一样多的 block 。如果您需要更多 block ，您可以重用已完成工作的 block 的线程。
例如转换

GPU_kernel<<<1024,512>>>();

其中

__global__ void GPU_kernel(void){

    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    //rest of code
}

进入

GPU_kernel<<<(number_of_SMs*number_of_blocks_per_SM),512>>>();

其中

__global__ void GPU_kernel(void){

    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for (; tid < 1024*512; tid +=  blockIdx.x* gridDim.x ) {
        //rest of code
    }
} 

通常会带来更好的性能。

另请注意，在上面的代码片段中，我没有包含正确的 CUDA 错误检查。请应用您自己的方法来处理可能的错误。说明here .