Здравствуйте, elmal, Вы писали:

E>Выкинул ретурн и условие — никаких изменений, гонки.
Как именно гонка проявляется? Я у тебя вижу явный глюк

 const unsigned int curVal     = vals[myId];
 if (tid == curVal) // зачем?? ты же один раз проходишь по данным, значение просто потеряется
{
        atomicAdd(&(sdata[curVal]), 1);
  }

Может ты его как гонку засчитал?

Здравствуйте, elmal, Вы писали:

E>Не очень мне кстати действительно понятно во что именно компилится этот return, после которого идут барьеры. Соответственно я очень часто на этот return грешил, но как то не заметно чтоб он негативно влиял до сих пор. Я его убирал вообще, эффекта ноль.
Барьер компилируется в инструкцию bar.sync. Просто если внимательно сравнить описание bar.synс и __syncthreads, то окажется что у них разная семантика. А именно барьер ждёт когда хотя бы один поток из каждого варпа дойдёт до него, а не когда все потоки дойдут до барьера (как написано в описании __syncthreads). И экспериментами это подтверждается.
Соответственно, иногда вызов __syncthreads() не будет ждать потоков, которые сделали return, например, из-за того, что в варпе остались потоки, которые этот return не сделали, — и всё будет как бы работать. Но нет гарантий, что так повезёт для всех комбинаций размеров данных (как в одном из примеров из эксперимента). Или что такое поведение вообще сохранится в дальнейшем (__syncthreads может начать компилироваться в другую инструкцию).

Здравствуйте, elmal, Вы писали:

E> atomicAdd(&(histo[tid]), blockHistoVal);

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

    if (myId < numVals) {
        const unsigned int curVal     = vals[myId];
        atomicAdd(&(sdata[curVal]), 1);
    }

early return делать нельзя поскольку поток должен всё же обработать свой элемент sdata[] для обеспечения работы остальных потоков грида

грида на 1024 элемента хватит на то чтобы заполнить один smx. но этих smx в gpu минимум 5 штук, так что минимум 5 гридов у тебя крутится одновременно

Играюсь тут с CUDA, хочу сделать параллельную гистограмму, в перспективе вообще без атомиков.

__global__
void histo(const unsigned int* const vals,    //INPUT
                 unsigned int* const histo,   //OUPUT
                          int        numVals)
{
    extern __shared__ unsigned int sdata[];

    const unsigned int tid        = threadIdx.x;
    const unsigned int blockId    = blockDim.x * blockIdx.x;
    const unsigned int myId       = tid + blockId;

    if (myId > numVals) {
        return;
    }

    sdata[tid] = 0;
    __syncthreads();

    const unsigned int curVal     = vals[myId];
    if (tid == curVal) {
        atomicAdd(&(sdata[curVal]), 1);
    }

    __syncthreads();
    const unsigned int blockHistoVal = sdata[tid];
    if (blockHistoVal != 0) {
        atomicAdd(&(histo[tid]), blockHistoVal);
    }
}

histo<<<numBlocksForElements, MAX_THREADS, MAX_THREADS*sizeof(unsigned int)>>> (d_vals, d_histo, numElems);

То есть идея такая. Запускаю я этот kernel с 1024 потоками, с соответствующими параметрами запуска. Предполагаю, что это будет сгруппировано по блокам с 1024 потоками. Соответственно для каждого блока я соответствующие очищаю shared memory, синхронизую, делаю пока просто атомиком тривиальную гистограмму над shared memory, синхронизую, и затем атомиком уже из всех shared memory агрегирую в результат. Знаю, что это все тормознуто, но дело не в этом. Какого то черта здест гонки явные возникают, а вот почему — понять не могу. Предполагаю, что __syncthreads() синхронизует потоки внутри thread block с которым связана shared memory, соответственно ну никак не ожидаю я здесь гонок. А они есть. Явно я что то не понимаю в устройстве как это все работает.

Здравствуйте, elmal, Вы писали:

E>{
E>    if (myId > numVals) {
E>        return;
E>    }
E>    __syncthreads();
E>}
E>

На первый взгляд. Про куду не помню, но в opencl сonditional syncthreads (барьер) самое UB, которое только может быть, которое в большинстве случаев приводит к гонке.

