在这个并行GPU代码中，什么可能导致“未定义的行为”？

Question

假设core1和core2尝试将它们的变量a和b写入相同的内存位置。

在这里怎么解释？

• 我们不知道a或b是否被写入到那个内存位置(作为最后的操作)。
• 我们甚至不知道里面写的是什么(垃圾)
• 甚至目标内存地址也可能计算错误(分段错误？)。
• 一些逻辑门会产生错误的电流，而CPU则会使自己无法工作。
• CPU的频率信息变得损坏，超时(并自行中断)

我是否可以假设只有第一个选项对CPU(和GPU)的所有供应商有效？

我刚刚把下面的代码转换成一个并行的GPU代码，它看起来工作得很好。

通用代码：

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);"

这是一些罕见的2D NDrangeKernel的边界UB的文物。这些东西能杀死我的GPU吗？

​

​

Answer 1

在xf86和xf86_64体系结构上，这意味着我们不知道a或b是否写入到该内存位置(作为最后的操作)，因为32 (两者都适用)或64位(仅针对xf86_64)的内存对齐数据类型的加载/存储操作是原子的。

在其他架构上，我们甚至不知道在那里写了什么(垃圾)是一个有效的答案--当然，在RISC架构上，我目前还不知道在GPU上。

请注意，代码工作的事实并不意味着它是正确的，在99%的时间里，它是句子的来源，例如“存在编译器错误，代码一直工作到上一个版本”，或者“代码在开发机器上工作。选择用于生产的服务器被破坏了”) :)

编辑：

在NVidia图形处理器上，我们建立了弱有序内存模型.在关于Cuda C编程指南的说明中，没有明确说明存储操作是原子的。写操作来自同一个线程，所以这并不意味着加载/存储操作是原子操作。

Answer 2

对于上面的代码，IMHO的第一个选项是唯一可能的。基本上，如果假设有足够的线程/处理器并行执行所有循环，则内部嵌套循环( x和y循环)将有未定的值。

例如，如果我们只考虑

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

节中，fire_r[j+y][i+x]处的值可以是原始值，也可以是在另一个线程中完成相同循环的另一个实例的结果。