| 67 | |
| 68 | |
| 69 | __global__ void BlockJacobiTransKernel (double s, FlatArray<Dev<BlockJacobiCtr>> ctrs, |
| 70 | BareVector<Dev<double>> x, BareVector<Dev<double>> y) |
| 71 | { |
| 72 | for (int i = blockIdx.x*blockDim.y+threadIdx.y; i < ctrs.Size(); i += gridDim.x*blockDim.y) |
| 73 | { |
| 74 | BlockJacobiCtr mv = ctrs[i]; |
| 75 | size_t h = mv.mat.Height(); |
| 76 | size_t w = mv.mat.Width(); |
| 77 | |
| 78 | for (int r = threadIdx.x; r < w; r += blockDim.x) |
| 79 | { |
| 80 | double sum = 0; |
| 81 | for (int c = 0; c < h; c++) |
| 82 | sum += mv.mat(c,r) * x(mv.indices[c]); |
| 83 | atomicAdd((double*)&y(mv.indices[r]), s*sum); |
| 84 | } |
| 85 | } |
| 86 | } |
| 87 | |
| 88 | |
| 89 | |