如何以编程方式确定持久内核的正确启动参数

Question

以编程方式确定持久内核的启动参数的正确方法是什么？我发现的所有示例都使用硬编码值。

以下内容正确吗？

cudaDeviceProp props;
cudaGetDeviceProperties(&props, 0);
int blockCount = props.maxBlocksPerMultiProcessor * props.multiProcessorCount;
int blockThreadCount = props.maxThreadsPerMultiProcessor / props.maxBlocksPerMultiProcessor;
//  Gives <<<1312, 96>>> on a RTX 3090
PersistentKernel<<<blockCount, blockThreadCount>>>(...);

Answer 1

以下内容正确吗？

否。

使用cudaOccupancyMaxPotentialBlockSize。这将为您提供当前设备的网格大小和块大小，从而以最小的块数最大化给定内核的占用率。这是给定持久内核的最佳启动参数。

请注意，返回的块和网格维度是标量。您可以自由地将它们重新整形为多维dim3块和/或网格维度，以保留每个块和API返回的块的线程总数。