Выполняет ли __syncthreads() синхронизацию всех потоков в сетке?

... или просто потоки в текущей деформации или блоке?

Кроме того, когда потоки в конкретном блоке встречаются (в ядре), следующая строка

__shared__  float srdMem[128];

будут ли они просто объявлять это пространство один раз (за каждый блок)?

Все они, очевидно, работают асинхронно, поэтому, если Thread 23 в блоке 22 является первым потоком для достижения этой строки, а затем Thread 69 в блоке 22 является последним для достижения этой строки, Thread 69 будет знать, что он уже был объявлен

Ответ 1

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

Пример использования __syncthreads(): (источник)

__global__ void globFunction(int *arr, int N) 
{
    __shared__ int local_array[THREADS_PER_BLOCK];  //local block memory cache           
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;

    //...calculate results
    local_array[threadIdx.x] = results;

    //synchronize the local threads writing to the local memory cache
    __syncthreads();

    // read the results of another thread in the current thread
    int val = local_array[(threadIdx.x + 1) % THREADS_PER_BLOCK];

    //write back the value to global memory
    arr[idx] = val;        
}

Чтобы синхронизировать все потоки в сетке, в настоящее время существует собственный вызов API не. Одним из способов синхронизации потоков на уровне сетки является использование последовательных вызовов ядра, так как в этот момент все потоки завершаются и снова начинаются с одной и той же точки. Он также обычно называется синхронизацией процессора или неявной синхронизацией. Таким образом, все они синхронизированы.

Пример использования этой техники (источник):

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

Ответ 2

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

Если вы объявляете общую память в ядре, массив будет видимым только для одного потокаблока. Таким образом, каждый блок будет иметь свой собственный блок разделяемой памяти.

Ответ 3

Я согласен со всеми ответами здесь, но я думаю, что мы упускаем один важный момент здесь, первый вопрос w.r.t. Я не отвечаю на второй ответ, так как в ответах на этот вопрос он отлично ответил.

Выполнение на GPU происходит в единицах перекосов. Варп представляет собой группу из 32 нитей и один экземпляр экземпляра каждый поток определенной основы выполняет одну и ту же инструкцию. Если вы выделите 128 потоков в блоке, это (128/32 =) 4 искажения для графического процессора.

Теперь возникает вопрос: "Если все потоки выполняют одну и ту же инструкцию, то зачем нужна синхронизация?". Ответ: нам нужно синхронизировать перекосы, принадлежащие блоку SAME. __syncthreads не синхронизирует потоки в warp, они уже синхронизированы. Он синхронизирует перекосы, принадлежащие одному и тому же блоку.

Вот почему ответ на ваш вопрос: __syncthreads не синхронизирует все потоки в сетке, но потоки, принадлежащие одному блоку, как каждый блок, выполняются независимо.

Если вы хотите синхронизировать сетку, разделите ядро (K) на два ядра (K1 и K2) и вызовите оба. Они будут синхронизированы (K2 будет выполнен после завершения K1).