CUDA: __syncthreads() внутри операторов if

Question

CUDA: __syncthreads() внутри операторов if

У меня есть вопрос о синхронизации CUDA. В частности, мне нужны пояснения по синхронизации операторов if. Я имею в виду, что если я помещу __syncthreads() в область действия оператора if, затронутого частью потоков внутри блока, что произойдет? Я думал, что некоторые потоки останутся "навсегда" в ожидании других потоков, которые не достигнут точки синхронизации. Итак, я написал и выполнил пример кода для проверки:

__global__ void kernel(float* vett, int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();

        index += gridSize;
    }
}

Удивительно, но я заметил, что вывод был довольно "нормальным" (64 элемента, размер блока 32):

100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2

Поэтому я немного изменил свой код следующим образом:

__global__ void kernel(float* vett, int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();
            vett[index] = 3;
        __syncthreads();

        index += gridSize;
    }
}

И вывод был:

3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3

Опять же, я был неправ: я думал, что потоки внутри оператора if после изменения элемента вектора останутся в состоянии ожидания и никогда не выйдут из области действия if. Итак... не могли бы вы уточнить, что случилось? Разблокирует ли поток, получаемый после точки синхронизации, потоки, ожидающие на барьере? Если вам нужно воспроизвести мою ситуацию, я использовал CUDA Toolkit 5.0 RC с SDK 4.2. Заранее большое спасибо.

20

c parallel-processing cuda synchronizing

Источник

user1001693 20 сен '12 в 19:27

4 ответа

Другие вопросы по тегам c parallel-processing cuda synchronizing

user1662425 20 сен '12 в 20:28 2012-09-20 20:28 · Answer 1 · 2012-09-20 20:28

Короче говоря, поведение не определено. Так что иногда он может делать то, что вы хотите, а может и нет, или (вполне вероятно) просто зависнет или сломает ваше ядро.

Если вам действительно любопытно, как все работает внутри, вы должны помнить, что потоки не выполняются независимо, а деформируются (группа из 32 потоков) одновременно.

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

Это отдельное выполнение путей продолжается до некоторой точки, для которой компилятор может определить, что он гарантированно будет достигнут всеми потоками двух отдельных путей выполнения ("точка сходимости" или "точка синхронизации"). Когда выполнение первого пути кода достигает этой точки, оно останавливается, и вместо этого выполняется второй путь кода. Когда второй путь достигает точки синхронизации, все потоки снова включаются, и выполнение оттуда продолжается равномерно.

Ситуация усложняется, если перед синхронизацией встречается другая условная ветвь. Эта проблема решается с помощью стека путей, которые все еще должны быть выполнены (к счастью, рост стека ограничен, поскольку мы можем иметь не более 32 различных путей кода для одной деформации).

Место вставки точек синхронизации не определено и даже немного различается в зависимости от архитектуры, поэтому опять же нет никаких гарантий. Единственный (неофициальный) комментарий, который вы получите от Nvidia, заключается в том, что компилятор довольно хорошо находит оптимальные точки синхронизации. Однако часто возникают тонкие проблемы, которые могут сдвинуть оптимальную точку дальше, чем вы могли бы ожидать, особенно если потоки выходят рано.

Теперь, чтобы понять поведение директивы __syncthreads(), которая переводит в bar.sync инструкции в PTX) важно понимать, что эта инструкция выполняется не для каждого потока, а для всей деформации сразу (независимо от того, отключены ли какие-либо потоки), потому что синхронизировать нужно только деформации блока. Потоки деформации уже выполняются синхронно, и дальнейшая синхронизация либо не будет иметь эффекта (если все потоки включены), либо приведет к тупику при попытке синхронизировать потоки из разных путей условного кода.

Вы можете перейти от этого описания к поведению вашего конкретного фрагмента кода. Но имейте в виду, что все это не определено, никаких гарантий нет, и опора на конкретное поведение может повредить ваш код в любое время.

Возможно, вы захотите взглянуть на руководство по PTX для более подробной информации, особенно для bar.sync инструкция, что __syncthreads() компилируется в. Документ Генри Вонга "Демистификация микроархитектуры графических процессоров с помощью микробенчмаркинга", на который ссылается Ахмад, также стоит прочитать. Хотя на данный момент устаревшая архитектура и версия CUDA, разделы об условном ветвлении и __syncthreads() кажется, все еще в целом действителен.

user1501388 20 сен '12 в 20:27 2012-09-20 20:27 · Answer 2 · 2012-09-20 20:27

Модель CUDA - это MIMD, но в современных графических процессорах NVIDIA реализована __syncthreads() при зернистости основы вместо резьбы. Это означает, что это warps inside a thread-block которые синхронизируются не обязательно threads inside a thread-block, __syncthreds() ждет, пока все "перекосы" потокового блока не достигнут барьера или не выйдут из программы. Обратитесь к документу "Разоблачение" Генри Вонга для получения более подробной информации.

user596781 20 сен '12 в 20:07 2012-09-20 20:07 · Answer 3 · 2012-09-20 20:07

Вы не должны использовать __syncthreads() если оператор не достигнут во всех потоках в одном блоке потоков, всегда. Из руководства по программированию (B.6):

__syncthreads() допускается в условном коде, но только в том случае, если условное вычисление выполняется одинаково по всему блоку потока, в противном случае выполнение кода может зависнуть или привести к непреднамеренным побочным эффектам.

По сути, ваш код не является правильно сформированной программой CUDA.

user1547363 21 сен '12 в 07:28 2012-09-21 07:28 · Answer 4 · 2012-09-21 07:28

__syncthreads() используется для синхронизации потоков внутри блока. Это означает, что все потоки в блоке будут ждать завершения всех, прежде чем продолжить.

Рассмотрим случай, когда в блоке есть несколько потоков, которые входят в оператор if, а некоторые нет. Эти темы, ожидающие, будут заблокированы; вечно жду.

Обычно не очень удобно помещать синхронизацию в условный оператор if. Лучше всего избегать этого и перепроектировать свой код, если он у вас есть. Цель синхронизации - убедиться, что все потоки работают вместе. Почему вы сначала отфильтровываете их с помощью оператора if?

Чтобы добавить, если синхронизация требуется между блоками. Требуется перезапуск ядра.

user10163119 16 сен '20 в 17:11 2020-09-16 17:11 · Answer 5 · 2020-09-16 17:11

Лучше избегать __syncthreads()в условном if. Вы можете переписать код с помощью цикла for и __syncthreads() после цикла.

0

Источник

user10163119 16 сен '20 в 17:11