¿La programación warp de CUDA es determinista?

gpu-warp

Question

Me pregunto si el orden de programación warp de una aplicación CUDA es determinista.

Específicamente, me pregunto si el orden de la ejecución de warp seguirá siendo el mismo con múltiples ejecuciones del mismo kernel con los mismos datos de entrada en el mismo dispositivo. De lo contrario, ¿hay algo que pueda forzar la ordenación de la ejecución de warp (por ejemplo, en el caso de depurar un algoritmo dependiente de la orden)?

Answer 1

El comportamiento preciso de la programación warp de CUDA no está definido. Por lo tanto, no puedes depender de que sea determinista. En particular, si múltiples urdimbres están listas para ser ejecutadas en un intervalo de emisión dado, no hay una descripción de qué urdimbre será seleccionado por el planificador de warp.

No existe un método externo para controlar con precisión el orden de ejecución de warp.

Ciertamente es posible construir un código que determine la ID de urdimbre y fuerce las deformaciones para que se ejecuten en un orden particular. Algo como esto:

#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;
}

permitir solo un warp para ejecutar a la vez será muy ineficiente, por supuesto.

En general, los mejores algoritmos paralelizables tienen poca o ninguna dependencia de orden.