Loser’s “Brute-forced Cholesky Factorization for Sparse Matrix on CUDA”
2012-04-23 20:19
323 查看
Cholesky Factorization (CF) needs to access matrix elements along both column and row, so that to the sparse matrix, addressing would always be the greatest problem, this utility for dense matrix is quite easy. My brute-forced implementation is much slower than MKL's dcsrilu0.
CF method is a highly serialized algorithm, it processes the whole matrix row by row, and to the 1st row, its complexity is log(n*n/2), and then it would be faster and faster.
In fact I thought several methods as following,
Gathering and Scattering. As the bottleneck is the loop in Phase2, it need to access column and row, we may use gathering and scattering to simplify memeory access, but that’s as the same as CPU’s.
Transpose the matrix as a backup, so that access column will countious, but need to sync both matrix at the end of each step, still a overhead.
Here is the code, workable, slow, you may improve it by yourself, I switched to CPU, no sh*t & fu*k GPU.
Apl. 24
Found a paper used ELLPACK-R instead of CSR. I will take a try.
CF method is a highly serialized algorithm, it processes the whole matrix row by row, and to the 1st row, its complexity is log(n*n/2), and then it would be faster and faster.
In fact I thought several methods as following,
Gathering and Scattering. As the bottleneck is the loop in Phase2, it need to access column and row, we may use gathering and scattering to simplify memeory access, but that’s as the same as CPU’s.
Transpose the matrix as a backup, so that access column will countious, but need to sync both matrix at the end of each step, still a overhead.
Here is the code, workable, slow, you may improve it by yourself, I switched to CPU, no sh*t & fu*k GPU.
#include <iostream>
__device__ bool ReadCSR(const double *vals, const int *rows, const int *cols, const int r, const int c, double * v)
{
int i, j;
i = rows[r];
j = rows[r + 1];
i = cols[i];
j = cols[j - 1];
if (c < i)
{
return -1;
}
if (c > j)
{
return 1;
}
for (i = rows[r]; i < rows[r + 1]; ++ i)
{
j = cols[i];
if (j == c)
{
*v = vals[i];
return 0;
}
}
return -1;
}
__device__ void WriteCSR(double * vals, const int * rows, const int * cols, const int r, const int c, const double v)
{
int i, j;
for (i = rows[r]; i < rows[r + 1]; ++ i)
{
j = cols[i];
if (j == c)
{
vals[i] = v;
}
}
}
__global__ void Phase1(const int k, double *vals, const int *rows, const int *cols)
{
int i = blockIdx.x;
if (i > k)
{
double Akk = 0.0;
ReadCSR(vals, rows, cols, k, k, &Akk);
Akk = sqrt(Akk);
double Aik = 0.0f;
if (ReadCSR(vals, rows, cols, i, k, &Aik) == 0)
{
WriteCSR(vals, rows, cols, i, k, Aik / Akk);
}
}
}
__global__ void Phase2(const int k, double *vals, const int *rows, const int *cols)
{
int j = blockIdx.x;
int r = gridDim.x;
if (j > k)
{
for (int i = j; i < r; ++ i)
{
double Aij = 0.0;
int a = ReadCSR(vals, rows, cols, i, j, &Aij);
if (a == 0)
{
double Aik = 0.0, Ajk = 0.0;
ReadCSR(vals, rows, cols, i, k, &Aik);
ReadCSR(vals, rows, cols, j, k, &Ajk);
WriteCSR(vals, rows, cols, i, j, Aij - Aik * Ajk);
}
else if (a > 0)
{
break;
}
}
}
}
__global__ void Phase3(const int k, double *vals, const int *rows, const int *cols)
{
double Akk = 0.0;
ReadCSR(vals, rows, cols, k, k, &Akk);
Akk = sqrt(Akk);
WriteCSR(vals, rows, cols, k, k, Akk);
}
void test(const int numRow, double *vals, const int *rows, const int *cols)
{
for (int k = 0; k < numRow; ++ k)
{
std::cout << k << std::endl;
Phase1<<<numRow, 1>>>(k, vals, rows, cols);
Phase2<<<numRow, 1>>>(k, vals, rows, cols);
Phase3<<<1, 1>>>(k, vals, rows, cols);
}
}
Apl. 24
Found a paper used ELLPACK-R instead of CSR. I will take a try.
相关文章推荐
- 论文笔记:Sparse Matrix Format Selection with Multiclass SVM for SpMV on GPU
- Cholesky decomposition for Matrix Inversion
- 分布式稀疏矩阵的数据结构(data structure for distributed sparse matrix)
- Eclipse调试Android App若选择“Use same device for future launches”就再也无法选择其他设备的问题
- ArcMap 连接SDE 出错“Failed to connect to the specified server. Entry for SDE instance no found in services file.”
- [论文学习]Convolutional matrix factorization for document context-aware recommendation
- ICCV Tutorial on Sparse Coding and Dictionary Learning for Image Analysis
- 论文<Algorithms for non-negative matrix Factorization>
- 矩阵分解笔记(Notes on Matrix Factorization)
- Step By Step Guide to configure the “Replicating directory changes” for SharePoint 2010 and 2013
- Use cloud data from “Dallas” with “PowerPivot” addin for excel
- #Paper Reading# Joint Matrix Factorization and Manifold-Ranking for Topic-Focused Multi-Document Sum
- 反射引发的错误“reflection Unable to load one or more of the requested types. Retrieve the LoaderExceptions property for more information.”
- “Razor” – a new view engine for ASP.NET
- ABAP “FOR ALL ENTRIES IN” 使用指南
- Ground Segmentation based on Loopy Belief Propagation for Sparse 3D Point Clouds (论文速读)
- Inaccurate values for “Currently allocated space” and “Available free space” in the Shrink File dialog for TEMPDB only
- Notes on “feature-base locomotion controllers”
- Eclipse调试Android App若选择“Use same device for future launches”
- MSDTC on server “Server Name” is unavailable