CUDA, mutex и atomicCAS()

Question

Недавно я начал разрабатывать CUDA и столкнулся с проблемой atomicCAS(). Чтобы сделать некоторые манипуляции с памятью в коде устройства, мне нужно создать мьютекс, чтобы только один поток мог работать с памятью в критической части кода.

Код устройства ниже работает на 1 блок и несколько потоков.

__global__ void cudaKernelGenerateRandomGraph(..., int* mutex)
{
    int i = threadIdx.x;
    ...

    do 
    {
        atomicCAS(mutex, 0, 1 + i);
    }
    while (*mutex != i + 1);

    //critical section
    //do some manipulations with objects in device memory

    *mutex = 0;

    ...
}

Когда первый поток выполняется

atomicCAS(mutex, 0, 1 + i);

мьютекс равен 1. После этого первый поток меняет свой статус с активного на неактивный, и строка

*mutex = 0;

не выполняется. Другие темы всегда остаются в цикле. Я пробовал много вариантов этого цикла, таких как while(){};, do{}while();, с временной переменной = *mutex внутри цикла, даже вариант с if(){} и goto. Но результат тот же.

Основная часть кода:

...
int verticlesCount = 5;
int *mutex;
cudaMalloc((void **)&mutex, sizeof(int));
cudaMemset(mutex, 0, sizeof(int));
cudaKernelGenerateRandomGraph<<<1, verticlesCount>>>(..., mutex);

Я использую Visual Studio 2012 с CUDA 5.5.

Устройство NVidia GeForce GT 240 с вычислительной возможностью 1.2.

Заранее спасибо.

---

UPD: Спустя некоторое время, работая над моим дипломным проектом этой весной, я нашел решение для критического раздела о cuda. Это сочетание механизмов без блокировки и мьютекса. Вот рабочий код. Использовал его для создания атомного динамически изменяемого размера массива.

// *mutex should be 0 before calling this function
__global__ void kernelFunction(..., unsigned long long* mutex) 
{
    bool isSet = false; 
    do 
    {
        if (isSet = atomicCAS(mutex, 0, 1) == 0) 
        {
            // critical section goes here
        }
        if (isSet) 
        {
            mutex = 0;
        }
    } 
    while (!isSet);
}

Answer 1

Цикл, о котором идет речь

do 
{
    atomicCAS(mutex, 0, 1 + i);
}
while (*mutex != i + 1);

будет работать нормально, если он работает на стороне хоста (ЦП); однажды поток 0 устанавливает *mutex до 1, другие потоки будут ждать точно, пока поток 0 устанавливает *mutex вернуться к 0.

Однако потоки GPU не так независимы, как их аналоги CPU. Потоки GPU сгруппированы в группы по 32, обычно называемые перекосами. Потоки в одной и той же деформации будут выполнять инструкции в полной блокировке. Если контрольный оператор, такой как if или же while приводит к тому, что некоторые из 32 потоков расходятся с остальными, остальные потоки будут ждать (то есть спит), пока расходящиеся потоки не завершатся. [1]

Возвращаясь к рассматриваемому циклу, поток 0 становится неактивным, потому что потоки 1, 2, ..., 31 все еще застряли в while петля. Таким образом, поток 0 никогда не достигает линии *mutex = 0и остальные 31 темы зацикливаются навсегда.

Потенциальное решение состоит в том, чтобы сделать локальную копию рассматриваемого общего ресурса, позволить 32 потокам изменить копию, а затем выбрать один поток, чтобы "перенести" изменение обратно в общий ресурс. __shared__ переменная идеальна в этой ситуации: она будет совместно использоваться потоками, принадлежащими тому же блоку, но не другими блоками. Мы можем использовать __syncthreads() для точного контроля доступа к этой переменной потоками-членами.

[1] Руководство по передовой практике CUDA - ветвление и расхождение

Избегайте разных путей выполнения в пределах одной и той же основы.

Любая инструкция управления потоком (если, switch, do, for, while) может существенно повлиять на пропускную способность команды, вызвав расхождение потоков одной и той же деформации; то есть следовать различным путям выполнения. Если это происходит, различные пути выполнения должны быть сериализованы, поскольку все потоки деформации совместно используют счетчик программ; это увеличивает общее количество инструкций, выполненных для этой деформации. Когда все различные пути выполнения завершены, потоки возвращаются к одному и тому же пути выполнения.