内核启动的子内核是否与主机启动的内核具有相同的 warp 大小?
Do kernel-launched child kernels have the same warp size as host-launched kernels?
当一个内核块从主机启动时,它的 warp 大小为 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 上浪费大量处理器。在此特定示例中,您可能希望首先对数据进行排序,然后在一个内核中分派可动态并行化的索引,在另一个内核中分派形状不佳的索引。
它们确实具有相同的经纱尺寸。但那是因为图形卡的扭曲大小是固定的。同一显卡上的所有内核 运行 将具有相同的扭曲大小。
今天几乎所有 GPU 都使用 32 的扭曲大小,但将来可能会改变。
您是否考虑过内核中的线程数,而不是 warp 大小?
如果是这样,那么不,它们不一定相同。您以与从主机启动它相同的方式启动具有动态并行性的新内核:
<<<blocks, threads>>>threadsafe_rand_uniform();
请注意,这与您当前正在执行的调用设备功能不同。
关于你是否值得的问题?好吧,如果不考虑替代方案就很难判断。如果备选方案是 return 数据到主机,以便主机可以启动新的合适的内核,那么它可能是值得的。但这完全取决于上下文。
当一个内核块从主机启动时,它的 warp 大小为 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 上浪费大量处理器。在此特定示例中,您可能希望首先对数据进行排序,然后在一个内核中分派可动态并行化的索引,在另一个内核中分派形状不佳的索引。
它们确实具有相同的经纱尺寸。但那是因为图形卡的扭曲大小是固定的。同一显卡上的所有内核 运行 将具有相同的扭曲大小。
今天几乎所有 GPU 都使用 32 的扭曲大小,但将来可能会改变。
您是否考虑过内核中的线程数,而不是 warp 大小? 如果是这样,那么不,它们不一定相同。您以与从主机启动它相同的方式启动具有动态并行性的新内核:
<<<blocks, threads>>>threadsafe_rand_uniform();
请注意,这与您当前正在执行的调用设备功能不同。
关于你是否值得的问题?好吧,如果不考虑替代方案就很难判断。如果备选方案是 return 数据到主机,以便主机可以启动新的合适的内核,那么它可能是值得的。但这完全取决于上下文。