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. 您在帖子中解决的第一步是棘手的​​部分。

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

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

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

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

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

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

所有上述方法都在 M. Carcenac 的网页上的演示文稿中进行了总结，该网页题为“应用：使用高斯方法的线性系统分辨率”。

此外，可下载的 Visual Studio 2010 项目通过一些性能测试实现了所有这些，可在CUDA后的高斯消元中获得。从可用代码中，您可以对感兴趣的架构进行自己的测试，并体验 M. Carcenac 提出的方法相对于其他方法的改进。