Здравствуйте, 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 в начале, у компилятора в принципе достаточно инфы чтоб понять и разрулить это все достаточно корректно.