Накопительная сумма в двух измерениях в массиве во вложенном цикле - реализация CUDA?

Question

Я думал о том, как выполнить эту операцию на CUDA, используя сокращения, но я немного растерялся, как ее выполнить. Код C ниже. Важно помнить, что переменная precalculatedValue зависит от обоих итераторов цикла. Кроме того, переменная ngo не уникальна для каждого значения m ... например. m = 0,1,2 может иметь ngo = 1, тогда как m = 4,5,6,7,8 может иметь ngo = 2 и т. Д. Я включил размеры итераторов цикла на случай, если это поможет обеспечить лучшие предложения по реализации.

// macro that translates 2D [i][j] array indices to 1D flattened array indices
#define idx(i,j,lda) ( (j) + ((i)*(lda)) )

int Nobs = 60480;
int NgS  = 1859;
int NgO  = 900;
// ngo goes from [1,900]

// rInd is an initialized (and filled earlier) as:
// rInd = new long int [Nobs];

for (m=0; m<Nobs; m++) {        
    ngo=rInd[m]-1;

    for (n=0; n<NgS; n++) {
            Aggregation[idx(n,ngo,NgO)] += precalculatedValue;
    }
}

В предыдущем случае, когда precalculatedValue был только функцией переменной внутреннего цикла, я сохранял значения в уникальных индексах массива и добавил их с параллельным сокращением (Thrust) после факта. Однако этот случай меня озадачил: значения m не однозначно сопоставляются со значениями ngo . Таким образом, я не вижу способа сделать этот код эффективным (или даже работоспособным) для использования сокращения. Любые идеи приветствуются.

Answer 1

Просто удар ...

Я подозреваю, что транспонирование ваших циклов может помочь.

for (n=0; n<NgS; n++) {
    for (m=0; m<Nobs; m++) {            
        ngo=rInd[m]-1;
        Aggregation[idx(n,ngo,NgO)] += precalculatedValue(m,n);
    }
}

Причина, по которой я это сделал, заключается в том, что idx меняется быстрее с ngo(функция m), чем с n, поэтому создание внутреннего цикла улучшает согласованность.Заметьте, я также сделал precalculatedValue функцией (m, n), потому что вы сказали, что это так - это делает псевдокод более понятным.

Затем вы можете начать с оставления внешнего цикла на хосте и созданияядро для внутреннего цикла (64,480-параллельного параллелизма достаточно для заполнения большинства современных графических процессоров).

Во внутреннем цикле просто начните с использования atomicAdd () для обработки коллизий.Если они нечасты, производительность графических процессоров Fermi не должна быть слишком плохой.В любом случае вы будете ограничены пропускной способностью, поскольку арифметическая интенсивность этого вычисления мала.Так что, как только это сработает, измерьте полосу пропускания, которую вы достигаете, и сравните с пиком для вашего GPU.Если вы далеко, , тогда подумайте о дальнейшей оптимизации (возможно, путем распараллеливания внешнего цикла - одна итерация на блок потока и выполните внутренний цикл с использованием некоторой оптимизации совместной памяти и взаимодействия потоков).

Ключ: начать с простого, измерить производительность, а затем решить, как оптимизировать.

Обратите внимание, что этот расчет выглядит очень похоже на расчет гистограммы, который имеет аналогичные проблемы, поэтому вы можете захотеть использовать Google для GPUгистограммы, чтобы увидеть, как они были реализованы.

Одна идея состоит в том, чтобы отсортировать (rInd[m], m) пары, используя thrust::sort_by_key(), а затем (так как дубликаты rInd будут сгруппированы вместе), выможете перебирать их и делать сокращения без коллизий.(Это один из способов создания гистограмм.) Вы могли бы сделать это даже с thrust::reduce_by_key().