内核启动的子内核是否与主机启动的内核具有相同的扭曲大小？

Question

从主机启动内核块时，它的扭曲大小为 32。通过动态并行启动的子内核是否也一样？我的猜测是肯定的，但我没有在文档中看到它。

当然，更大的问题是：值得吗？

__global__ void kernel(const int * vec, float * outvec){
    int idx = threadIdx.x;
    float random_sum=0;
    for(int j=0; j<vec[idx]; j++){
        random_sum+=threadsafe_rand_uniform();
    }
    outvec[idx] = random_sum;
}

好的，这个例子有点做作。不过，关键是，如果线程之间有一个不同长度的循环，那么尝试动态并行化它是很诱人的。但是，如果 warp 仍然是 32，您最终会在大小不均的 warp 上浪费大量处理器。在这个特定的示例中，您可能希望首先对数据进行排序，然后将动态可并行化的索引分派到一个内核中，并将形状不佳的索引分派到另一个内核中。

Answer 1

它们确实具有相同的经线尺寸。但那是因为图形卡的扭曲尺寸是固定的。在同一显卡上运行的所有内核都将具有相同的 warp 大小。

如今几乎所有 GPU 都使用 32 的 warp 大小，但未来可能会发生变化。

您是否考虑过内核中的线程数，而不是 warp 大小？ 如果是这样，那么不，它们不一定相同。您启动具有动态并行性的新内核的方式与从主机启动它的方式相同：

<<<blocks, threads>>>threadsafe_rand_uniform();

请注意，这与仅调用设备函数不同，这是您当前正在执行的操作。

关于是否值得的问题？好吧，如果不考虑替代方案，就很难说。如果替代方案是将数据返回给主机，以便主机可以启动新的适当内核，那么它​​可能是值得的。但这一切都取决于上下文。