Method Mult

ngscuda/cuda_linalg.cpp:384–448

  void DevConstantElementByElementMatrix ::
  Mult (const BaseVector & x, BaseVector & y) const
  {
    static Timer t("DevConstantEBEMatrix::Mult"); RegionTimer reg(t);
    static Timer tmult("DevConstantEBEMatrix::Mult - mult");
    static Timer tcopyin("DevConstantEBEMatrix::Mult - copyin");
    static Timer tcopyout("DevConstantEBEMatrix::Mult - copyout");
    
    UnifiedVectorWrapper ux(x);
    UnifiedVectorWrapper uy(y);

    ux.UpdateDevice();
    uy.UpdateDevice();

    if (synckernels) cudaDeviceSynchronize();


    if (!output_onto)
      {
        uy = 0.0;
        MultAdd (1, x, y);
        return;
      }

    
      {
        DevStackArray<double> dev_hx(numblocks*devmat.Width());
        DevStackArray<double> dev_hy(numblocks*devmat.Height());

        tcopyin.Start();
        // ConstEBEKernelCopyIn (numblocks, devmat.Width(), rowdnums.DevData(), ux.DevData(), dev_hx.DevData());
	DeviceParallelFor
          (numblocks*devmat.Width(),
           [locx=dev_hx.DevData(), globx=ux.DevData(), idx=rowdnums.DevData()] DEVICE_LAMBDA (auto tid)
           {
             locx[tid] = globx[idx[tid]];
           });
        if (synckernels) cudaDeviceSynchronize();
        tcopyin.Stop();
        
        // dev_hy = dev_hx * Trans(mat)
        tmult.Start();
        FlatMatrix<Dev<double>> matx(numblocks, devmat.Width(), dev_hx.Data());
        FlatMatrix<Dev<double>> maty(numblocks, devmat.Height(), dev_hy.Data());
        // MultMatMat (matx, Trans(devmat), maty, s, 0);
        maty = matx * Trans(devmat);
        if (synckernels) cudaDeviceSynchronize();
        tmult.Stop();
        
        tcopyout.Start();        
        // ConstEBEKernelCopyOut (numblocks, devmat.Height(), coldnums.DevData(), dev_hy.DevData(), uy.DevData());
        DeviceParallelFor
          (numblocks*devmat.Height(),
           [globy=uy.DevData(), locy=dev_hy.DevData(), idx=coldnums.DevData() ] DEVICE_LAMBDA (auto tid)
           {
             // atomicAdd((double*)globy+idx[tid], locy[tid]);
             globy[idx[tid]] = locy[tid];
           });