Newer
Older
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
#include "cs.h"
///* solve Lx=b where x and b are dense. x=b on input, solution on output. */
///*
//csi cs_lsolve (const cs *L, double *x)
//{
// csi p, j, n, *Lp, *Li ;
// double *Lx ;
// if (!CS_CSC (L) || !x) return (0) ; /* check inputs */
// n = L->n ; Lp = L->p ; Li = L->i ; Lx = L->x ;
// for (j = 0 ; j < n ; j++)
// {
// x [j] /= Lx [Lp [j]] ;
// for (p = Lp [j]+1 ; p < Lp [j+1] ; p++)
// {
// x [Li [p]] -= Lx [p] * x [j] ;
// }
// }
// return (1) ;
//}
//*/
/* solve Lx=b where x and b are dense. x=b on input, solution on output. */
//csi cs_lsolve_gpu (const cs *L, double *x)
// typedef struct cs_sparse /* matrix in compressed-column or triplet form */
// {
// csi nzmax ; /* maximum number of entries */
// csi m ; /* number of rows */
// csi n ; /* number of columns */
// csi *p ; /* column pointers (size n+1) or col indices (size nzmax) */
// csi *i ; /* row indices, size nzmax */
// double *x ; /* numerical values, size nzmax */
// csi nz ; /* # of entries in triplet matrix, -1 for compressed-col */
// } cs ;
__global__
void cs_lsolve_gpu (int n, const int *Lp, const int *Li, const double *Lx, double *x, int n_rhs)
int tid = threadIdx.x; // + blockIdx.x * blockDim.x;
if (tid < n_rhs) {
int p, j;
// double *Lx ;
// n = L->n ;
// Lp = L->p ;
// Li = L->i ;
// Lx = L->x ;
int rhs_offset = tid * n;
for (j = 0 ; j < n ; j++)
x [rhs_offset + j] = x [rhs_offset + j] / Lx [Lp [j]] ;
for (p = Lp [j]+1 ; p < Lp [j+1] ; p++)
{
x [rhs_offset + Li [p]] = x [rhs_offset + Li [p]] - Lx [p] * x [rhs_offset + j] ;
printf("GPU; tid=%d Li[p]=%d addr=%d %f\n", tid, Li [p], rhs_offset + Li [p],x [rhs_offset + Li [p]]);
}
__global__
void cs_lsolve_gpu_trans (int n, const int *Lp, const int *Li, const double *Lx, double *x, int n_rhs)
{
int tid = threadIdx.x; // + blockIdx.x * blockDim.x;
if (tid < n_rhs) {
int p, j;
// double *Lx ;
// n = L->n ;
// Lp = L->p ;
// Li = L->i ;
// Lx = L->x ;
int rhs_offset = tid;;
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
for (j = 0 ; j < n ; j++)
{
x [rhs_offset + j * n_rhs] = x [rhs_offset + j * n_rhs] / Lx [Lp [j]] ;
for (p = Lp [j]+1 ; p < Lp [j+1] ; p++)
{
x [rhs_offset + Li[p] * n_rhs] = x [rhs_offset + Li[p] * n_rhs] - Lx [p] * x [rhs_offset + j * n_rhs] ;
//printf("GPU; tid=%d Li[p]=%d addr=%d x=%f\n", tid, Li [p], rhs_offset + Li [p]*n_rhs, x [rhs_offset + Li [p]*n_rhs]);
}
}
}
}
extern "C" {
void cs_lsolve_acc (int n, const int *Lp, const int *Li, const double *Lx, double *x, int n_rhs, int threads, int blocks)
{
cs_lsolve_gpu <<<blocks, threads>>> (n, Lp, Li, Lx, x, n_rhs);
}
}
void cs_lsolve_acc_trans (int n, const int *Lp, const int *Li, const double *Lx, double *x, int n_rhs, int threads, int blocks)
cs_lsolve_gpu_trans <<<blocks, threads>>> (n, Lp, Li, Lx, x, n_rhs);