CUDA: Макс массива, как предотвратить конфликты записи?

Question

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

Предположим, что local_max = 0.0 в начале выполнения блока.Затем каждый поток считывает свое значение из входного вектора, решает, что оно больше, чем local_max, и затем пытается записать свое значение в local_max.Когда все это происходит в одно и то же время (по крайней мере, когда внутри одной и той же деформации), как это может работать и в конечном итоге получить действительный максимум в этом блоке?блокировки или критического раздела будет необходимо, но я не видел, что это адресовано в ответах, которые я нашел.(напр. http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf)

Answer 1

Ответ на ваши вопросы содержится в том самом документе, на который вы ссылаетесь, и пример сокращения SDK показывает конкретные реализации концепции сокращения.

Для полноты приведем конкретный пример ядра сокращения:

template <typename T, int BLOCKSIZE>
__global__ reduction(T *inputvals, T *outputvals, int N)
{
    __shared__ volatile T data[BLOCKSIZE];

    T maxval = inputvals[threadIdx.x];
    for(int i=blockDim.x + threadIdx.x; i<N; i+=blockDim.x) 
    {
        maxfunc(maxval, inputvals[i]);
    }

    data[threadIdx.x] = maxval;
    __syncthreads();

    // Here maxfunc(a,b) sets a to the minimum of a and b
    if (threadIdx.x < 32) {

        for(int i=32+threadIdx.x; i < BLOCKSIZE; i+= 32) {
            maxfunc(data[threadIdx.x], data[i]);
        }

        if (threadIdx.x < 16) maxfunc(data[threadIdx.x], data[threadIdx.x+16]);
        if (threadIdx.x < 8) maxfunc(data[threadIdx.x], data[threadIdx.x+8]);
        if (threadIdx.x < 4) maxfunc(data[threadIdx.x], data[threadIdx.x+4]);
        if (threadIdx.x < 2) maxfunc(data[threadIdx.x], data[threadIdx.x+2]);
        if (threadIdx.x == 0) {
            maxfunc(data[0], data[1]);
            outputvals[blockIdx.x] = data[0];
        }
    }
}

Ключевым моментом является использование неявной синхронизации внутри деформации для выполнения сокращения общей памяти.Результатом является максимальное значение для каждого блока.Второй проход сокращения необходим, чтобы уменьшить набор максимумов блоков до глобального максимума (часто это происходит быстрее на хосте).В этом примере maxvals - это функция «сравнить и установить», которая может быть простой:

template<T>
__device__ void maxfunc(T & a, T  & b)
{
    a = (b > a) ? b : a;
}

Answer 2

Не готовьте свой собственный код, используйте тягу (включена в версию 4.0 Cuda SDK):

#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <iostream>

int main(void)
{

    thrust::host_vector<int> h_vec(10000);
    thrust::sequence(h_vec.begin(), h_vec.end());
    // show hvec
    thrust::copy(h_vec.begin(), h_vec.end(), 
                 std::ostream_iterator<int>(std::cout, "\n"));

    // transfer to device
    thrust::device_vector<int> d_vec = h_vec;

    int max_dvec_value = *thrust::max_element(d_vec.begin(), d_vec.end()); 

    std::cout << "max value: " << max_dvec_value << "\n";
    return 0;    
}

И следите, чтобы thrust :: max_element возвращал указатель.

Answer 3

Ваш вопрос четко дан ответ в документе, на который вы ссылаетесь. Я думаю, вам просто нужно потратить больше времени на его чтение и понимание концепций CUDA, используемых в нем. В частности, я бы сосредоточился на разделяемой памяти, методе __syncthreads () и на том, как однозначно идентифицировать поток внутри ядра. Кроме того, вы должны попытаться понять, почему сокращение может быть выполнено за 2 прохода, чтобы найти глобальный максимум.