Что может вызвать "неопределенное поведение" в этом параллельном коде графического процессора?

Предположим, что core1 и core2 пытаются записать свои переменные a и b в одну и ту же область памяти.

Как можно объяснить UB здесь?

  • Мы не знаем, записаны ли a или b в эту ячейку памяти (как последнее действие).
  • Мы даже не знаем, что там написано (фигня)
  • Даже адрес целевой памяти может быть просчитан (segfault?).
  • Некоторые логические вентили создают неправильные токи, а процессор отключает себя
  • Информация о частоте процессора становится поврежденной и разгоняется (и ломается сама)

Можно ли предположить, что только первый вариант действителен для всех поставщиков процессоров (и графических процессоров)?

Я только что преобразовал приведенный ниже код в параллельный код графического процессора, и он, кажется, работает нормально.

Общий код:

for (j=0; j<YRES/CELL; j++) // this is parallelized
        for (i=0; i<XRES/CELL; i++) // this is parallelized
        {
            r = fire_r[j][i];
            g = fire_g[j][i];
            b = fire_b[j][i];
            if (r || g || b)
                for (y=-CELL; y<2*CELL; y++)
                    for (x=-CELL; x<2*CELL; x++)
                        addpixel(i*CELL+x, j*CELL+y, r, g, b, fire_alpha[y+CELL][x+CELL]);
   //addpixel accesses neighbour cells' informations and writes on them
   //and makes UB
            r *= 8;
            g *= 8;
            b *= 8;
            for (y=-1; y<2; y++)
                for (x=-1; x<2; x++)
                    if ((x || y) && i+x>=0 && j+y>=0 && i+x<XRES/CELL && j+y<YRES/CELL)
                    {
                        r += fire_r[j+y][i+x];
                        g += fire_g[j+y][i+x];
                        b += fire_b[j+y][i+x];
                    }
            r /= 16;
            g /= 16;
            b /= 16;
            fire_r[j][i] = r>4 ? r-4 : 0; // UB
            fire_g[j][i] = g>4 ? g-4 : 0; // UB
            fire_b[j][i] = b>4 ? b-4 : 0;
        }

OpenCL:

"   int i=get_global_id(0); int j=get_global_id(1);"
"   int VIDXRES="+std::to_string(kkVIDXRES)+";"
                        "   int VIDYRES="+std::to_string(kkVIDYRES)+";"
                        "   int XRES="+std::to_string(kkXRES)+";"
                        "   int CELL="+std::to_string(kkCELL)+";"
                        "   int YRES="+std::to_string(kkYRES)+";"

                        "   int x=0,y=0,r=0,g=0,b=0,nx=0,ny=0;"

                        "       r = fire_r[j*(XRES/CELL)+i];"
                        "       g = fire_g[j*(XRES/CELL)+i];"
                        "       b = fire_b[j*(XRES/CELL)+i];"

                        "       int counterx=0;"
                        "       if (r || g || b)"
                        "       for (y=-CELL; y<2*CELL; y++){"
                        "       for (x=-CELL; x<2*CELL; x++){"
                        "       addpixel(i*CELL+x, j*CELL+y, r, g, b, fire_alpha[(y+CELL)*(3*CELL)+(x+CELL)],vid,vido);"
                        "       }}"

                        "       r *= 8;"
                        "       g *= 8;"
                        "       b *= 8;"
                        "       for (y=-1; y<2; y++){"
                        "       for (x=-1; x<2; x++){"
                        "       if ((x || y) && i+x>=0 && j+y>=0 && i+x<XRES/CELL && j+y<YRES/CELL)"
                        "       {"
                        "           r += fire_r[(j+y)*(XRES/CELL)+(i+x)];"
                        "           g += fire_g[(j+y)*(XRES/CELL)+(i+x)];"
                        "           b += fire_b[(j+y)*(XRES/CELL)+(i+x)];"
                        "       }}}"
                        "       r /= 16;"
                        "       g /= 16;"
                        "       b /= 16;"
                        "       fire_r[j*(XRES/CELL)+i] = (r>4 ? r-4 : 0);"
                        "       fire_g[j*(XRES/CELL)+i] = (g>4 ? g-4 : 0);"
                        "       fire_b[j*(XRES/CELL)+i] = (b>4 ? b-4 : 0);"

Вот картина некоторых редких артефактов двумерной локальной границы NDrangeKernel UB. Могут ли они убить мой GPU?

Источник

2 ответа

Решение

В архитектурах xf86 и xf86_64 это означает, что мы не знаем, записаны ли a или b в эту область памяти (как последнее действие), потому что операции загрузки / сохранения 32-битных (для обеих) или 64-битных (только для xf86_64) типов данных, выровненных по памяти, являются атомарными,

На других архитектурах мы обычно даже не знаем, что там написано (мусор) является правильным ответом - наверняка на архитектурах RISC, я в настоящее время не знаю на GPU.

Обратите внимание, что тот факт, что код работает, не означает, что он правильный, и в 99% случаев он является источником предложений типа "есть ошибка компилятора, код работал до предыдущей версии" или "код работает над машина разработки. Сервер, выбранный для производства, сломан ":)

РЕДАКТИРОВАТЬ:

На графических процессорах NVidia у нас слабо упорядоченная модель памяти. В описании руководства по программированию на Cuda C прямо не указано, что операции с хранилищем являются атомарными. Операции записи происходят из одного и того же потока, поэтому это не означает, что операции загрузки / сохранения являются атомарными.

Источник

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

Например, если мы рассмотрим только

r += fire_r[j+y][i+x];

раздел, значение в fire_r[j+y][i+x] может быть как исходным, так и результатом завершения другого экземпляра того же цикла в другом потоке.

Источник
Другие вопросы по тегам