Многократное копирование всего буфера глобальной памяти в буфер разделяемой памяти

Question

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

Как это сделать?

Я знаю размер буфера только во время выполнения:

__global__ void foo( int *globalMemArray, int N )
{
    extern __shared__ int s_array[];

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if( idx < N )
    {

       ...?
    }
}

Answer 1

Первое, что нужно сделать, это то, что общая память ограничена максимумом 16 КБ или 48 КБ на потоковый мультипроцессор (SM), в зависимости от того, какой графический процессор вы используете и как он настроен, так что если ваш буфер глобальной памяти не оченьмаленький, вы не сможете загрузить все это в общую память одновременно.

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

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

template <int blocksize>
__global__ void foo( int *globalMemArray, int *globalMemOutput, int N ) 
{ 
    __shared__ int s_array[blocksize]; 
    int npasses = (N / blocksize) + (((N % blocksize) > 0) ? 1 : 0);

    for(int pos = threadIdx.x; pos < (blocksize*npasses); pos += blocksize) { 
        if( pos < N ) { 
            s_array[threadIdx.x] = globalMemArray[pos];
        }
        __syncthreads(); 

        // Calculations using partial buffer contents
        .......

        __syncthreads(); 
    }

    // write final per thread result to output
    globalMemOutput[threadIdx.x + blockIdx.x*blockDim.x] = .....;
} 

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

CUDA SDK содержит несколько хороших примеров кодов, которые демонстрируют различные способы использования разделяемой памяти в ядрах для улучшения чтения и записи памяти.rformance.Примеры методов транспонирования матрицы, редукции и трехмерного метода конечных разностей являются хорошими моделями использования общей памяти.У каждого также есть хороший документ, в котором обсуждаются стратегии оптимизации использования общей памяти в кодах.Вы будете хорошо изучены, пока не поймете, как и почему они работают.