| 450 | |
| 451 | |
| 452 | void DevConstantElementByElementMatrix :: |
| 453 | MultAdd (double s, const BaseVector & x, BaseVector & y) const |
| 454 | { |
| 455 | static Timer t("DevConstantEBEMatrix::MultAdd"); RegionTimer reg(t); |
| 456 | static Timer tmult("DevConstantEBEMatrix::MultAdd - mult"); |
| 457 | static Timer tcopyin("DevConstantEBEMatrix::MultAdd - copyin"); |
| 458 | static Timer tcopyout("DevConstantEBEMatrix::MultAdd - copyout"); |
| 459 | |
| 460 | UnifiedVectorWrapper ux(x); |
| 461 | UnifiedVectorWrapper uy(y); |
| 462 | |
| 463 | ux.UpdateDevice(); |
| 464 | uy.UpdateDevice(); |
| 465 | |
| 466 | if (synckernels) cudaDeviceSynchronize(); |
| 467 | |
| 468 | // if (disjoint_cols) |
| 469 | if (true) |
| 470 | { |
| 471 | DevStackArray<double> dev_hx(numblocks*devmat.Width()); |
| 472 | DevStackArray<double> dev_hy(numblocks*devmat.Height()); |
| 473 | |
| 474 | tcopyin.Start(); |
| 475 | // ConstEBEKernelCopyIn (numblocks, devmat.Width(), rowdnums.DevData(), ux.DevData(), dev_hx.DevData()); |
| 476 | DeviceParallelFor |
| 477 | (numblocks*devmat.Width(), |
| 478 | [locx=dev_hx.DevData(), globx=ux.DevData(), idx=rowdnums.DevData()] DEVICE_LAMBDA (auto tid) |
| 479 | { |
| 480 | locx[tid] = globx[idx[tid]]; |
| 481 | }); |
| 482 | if (synckernels) cudaDeviceSynchronize(); |
| 483 | tcopyin.Stop(); |
| 484 | |
| 485 | // dev_hy = dev_hx * Trans(mat) |
| 486 | tmult.Start(); |
| 487 | FlatMatrix<Dev<double>> matx(numblocks, devmat.Width(), dev_hx.Data()); |
| 488 | FlatMatrix<Dev<double>> maty(numblocks, devmat.Height(), dev_hy.Data()); |
| 489 | // MultMatMat (matx, Trans(devmat), maty, s, 0); |
| 490 | maty = s * matx * Trans(devmat); |
| 491 | if (synckernels) cudaDeviceSynchronize(); |
| 492 | tmult.Stop(); |
| 493 | |
| 494 | tcopyout.Start(); |
| 495 | // ConstEBEKernelCopyOut (numblocks, devmat.Height(), coldnums.DevData(), dev_hy.DevData(), uy.DevData()); |
| 496 | DeviceParallelFor |
| 497 | (numblocks*devmat.Height(), |
| 498 | [globy=uy.DevData(), locy=dev_hy.DevData(), idx=coldnums.DevData() ] DEVICE_LAMBDA (auto tid) |
| 499 | { |
| 500 | atomicAdd((double*)globy+idx[tid], locy[tid]); |
| 501 | }); |
| 502 | |
| 503 | if (synckernels) cudaDeviceSynchronize(); |
| 504 | tcopyout.Stop(); |
| 505 | } |
| 506 | else |
| 507 | { |
| 508 | for (auto c : col_coloring) |
| 509 | { |
nothing calls this directly
no test coverage detected