| 370 | |
| 371 | |
| 372 | void DevSparseCholesky :: |
| 373 | MultAdd (double s, const BaseVector & x, BaseVector & y) const |
| 374 | { |
| 375 | |
| 376 | static Timer t("DevSparseCholesky::MultAdd"); |
| 377 | static Timer tL("DevSparseChol, L-fact"); |
| 378 | static Timer tLT("DevSparseChol, LT-fact"); |
| 379 | static Timer tR("DevSparseChol, Reorder"); |
| 380 | static Timer tRA("DevSparseChol, Reorder Add"); |
| 381 | static Timer tD("DevSparseChol, Diag"); |
| 382 | |
| 383 | RegionTimer rt(t); |
| 384 | |
| 385 | UnifiedVectorWrapper ux(x); |
| 386 | UnifiedVectorWrapper uy(y); |
| 387 | ux.UpdateDevice(); |
| 388 | uy.UpdateDevice(); |
| 389 | if (synckernels) cudaDeviceSynchronize(); |
| 390 | |
| 391 | DevStackArray<double> mem_hx(x.Size()); |
| 392 | FlatVector<Dev<double>> hx(x.Size(), mem_hx.Data()); |
| 393 | |
| 394 | CudaRegionTimer rtR(tR); |
| 395 | DeviceSparseCholeskyReorderKernel<<<512,256>>> (ux.FVDev(), hx, order); |
| 396 | rtR.Stop(); |
| 397 | |
| 398 | // cout << "reordered[:10] = " << endl << D2H(hx.Range(10)) << endl; |
| 399 | |
| 400 | // Array<Dev<int>> incomingdep(host_incomingdep); |
| 401 | // Array<Dev<int>> incomingdep_trans(host_incomingdep_trans); |
| 402 | |
| 403 | DevStackArray<int> incomingdep(host_incomingdep.Size()); |
| 404 | H2D(incomingdep, host_incomingdep); |
| 405 | DevStackArray<int> incomingdep_trans(host_incomingdep_trans.Size()); |
| 406 | H2D(incomingdep_trans, host_incomingdep_trans); |
| 407 | |
| 408 | |
| 409 | CudaRegionTimer rtL(tL); |
| 410 | Dev<int> * pcnt = Dev<int>::Malloc(2); |
| 411 | pcnt->H2D(0); |
| 412 | (pcnt+1)->H2D(0); |
| 413 | |
| 414 | DeviceSparseCholeskySolveLKernel<<<512,dim3(32,8)>>> |
| 415 | ( |
| 416 | micro_dependency, incomingdep, hx, *(int*)pcnt, |
| 417 | microtasks, blocks, rowindex2, firstinrow_ri, firstinrow, lfact |
| 418 | ); |
| 419 | rtL.Stop(); |
| 420 | |
| 421 | // cout << "SolveL[:10] = " << endl << D2H(hx.Range(10)) << endl; |
| 422 | |
| 423 | // TODO : Diag |
| 424 | { |
| 425 | CudaRegionTimer rtD(tD); |
| 426 | DeviceSparseCholeskyMultDiagKernel<<<512,256>>> (diag, hx); |
| 427 | } |
| 428 | |
| 429 | // cout << "SolveDL[:10] = " << endl << D2H(hx.Range(10)) << endl; |