CUDA, поиск Макса с использованием сокращения, ошибка - PullRequest
Новая
       17

CUDA, поиск Макса с использованием сокращения, ошибка

1 голос
/ 28 июня 2011

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

Для нитей 1-31 у меня правильная распечатка maxVal, но для нитей 32-49 это абсолютно случайное число. Я не знаю, что я делаю неправильно.

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

Заранее спасибо за любую помощь. 

//block size = 50


__syncthreads();

if (tid<32){

    cptmp[tid]=(cptmp[tid]< cptmp[tid+32]) ? cptmp[tid+32] : cptmp[tid];__syncthreads();    
    cptmp[tid]=(cptmp[tid]< cptmp[tid+16]) ? cptmp[tid+16] : cptmp[tid];__syncthreads();
    cptmp[tid]=(cptmp[tid]< cptmp[tid+8]) ? cptmp[tid+8] : cptmp[tid];  __syncthreads();    
    cptmp[tid]=(cptmp[tid]< cptmp[tid+4]) ? cptmp[tid+4] : cptmp[tid];  __syncthreads();
    cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid];  __syncthreads();    
    cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid];  __syncthreads();

}

__syncthreads();

//if (tid==0) {
    maxVal=cptmp[0];
    if(bix==0 && biy==0) cuPrintf(" max:%f x:%d y:%d\n", maxVal, blockIdx.x, blockIdx.y);
//}

Ответы [ 3 ]

3 голосов
/ 30 июня 2011

Вот более эффективный (по крайней мере на графических процессорах Fermi) и правильный код с использованием volatile.Замените T своим типом (или используйте шаблон): 

if (tid<32) {
    volatile T *c = cptmp;
    T t = c[tid];
    c[tid] = t = (t < c[tid+32]) ? c[tid+32] : t;
    c[tid] = t = (t < c[tid+16]) ? c[tid+16] : t;
    c[tid] = t = (t < c[tid+ 8]) ? c[tid+ 8] : t;
    c[tid] = t = (t < c[tid+ 4]) ? c[tid+ 4] : t;
    c[tid] = t = (t < c[tid+ 2]) ? c[tid+ 2] : t;
    c[tid] = t = (t < c[tid+ 1]) ? c[tid+ 1] : t;
}

Почему это более эффективно?Что ж, для правильности при отсутствии __syncthreads() мы должны использовать энергозависимый указатель на разделяемую память.Но это заставляет компилятор «соблюдать» все операции чтения и записи в общую память - он не может оптимизировать и хранить что-либо в регистрах.Таким образом, явно сохраняя c[tid] во временном t, мы сохраняем одну загрузку разделяемой памяти на строку кода.А поскольку Fermi является архитектурой загрузки / хранения, которая может использовать регистры только в качестве операндов инструкций, это означает, что мы сохраняем инструкцию на строку или всего 6 инструкций (всего около 25%, я ожидаю).

На старомВ архитектуре T10 / GT200 и более ранних версиях ваш код (с volatile и без __syncthreads ()) был бы одинаково эффективен, поскольку эта архитектура могла бы получать один операнд на команду непосредственно из общей памяти.

Этот код должен быть эквивалентен, если вы предпочитаетеif более ?:: 

if (tid<32) {
    volatile T *c = cptmp;
    T t = c[tid];
    if (t < c[tid+32]) c[tid] = t = c[tid+32];
    if (t < c[tid+16]) c[tid] = t = c[tid+16];
    if (t < c[tid+ 8]) c[tid] = t = c[tid+ 8];
    if (t < c[tid+ 4]) c[tid] = t = c[tid+ 4];
    if (t < c[tid+ 2]) c[tid] = t = c[tid+ 2];
    if (t < c[tid+ 1]) c[tid] = t = c[tid+ 1];
}
2 голосов
/ 29 июня 2011

Не используйте __syncthreads() в расходящемся коде!Все потоки или никакие потоки из данного блока должны достигать каждого __syncthreads() в одном и том же месте.

Все потоки из одной основы (32 потока) неявно синхронизируются, поэтому вам не нужно __syncthreads()собрать их всех вместе.Однако, если вы беспокоитесь о том, что записи в общую память одного потока могут быть не видны другому потоку той же основы, используйте __threadfence_block().

, чтобы прояснить важность __threadfence_block().Рассмотрим следующие две строки: 

cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid];
cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid];

Он может скомпилироваться во что-то вроде этого: 

int tmp; //assuming that cptmp is an array of int-s
tmp=cptmp[tid];
tmp=(tmp<cptmp[tid+2])?cptmp[tid+2]:tmp;
tmp=(tmp<cptmp[tid+1])?cptmp[tid+1]:tmp;
cptmp[tid]=tmp;

Хотя это было бы правильно для однопоточного кода, очевидно, что оно не работает для CUDA.

Чтобы избежать подобных оптимизаций, вы должны либо объявить массив cptmp как volatile, либо добавить это __threadfence_block() между строками.Эта функция гарантирует, что все потоки одного и того же блока увидят общую память, записанную в текущем потоке, до того, как функция будет существовать.

Аналогичная функция __threadfence() существует для обеспечения видимости глобальной памяти.

1 голос
/ 13 февраля 2014

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

auto localMax = max(c[tid], c[tid + 32]);    
for (auto i = 16; i >= 1; i /= 2)
{
    localMax = max(localMax, __shfl_xor(localMax, i));
}
c[tid] = localMax;

Требуются только два чтения и одна запись из глобальной памяти, так что это довольно аккуратно.

...