| 47 | |
| 48 | |
| 49 | __global__ void BlockJacobiKernel (double s, FlatArray<Dev<BlockJacobiCtr>> ctrs, |
| 50 | BareVector<Dev<double>> x, BareVector<Dev<double>> y) |
| 51 | { |
| 52 | for (int i = blockIdx.x*blockDim.y+threadIdx.y; i < ctrs.Size(); i += gridDim.x*blockDim.y) |
| 53 | { |
| 54 | BlockJacobiCtr mv = ctrs[i]; |
| 55 | size_t h = mv.mat.Height(); |
| 56 | size_t w = mv.mat.Width(); |
| 57 | |
| 58 | for (int r = threadIdx.x; r < h; r += blockDim.x) |
| 59 | { |
| 60 | double sum = 0; |
| 61 | for (int c = 0; c < w; c++) |
| 62 | sum += mv.mat(r,c) * x(mv.indices[c]); |
| 63 | atomicAdd((double*)&y(mv.indices[r]), s*sum); |
| 64 | } |
| 65 | } |
| 66 | } |
| 67 | |
| 68 | |
| 69 | __global__ void BlockJacobiTransKernel (double s, FlatArray<Dev<BlockJacobiCtr>> ctrs, |