高斯消元并行

Question

我已经成功地在CUDA中实现了用于高斯消除的单线程程序，并希望实现并行性。到目前为止，并行代码如下所示：

__global__ void ParallelGaussian(double* A)
{
    int index = threadIdx.x;
    int stride = blockDim.x;

    if (index < ROWS) //Skip additional threads
    {
        for (unsigned int r = index; r < ROWS; r += stride)
        {
            //Forward elimination to reduce to row echelon form
            for (unsigned int k = r + 1; k < ROWS; ++k)
            {
                double c = -A[(ROWS + 1) * k + r] / A[(ROWS + 1) * r + r];
                for (unsigned int j = r; j < ROWS + 1; ++j)
                {
                    if (r == j)
                        A[(ROWS + 1) * k + j] = 0.0;
                    else
                        A[(ROWS + 1) * k + j] += c * A[(ROWS + 1) * r + j];
                }
            }
        }
    }
}

我们可以看到GPU上的代码会将1D阵列（矩阵）转换为下三角矩阵，然后在CPU上我会继续进行反向替换以获得最终结果。在这种方法中没有进行旋转，因为它不是完全需要的，但确实提高了算法的数值稳定性。

使用单个线程和块启动内核可以将矩阵转换为行梯形形式：

ParallelGaussian << < 1, 1 >> >(dev_a);

但是，如果我想增加线程数，就像

ParallelGaussian << < 1, 32 >> >(dev_a);

它将无法产生下三角矩阵。现在将__syncthreads（）调用添加到代码中以便同步块中的线程并不能改善这种情况，我无法弄清楚原因。

Answer 1

考虑你的内循环。每个线程访问A，并且由于k和j从r运行到矩阵的末尾，因此多线程可能会修改相同的A[(ROWS + 1) * k + j]值。

您还可能有一些线程访问A[(ROWS + 1) * r + j]而其他线程正在更新该值。

一种可能的解决方案是让每个线程累积到单独的结果数组中，然后在最后组合它们。这是内存密集型的。

另一种方法是对其进行重组，以便只有一个线程将写入特定值，并将这些值存储在新矩阵中（这样您就不会更改其他线程可能需要的任何值）。