Здравствуйте, С3141566=Z, Вы писали:

СZ>На первый взгляд. Про куду не помню, но в opencl сonditional syncthreads (барьер) самое UB, которое только может быть, которое в большинстве случаев приводит к гонке.
Конечно есть такая вероятность. Но подобный код весьма распространен, я когда игрался, весьма активно использовал, и проблем не было.

Брал за основу кода статью: http://isaac.gelado.cat/sites/isaac.gelado.cat/files/publications/samos_2013_histogramming.pdf
Там тоже подобное есть.

Правда ни хрена не понимаю зачем там цикл, почему не сделать как у меня.

Здравствуйте, elmal, Вы писали:

E>Здравствуйте, С3141566=Z, Вы писали:

СZ>>На первый взгляд. Про куду не помню, но в opencl сonditional syncthreads (барьер) самое UB, которое только может быть, которое в большинстве случаев приводит к гонке.
E>Конечно есть такая вероятность. Но подобный код весьма распространен, я когда игрался, весьма активно использовал, и проблем не было.
Нет он просто не должен работать. Смотри в точке барьера все потоки должны иметь одинаковое состояние. А те, потоки которые у тебя вышли по ретурну такого иметь не могут по определению. Соответственно, что делать остальным? Самое простое просто выкинуть барьер из кода, что часто и происходит.

E>Брал за основу кода статью: http://isaac.gelado.cat/sites/isaac.gelado.cat/files/publications/samos_2013_histogramming.pdf
Там все чисто в этом плане.

Здравствуйте, С3141566=Z, Вы писали:

СZ>Нет он просто не должен работать. Смотри в точке барьера все потоки должны иметь одинаковое состояние. А те, потоки которые у тебя вышли по ретурну такого иметь не могут по определению. Соответственно, что делать остальным? Самое простое просто выкинуть барьер из кода, что часто и происходит.
Выкинул ретурн и условие — никаких изменений, гонки.

Здравствуйте, С3141566=Z, Вы писали:

СZ>Как именно гонка проявляется? Я у тебя вижу явный глюк
Вызываю код 2 раза — получаю разные результаты. Мать, точно, все из за этого дерьма. Убрал этот if чтоб не терялось значение — теперь корректно все, спасибо. Тяжеловато мышление под параллельность перестраивать, да еще и без отладки нормальной.

Здравствуйте, elmal, Вы писали:

E>Правда ни хрена не понимаю зачем там цикл, почему не сделать как у меня.

для того чтобы уменьшить конкуренцию при обновлении глобального массива. одно дело — когда у тебя одно обновление на 1 входной элемент, другое — когда на 10

Здравствуйте, elmal, Вы писали:

СZ>>Как именно гонка проявляется? Я у тебя вижу явный глюк
E>Вызываю код 2 раза — получаю разные результаты.

разные результаты ты получал скорей всего из-за return — sdata оставалось неинициализированным. гонка между атомиками сама по сбе к такому результату приводить не должна (ни в каком api/языке)

Здравствуйте, elmal, Вы писали:

E>Здравствуйте, С3141566=Z, Вы писали:

СZ>>Как именно гонка проявляется? Я у тебя вижу явный глюк
E>Вызываю код 2 раза — получаю разные результаты. Мать, точно, все из за этого дерьма. Убрал этот if чтоб не терялось значение — теперь корректно все, спасибо. Тяжеловато мышление под параллельность перестраивать, да еще и без отладки нормальной.
Напиши эмулятор простейший и гоняйся под ним сперва. Хотя для твоего случая, это не совсем простейший уже будет. Хотя была какая то приблуда Nsight которая вроде позволяла отлаживаться.

Здравствуйте, BulatZiganshin, Вы писали:

BZ>разные результаты ты получал скорей всего из-за return — sdata оставалось неинициализированным. гонка между атомиками сама по сбе к такому результату приводить не должна (ни в каком api/языке)
Вот не понятен 1 момент мне:
Я тупо убрал явно косячное условие if (tid == curVal). Не убирал return. Все работает как часы. Я специально уменьшил количество элементов на минус 123, чтоб гарантированно не было кратно 1024 (изначально данные были кратны, соответственно этот return вообще не вызывался) — один хрен никаких гонок.

