synchronization - Sincronización global de OpenCL y GPU

Question

¿Alguien ha probado las funciones gpu_sync descritas en el artículo "Comunicación entre bloques de la GPU a través de la sincronización rápida de barreras"? Todos los códigos descritos parecen bastante simples y fáciles de implementar, pero continúan congelando mi GPU. Estoy seguro de que estoy haciendo algo estúpido, pero no puedo ver qué. ¿Alguien puede ayudarme?

La estrategia que estoy usando es la que se describe en la sección "Sincronización de bloqueo de GPU" y aquí está el código fuente de OpenCL que implementé:

static void globalSync(uint iGoalValue,
                   volatile __global int *globalSyncFlagsIN,
                   volatile __global int *globalSyncFlagsOUT)
{
 const size_t iLocalThreadID  = get_local_id(0);
 const size_t iWorkGroupID    = get_group_id(0);
 const size_t iWorkGroupCount = get_num_groups(0);

 //Only the first thread on each SM is used for synchronization
 if (iLocalThreadID == 0)
 { globalSyncFlagsIN[iWorkGroupID] = iGoalValue; }

 if (iWorkGroupID == 0)
 {
  if (iLocalThreadID < iWorkGroupCount)
  {
   while (globalSyncFlagsIN[iLocalThreadID] != iGoalValue) {
    // Nothing to do here
   }
  }

  barrier(CLK_GLOBAL_MEM_FENCE);

  if (iLocalThreadID < iWorkGroupCount)
  { globalSyncFlagsOUT[iLocalThreadID] = iGoalValue; }
 }

 if (iLocalThreadID == 0)
 {
  while (globalSyncFlagsOUT[iWorkGroupID] != iGoalValue) {
   // Nothing to do here 
  }
 }

 barrier(CLK_GLOBAL_MEM_FENCE);
} 

Gracias por adelantado.

Answer 1

No intenté ejecutar el código, pero la traducción directa de CUDA a OpenCL del código del artículo mencionado anteriormente sería:

{  
    int tid_in_blk = get_local_id(0) * get_local_size(1)
        + get_local_id(1);
    int nBlockNum = get_num_groups(0) * get_num_groups(1);
    int bid = get_group_id(0) * get_num_groups(1) + get_group_id(1);


    if (tid_in_blk == 0) {
        Arrayin[bid] = goalVal;
    }

    if (bid == 1) {
        if (tid_in_blk < nBlockNum) {
            while (Arrayin[tid_in_blk] != goalVal){

            }
        }
        barrier(CLK_LOCAL_MEM_FENCE);

        if (tid_in_blk < nBlockNum) {
            Arrayout[tid_in_blk] = goalVal;
        }
    }

    if (tid_in_blk == 0) {
        while (Arrayout[bid] != goalVal) {

        }
    }
}

Tenga en cuenta la diferencia en los ID de subprocesos y grupos y en el uso de la barrera de memoria local en lugar de la global.