cuda - CUDA 扭曲调度是确定性的吗？

Question

我想知道 CUDA 应用程序的扭曲调度顺序是否是确定性的。

具体来说，我想知道在同一设备上使用相同输入数据多次运行同一内核时，warp 执行的顺序是否会保持不变。如果没有，是否有任何东西可以强制执行扭曲执行的顺序（比如在调试依赖于顺序的算法时）？

Answer 1

CUDA warp 调度的精确行为没有定义。因此，您不能依赖它是确定性的。特别是，如果多个warp 准备好在给定的问题槽中执行，则没有关于warp 调度程序将选择哪个warp 的描述。

没有外部方法可以精确控制扭曲执行的顺序。

当然可以构建确定 warp ID 并强制 warp 以特定顺序执行的代码。像这样的东西：

#include <stdio.h>

#define N_WARPS 16
#define nTPB (32*N_WARPS)

__device__ volatile int my_next = 0;
__device__ int warp_order[N_WARPS];

__global__ void my_kernel(){

  __shared__ volatile int warp_num;
  unsigned my_warpid = (threadIdx.x & 0x0FE0U)>>5;
  if (!threadIdx.x) warp_num = 0;
  __syncthreads();  // don't use syncthreads() after this point
  while (warp_num != my_warpid);
  // warp specific code here
  if ((threadIdx.x & 0x01F) == 0){
    warp_order[my_next++] = my_warpid;
    __threadfence();
    warp_num++; // release next warp
    } // could use syncthreads() after this point, if more code follows
}


int main(){

  int h_warp_order[N_WARPS];
  for (int i = 0; i < N_WARPS; i++) h_warp_order[i] = -1;
  cudaMemcpyToSymbol(warp_order, h_warp_order, N_WARPS*sizeof(int));
  my_kernel<<<1,nTPB>>>();
  cudaDeviceSynchronize();
  cudaMemcpyFromSymbol(h_warp_order, warp_order, N_WARPS*sizeof(int));
  for (int i = 0; i < N_WARPS; i++) printf("index: %d, warp_id: %d\n", i, h_warp_order[i]);
  return 0;
}

当然，一次只允许执行一个 warp 是非常低效的。

一般来说，最好的可并行化算法很少或没有顺序依赖性。