Как синхронизировать потоки в CUDA без использования атома

Question

Я изучаю программирование CUDA с курса онлайн UDACITY. Образец кода был указан во втором уроке, который имеет basicaaly два ядра, первый __global__ void increment_naive(int *g) просто добавляет 1 к элементу массива *g, который находится в глобальной памяти.

Весь код в соответствии с Udacity выглядит следующим образом:

#include <stdio.h> 
#include "gputimer.h" 

#define NUM_THREADS 1000000 
#define ARRAY_SIZE 100 

#define BLOCK_WIDTH 1000 

void print_array(int *array, int size) 
{ 
    printf("{ "); 
    for (int i = 0; i < size; i++) { printf("%d ", array[i]); } 
    printf("}\n"); 
} 

__global__ void increment_naive(int *g) 
{ 
    // which thread is this? 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 

    // each thread to increment consecutive elements, wrapping at ARRAY_SIZE 
    i = i % ARRAY_SIZE; 
    g[i] = g[i] + 1; 
} 

__global__ void increment_atomic(int *g) 
{ 
    // which thread is this? 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 

    // each thread to increment consecutive elements, wrapping at ARRAY_SIZE 
    i = i % ARRAY_SIZE; 
    atomicAdd(& g[i], 1); 
} 

int main(int argc,char **argv) 
{ 
    GpuTimer timer; 
    printf("%d total threads in %d blocks writing into %d array elements\n", 
      NUM_THREADS, NUM_THREADS/BLOCK_WIDTH, ARRAY_SIZE); 

    // declare and allocate host memory 
    int h_array[ARRAY_SIZE]; 
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int); 

    // declare, allocate, and zero out GPU memory 
    int * d_array; 
    cudaMalloc((void **) &d_array, ARRAY_BYTES); 
    cudaMemset((void *) d_array, 0, ARRAY_BYTES); 

    // launch the kernel - comment out one of these 
    timer.Start(); 


    increment_naive<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array); 
    //increment_atomic<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array); 
    timer.Stop(); 

    // copy back the array of sums from GPU and print 
    cudaMemcpy(h_array, d_array, ARRAY_BYTES, cudaMemcpyDeviceToHost); 
    print_array(h_array, ARRAY_SIZE); 
    printf("Time elapsed = %g ms\n", timer.Elapsed()); 

    // free GPU memory allocation and exit 
    cudaFree(d_array); 
    return 0; 
} 

В соответствии с программой, миллион нитей с 1000 блоков записи в 10 элементов массива. Таким образом, каждый элемент массива будет иметь результат 100000.

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

Второй kernell отлично работает и производит правильный вывод, который выглядит следующим образом:

1000000 total threads in 1000 blocks writing into 100 array elements 
{ 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 10000 } 
Time elapsed = 0.367648 ms 

Как говорилось ранее, первый kernell производит неправильный выход каждый раз.

1000000 total threads in 1000 blocks writing into 100 array elements 
{ 75 75 75 75 78 78 78 78 73 73 73 73 82 82 82 82 85 85 85 85 92 92 92 92 104 104 104 104 107 107 107 107 89 89 89 89 88 88 88 88 95 95 95 95 103 103 103 103 106 106 106 106 107 107 107 107 105 105 105 105 113 113 113 113 96 96 96 96 95 95 95 95 95 95 95 95 100 100 100 100 98 98 98 98 104 104 104 104 110 110 110 110 126 126 126 126 90 90 90 90 } 
Time elapsed = 0.23392 ms 

Я пытаюсь исправить первый kernell пути размещения барьеров на разных этапах вычислений, но я не в состоянии получить необходимую продукцию. Моя попытка фиксации первой kernell выглядит следующим образом:

__global__ void increment_naive(int *g) 
{ 
    // which thread is this? 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    __syncthreads(); 
    // each thread to increment consecutive elements, wrapping at ARRAY_SIZE 
    //i = i % ARRAY_SIZE; 
    int temp = i%ARRAY_SIZE; 
    __syncthreads(); 
    i = temp; 
    __syncthreads(); 
    //g[i] = g[i] + 1; 
    int temp1 = g[i]+1; 
    __syncthreads(); 
    g[i] = temp1; 
    __syncthreads(); 

} 

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

Answer 1

Функция __syncthreads() гарантирует, что все потоки в блоке находятся на одном и том же месте в коде. Используя их, вы не достигнете того, чего хотите. Еще хуже - предположим, что CUDA была идеальной параллельной машиной, при этом все потоки работали в замке. Вам никогда не понадобится __syncthreads. Тем не менее, у вас будет другой результат. Рассмотрим следующий псевдокод и объяснение о том, что происходит:

__perfect_parallel_machine__ void increment_naive(int *g) 
{ 
    int idx = thisThreadIdx % ARRAY_SIZE; 
    int local = g[idx]; 
           //*all* threads load the initial value of g[idx] 
           //each thread holds a separate copy of 'local' variable 
           //local=0 in each thread 
    local = local + 1; 
           //each thread increment its own private copy of 'local' 
           //local=1 for all threads 
    g[idx] = local; 
           //each thread stores the same value (1) into global array 
           //g = {1, 1, 1, 1, 1, ...., 1} 
} 

Поскольку CUDA является не идеальный параллельно машина, вещи происходят из строя, и вы в конечном итоге получить более высокие значения в массиве. Помещение большего количества синхронизационных барьеров приблизит вас к идеальному результату {1, 1, ... , 1}.

Существуют и другие барьерные функции, такие как __threadfence(). Это останавливает текущий поток (только текущий!), Пока хранилище в глобальном массиве не будет гарантировано быть видимым другими потоками. Это связано с кэшированием L1/L2 и не имеет ничего общего с синхронизацией потоков. Обычно используется, например, __threadfence в сочетании с атомикой, чтобы отметить, что вы закончили заполнение некоторых данных.

Я думаю, что между вами и вашим наставником должно быть какое-то недоразумение. Я бы предложил поговорить с ним, чтобы уточнить это ...