Как предотвратить тупик в семафоре OpenCL?

Question

Я следую этому уроку, который великолепен на всем протяжении, кроме самого последнего примера о том, как создать семафор, у меня не работает. Логика довольно проста, но я не могу понять, почему это ядро ​​приводит к бесконечному циклу.

myKernel.cl

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
void GetSemaphor(__global int * semaphor, __global int * data) {
   int occupied = atom_xchg(semaphor, 1);
   int realityCheck = 0;
   while(occupied == 1 && realityCheck++ < 100000)
        occupied = atom_xchg(semaphor, 1);
}

void ReleaseSemaphor(__global int * semaphor)
{
   int prevVal = atom_xchg(semaphor, 0);
}

__kernel void myKernel(__global int* data, __global int* semaphor)
{
    // semaphor[0] is set to 0 on the host.
    GetSemaphor(&semaphor[0], data);
    data[0]++;
    ReleaseSemaphor(&semaphor[0]);
}

Это с:

OpenCL 1.2

FULL_PROFILE

на Quadro NVS 290, который имеет

*cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics

Answer 1

Этот учебник, который вы описываете, неверен и никогда не будет работать на устройствах с графическим процессором. Из-за архитектуры HW.

Любой вид механизма синхронизации, который блокирует рабочий элемент внутри рабочей группы, просто не будет работать. Поскольку состояние блокировки будет влиять на всю рабочую группу, образуется бесконечный цикл.

Вы можете делать такие вещи только с размером рабочей группы 1. Или для разных рабочих групп.