И дополнительно. Да, с return как минимум потенциальный косяк я вижу, sdata действительно может быть неинициализированным.

    sdata[tid] = 0;    
    if (myId > numVals) {
        return;
    }

    __syncthreads();

А вот в этом случае косяков не будет? Ибо не люблю я уровни вложенности на пустом месте. Если так допустимо, и оставить if (tid == curVal) — все время разные значения. Убрать — все нормально. То есть не изменилось ни черта.

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

Текущее мое представление о работе всего этого следующее. Эти все SMX — это обычные SIMD операции. Соответственно при компиляции это все загоняется в данные и выполняются инструкции. А барьер — это просто указание, что перед запуском команды чтоб все данные были готовы. Потоков в понимании как в CPU там нет ни черта, на низком уровне все примитивно, все навороты компилятор делает. Соответственно если я делаю return в начале, у компилятора в принципе достаточно инфы чтоб понять и разрулить это все достаточно корректно.

Здравствуйте, С3141566=Z, Вы писали:

СZ>Напиши эмулятор простейший и гоняйся под ним сперва. Хотя для твоего случая, это не совсем простейший уже будет. Хотя была какая то приблуда Nsight которая вроде позволяла отлаживаться.
Чтоб написать эмулятор, нужно понимать как работает то, что эмулируется . Сами то алгоритмы я по рабоче крестьянски отлаживаю на CPU, эмулирую параллельное выполнение набором циклов до барьера. Базовые вещи так неплохо отлавливаются, а вот с деталями — их блин понимать нужно.

Здравствуйте, elmal, Вы писали:

E>Я тупо убрал явно косячное условие if (tid == curVal). Не убирал return.

так ведь нужно не только обнулить sdata в начале, но и записать его значение в конце

E>Текущее мое представление о работе всего этого следующее. Эти все SMX — это обычные SIMD операции. Соответственно при компиляции это все загоняется в данные и выполняются инструкции. А барьер — это просто указание, что перед запуском команды чтоб все данные были готовы. Потоков в понимании как в CPU там нет ни черта, на низком уровне все примитивно, все навороты компилятор делает. Соответственно если я делаю return в начале, у компилятора в принципе достаточно инфы чтоб понять и разрулить это все достаточно корректно.

в одном smx — 4 ядра. каждое имеет 1 alu с органзиацией 32*32bit. но для эффективности одно ядро выполняет с десяток потоков одновременно. соответственно весь твой код преобразуется в simd-команды, если есть расхождения внутри одного варпа — то они закрываются маскированием соответствующих линий

1024 псевдопотока — это 32 реальных потока. вот их и приходится синхронизировать. а в целом titan x с его 24 smx может выполнять порядка 1000 (24*4*10) потоков одновременно, т.е. 30 тыщ скалярных псевдопотоков

вообще там вполне обычный процессор, от avx3 отличается уже не так сильно

Здравствуйте, BulatZiganshin, Вы писали:

BZ>1024 псевдопотока — это 32 реальных потока. вот их и приходится синхронизировать. а в целом titan x с его 24 smx может выполнять порядка 1000 (24*4*10) потоков одновременно, т.е. 30 тыщ скалярных псевдопотоков
Кстати о потоках. Теперь меня интересует, как избавиться от atomic, причем избавиться так, чтоб было быстрее чем текущий код с shared memory. Вижу одну лазейку небольшую, если бы ведер было 32, учитывая что потоки группируются как SIMD инструкции по 32, то на этапе обращений к shared memory уже никакие атомик были бы не нужны. Вернее можно было бы shared memory разбить на блоки по 32 элемента, и все нити, которые выполняются как SIMD инструкции, писали бы в свой блок. А при завершении уже агрегировать с глобальной памятью с помощью атомиков.

Но проблема в том, что ведер больше чем 32. Если их предварительно сортировать, можно было бы что сделать, но чтоб сортировать — нужна память. Выделение памяти занимает черти сколько времени, сразу получается в 20 раз медленнее только на одном выделении памяти. Есть идея не пытаться параллельно весь буфер обработать, а последовательно вызывать kernels из основного кода — тоже получается очень медленно.

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