Глобальная синхронизация OpenCL и GPU

Question

Кто-нибудь пробовал функции gpu_sync, описанные в статье «Межблочная передача графического процессора через Fast Barrier Synchronization»? Все описанные коды кажутся довольно простыми и легкими в реализации, но он продолжает замерзать мой графический процессор. Я уверен, что я делаю что-то глупое, но я не понимаю, что. Может кто-нибудь мне помочь?

Стратегии Я использую это один описан в разделе «GPU безблокировочной синхронизации» и вот исходный код OpenCL я реализовал:

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

Спасибо заранее.

Answer 1

Я не пытался запустить код, но прямой перевод с CUDA на OpenCL коды из упомянутой выше статьи будет:

{ 
    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) { 

     } 
    } 
} 

Пожалуйста, обратите внимание на разницу в нити и идентификаторы групп, так и в используя локальный барьер памяти вместо глобального.