cuda
Comunicación interbloque

Observaciones

Guardia de ultima cuadra

Considere una cuadrícula trabajando en alguna tarea, por ejemplo, una reducción paralela. Inicialmente, cada bloque puede hacer su trabajo de manera independiente, produciendo un resultado parcial. Sin embargo, al final, los resultados parciales deben combinarse y fusionarse. Un ejemplo típico es un algoritmo de reducción en un big data.

Un enfoque típico es invocar dos núcleos, uno para el cálculo parcial y el otro para la fusión. Sin embargo, si la fusión se puede realizar de manera eficiente mediante un solo bloque, solo se requiere una llamada al kernel. Esto se logra mediante una protección de lastBlock definida como:

__device__ bool lastBlock(int* counter) {
  __threadfence(); //ensure that partial result is visible by all blocks
  int last = 0;
  if (threadIdx.x == 0)
    last = atomicAdd(counter, 1);
  return __syncthreads_or(last == gridDim.x-1);
}

__device__ bool lastBlock(int* counter) {
  __shared__ int last;
  __threadfence(); //ensure that partial result is visible by all blocks
  if (threadIdx.x == 0) {
    last = atomicAdd(counter, 1);
  }
  __syncthreads();
  return last == gridDim.x-1;
}

Con tal protección, se garantiza que el último bloque vea todos los resultados producidos por todos los demás bloques y puede realizar la fusión.

__device__ void computePartial(T* out) { ... }
__device__ void merge(T* partialResults, T* out) { ... }

__global__ void kernel(int* counter, T* partialResults, T* finalResult) {
    computePartial(&partialResults[blockIdx.x]);
    if (lastBlock(counter)) {
      //this is executed by all threads of the last block only
      merge(partialResults,finalResult);
    }
}

Suposiciones

• El contador debe ser un puntero de memoria global, inicializado a 0 antes de invocar el kernel.
• La función lastBlock es invocada uniformemente por todos los hilos en todos los bloques
• El kernel se invoca en una cuadrícula unidimensional (para simplificar el ejemplo)
• T nombra cualquier tipo que te guste, pero el ejemplo no pretende ser una plantilla en el sentido de C ++

Cola de trabajo global

Considere una serie de elementos de trabajo. El tiempo necesario para completar cada elemento de trabajo varía mucho. Para equilibrar la distribución de trabajo entre bloques, puede ser prudente que cada bloque obtenga el siguiente elemento solo cuando el anterior esté completo. Esto contrasta con la asignación a priori de elementos a bloques.

class WorkQueue {
private:
  WorkItem* gItems;
  size_t totalSize;
  size_t current;
public:
  __device__ WorkItem& fetch() {
    __shared__ WorkItem item;
    if (threadIdx.x == 0) {
      size_t itemIdx = atomicAdd(current,1);
      if (itemIdx<totalSize)
        item = gItems[itemIdx];
      else
        item = WorkItem::none();
    }
    __syncthreads();
    return item; //returning reference to smem - ok
  }
}

Suposiciones

• El objeto WorkQueue, así como la matriz gItem, residen en la memoria global
• No se agregan nuevos elementos de trabajo al objeto WorkQueue en el kernel que se está recuperando
• WorkItem es una pequeña representación de la asignación de trabajo, por ejemplo, un puntero a otro objeto
• WorkItem::none() función miembro estática WorkItem::none() crea un objeto WorkItem que no representa ningún trabajo
• WorkQueue::fetch() debe llamarse de manera uniforme por todos los subprocesos del bloque
• No hay 2 invocaciones de WorkQueue::fetch() sin otras __syncthreads() entre ellas. De lo contrario aparecerá una condición de carrera!

El ejemplo no incluye cómo inicializar el WorkQueue o rellenarlo. Lo hace otro kernel o código de CPU y debería ser bastante sencillo.