CUDA 中的高斯消除（无旋转）

Question

我正在尝试使用 CUDA 解决高斯消元问题。

我有一个N*N 矩阵。为了得到这个矩阵的新元素，我使用下面的 CPU 代码，其中C.width=N:

for(int z=0; z< C.width-1; z++)
{
    for ( int c = z+1 ; c < C.width ; c++ )
    {
        for (int d = z ; d < C.width ; d++ )
        {
            C.elements[c*C.width+d]=C.elements[c*C.width+d] - (B.elements[c*C.width+z]*C.elements[z*C.width+d]);
        }
    }
}

我正在尝试使用 CUDA 来实现它。例如，对于N=512

dim3 dimBlock(16,16,1);

dim3 dimGrid(32,32,1); 

MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);

我认为对于每次迭代我都应该使用N-i*N 线程来计算元素更新，即

    if(idx>511 || idy>510)
        return;
        for(int i=1; i<512;i++)
        {
            if(idx>=i-1 && idy>=i-1)
                C.elements[(idy+1)*C.width+idx]=C.elements[(idy+1)*C.width+idx]-((C.elements[(idy+1)*C.width+(i-1)]/C.elements[(i-1)*C.width+(i-1)])*C.elements[(i-1)*C.width+idx]);

            __syncthreads();
        }
        }

在GPU和CPU上得到的结果是一样的，但是处理时间是Time(CPU)=2*Time(GPU)

对于N=512：Time(CPU) = 1900 ms； Time(GPU) = 980 ms

对于N=1024：Time(CPU) = 14000 ms；时间（GPU）= 7766 毫秒` . . .

我认为加速应该比我现在的要大。我的并行代码有什么错误吗？你能帮我改写我的代码吗？

感谢您的帮助！

Answer 1

高斯消元可以看作是一个两步过程。第一步旨在将线性系统转换为上三角线性系统，第二步包括求解如此获得的上三角线性系统。第二步在 CUDA 中是微不足道的，可以由cublasStrsm 高效执行。您在帖子中提到的第一步是棘手的​​部分。

有几种优化方法可以解决第一步。我认为您的方法有些幼稚，我建议您研究文献以实现不错的加速。

基本上，可以通过 tiling 方法执行将原始系统转换为上三角形的方法，从某些角度来看，该方法类似于用于执行矩阵的 tiling 方法-CUDA C 编程指南经典示例中的矩阵乘法。

可以通过专门编写的内核或大量使用 cuBLAS 例程来执行平铺方法。

上个月（2013 年 11 月），以下论文

Manuel Carcenac，“从平铺算法到条带算法：基于 CUBLAS 的高斯 GPU 并行实现方法，用于解决存储在固态设备阵列上的超大型密集线性系统”，超级计算杂志，DOI 10.1007/s11227-013-1043-3

提出了一种基于 cuBLAS 使用的平铺/剥离方法。

上面提到的所有方法都在 M. Carcenac 的网页上的一个名为 Application: linear system resolution with Gauss method 的演示文稿中进行了总结。

此外，Gaussian elimination with CUDA 帖子中提供了一个可下载的 Visual Studio 2010 项目，该项目实现了所有这些并进行了一些性能测试。从可用的代码中，您可以针对您感兴趣的架构进行自己的测试，并体验 M. Carcenac 引入的方法相对于其他方法的改进。