<html><head><meta http-equiv="Content-Type" content="text/html; charset=us-ascii"></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><br class=""><div><br class=""><blockquote type="cite" class=""><div class="">On Feb 3, 2022, at 12:11 AM, Rohan Yadav <<a href="mailto:rohany@alumni.cmu.edu" class="">rohany@alumni.cmu.edu</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div dir="ltr" class="">Hi All,<div class=""><br class=""></div><div class="">I'm trying to use the MatMatMult function with 1 sparse matrix B and two dense matrices A, C. I'm computing A = B * C.</div><div class=""><br class=""></div><div class="">My code is below:<br class=""><br class="">```<br class=""><pre style="white-space: pre-wrap;" class="">void spmm(Mat B, int warmup, int niter) {
  Mat A, C;
  PetscInt i, j = 32, k;
  MatGetSize(B, &i, &k);
  MatCreateDenseCUDA(PETSC_COMM_WORLD, PETSC_DECIDE, PETSC_DECIDE, i, j, NULL, &A);
  MatCreateDenseCUDA(PETSC_COMM_WORLD, PETSC_DECIDE, PETSC_DECIDE, k, j, NULL, &C);

  // Initialize entries in the output.
  MatZeroEntries(A);
  setMatToConstant(C, 1.0);
  
  // Finally, do the computation.
  auto avgTime = benchmarkWithWarmup(warmup, niter, [&]() {
    MatMatMult(B, C, MAT_REUSE_MATRIX, PETSC_DEFAULT, &A);
  });
  PetscPrintf(PETSC_COMM_WORLD, "Average time: %lf ms.\n", avgTime * 1000);
}</pre><pre style="white-space: pre-wrap;" class="">```</pre><pre style="overflow-wrap: break-word;" class=""><font face="Arial, Helvetica, sans-serif" class=""><span style="white-space:normal" class="">where benchmarkWithWarmup is a simple wrapper function that runs a lambda several times.</span></font></pre><pre style="overflow-wrap: break-word;" class=""><font face="Arial, Helvetica, sans-serif" class=""><span style="white-space:normal" class=""><br class=""></span></font></pre><pre style="overflow-wrap: break-word;" class=""><font face="Arial, Helvetica, sans-serif" class=""><span style="white-space:normal" class="">I'm running this function with arguments `-vec_type cuda -mat_type aijcusparse`, </span></font></pre></div></div></div></blockquote><div>   These arguments are not appropriate; they are only for certain examples, you shouldn't rely on them.</div>  <br class=""><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class=""><pre style="overflow-wrap: break-word;" class=""><font face="Arial, Helvetica, sans-serif" class=""><span style="white-space:normal" class="">and see that the performance is relatively slow. I'm wondering if I'm using the API incorrectly, or the computation is executing as expected. `nvprof` shows that much of the time is spent in a device to host memcpys:</span></font></pre></div></div></div></blockquote><div><br class=""></div>   Please send the code that builds the sparse B matrix and the <span style="white-space: pre-wrap;" class="">setMatToConstant() routine.</span></div><div><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div><br class=""></div><div><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class=""><pre style="overflow-wrap: break-word;" class=""><font face="Arial, Helvetica, sans-serif" class=""><span style="white-space:normal" class="">```<br class=""></span></font>            Type  Time(%)      Time     Calls       Avg       Min       Max  Name<br class=""> GPU activities:   87.32%  11.9978s        33  363.57ms  1.5040us  388.26ms  [CUDA memcpy DtoH]<br class="">                    8.71%  1.19611s        30  39.870ms  37.421ms  39.976ms  void cusparse::csrmm_kernel<cusparse::CsrMMPolicy<unsigned int=128, bool=0, bool=0, unsigned int=8, unsigned int=16, unsigned int=4, unsigned int=0>, int, double, double, double>(bool=0, cusparse::csrmm_kernel<cusparse::CsrMMPolicy<unsigned int=128, bool=0, bool=0, unsigned int=8, unsigned int=16, unsigned int=4, unsigned int=0>, int, double, double, double>, cusparse::csrmm_kernel<cusparse::CsrMMPolicy<unsigned int=128, bool=0, bool=0, unsigned int=8, unsigned int=16, unsigned int=4, unsigned int=0>, int, double, double, double>, bool=0 const *, bool=0 const , bool=0, bool=0, int, cusparseOperation_t, cusparse::csrmm_kernel<cusparse::CsrMMPolicy<unsigned int=128, bool=0, bool=0, unsigned int=8, unsigned int=16, unsigned int=4, unsigned int=0>, int, double, double, double> const *, cusparse::csrmm_kernel<cusparse::CsrMMPolicy<unsigned int=128, bool=0, bool=0, unsigned int=8, unsigned int=16, unsigned int=4, unsigned int=0>, int, double, double, double> const , unsigned int=8 const *, unsigned int=8 const , cusparse::csrmm_kernel<cusparse::CsrMMPolicy<unsigned int=128, bool=0, bool=0, unsigned int=8, unsigned int=16, unsigned int=4, unsigned int=0>, int, double, double, double>, unsigned int=16*, cusparse::csrmm_kernel<cusparse::CsrMMPolicy<unsigned int=128, bool=0, bool=0, unsigned int=8, unsigned int=16, unsigned int=4, unsigned int=0>, int, double, double, double>)<br class="">                    3.87%  531.56ms        14  37.968ms  1.0240us  227.29ms  [CUDA memcpy HtoD]<br class="">                    0.07%  9.7452ms         6  1.6242ms  1.0880us  3.2481ms  [CUDA memset]<br class="">                    0.02%  2.8727ms         1  2.8727ms  2.8727ms  2.8727ms  void thrust::cuda_cub::core::_kernel_agent<thrust::cuda_cub::__parallel_for::ParallelForAgent<thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<double>, double>, unsigned long>, thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<double>, double>, unsigned long>(thrust::device_ptr<double>, double)<br class="">                    0.01%  1.4953ms         2  747.67us  56.188us  1.4392ms  void thrust::cuda_cub::core::_kernel_agent<thrust::cuda_cub::__parallel_for::ParallelForAgent<thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<int>, int>, unsigned long>, thrust::cuda_cub::__uninitialized_fill::functor<thrust::device_ptr<int>, int>, unsigned long>(thrust::device_ptr<int>, int)
<font face="Arial, Helvetica, sans-serif" class=""><span style="white-space:normal" class="">```</span></font></pre><pre style="overflow-wrap: break-word;" class="">The logview output is:</pre><pre style="overflow-wrap: break-word;" class="">```<span style="white-space: pre-wrap;" class="">-----------------------------------------------------------------------------------------------------------------------</span>
</pre></div></div></div></blockquote><div><span style="white-space: pre-wrap; font-style: normal; font-size: 12px;" class=""><font face="Monaco" class="">Event                Count      Time (sec)     Flop                              --- Global ---  --- Stage ----  Total   GPU    - CpuToGpu -   - GpuToCpu - GPU
                   Max Ratio  Max     Ratio   Max  Ratio  Mess   AvgLen  Reduct  %T %F %M %L %R  %T %F %M %L %R Mflop/s Mflop/s Count   Size   Count   Size  %F
---------------------------------------------------------------------------------------------------------------------------------------------------------------</font></span></div><div><span style="white-space: pre-wrap; font-style: normal; font-size: 12px;" class=""><font face="Monaco" class="">MatMatMultSym         60 1.0 9.2215e-01 2.6 0.00e+00 0.0 4.0e+00 7.3e+05 3.2e+01  1  0  6  0 40   1  0  6  0 51     0       0      0 0.00e+00    0 0.00e+00  0
MatMatMultNum         30 1.0 4.2967e+01 1.0 6.34e+11 1.1 6.0e+01 9.4e+07 0.0e+00 37100 86110  0  37100 86110  0 28598  920026      2 6.71e+03   30 8.73e+04 98
MatCUSPARSCopyTo       1 1.0 4.4761e-01 1.0 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00  0  0  0  0  0   0  0  0  0  0     0       0      1 3.80e+03    0 0.00e+00  0
MatDenseCopyTo         1 1.0 2.2742e-01 1.0 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00  0  0  0  0  0   0  0  0  0  0     0       0      1 2.91e+03    0 0.00e+00  0
MatDenseCopyFrom      31 1.0 1.2006e+01 1.0 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00 10  0  0  0  0  10  0  0  0  0     0       0      0 0.00e+00   31 9.02e+04  0</font></span><br class=""></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">From the third line we see the sparse matrix is copied to the GPU once, this is good.</span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">From line 4 a dense matrix is copied to the GPU once, this is good. </span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">But from line 5 we see a dense matrix is copied from the GPU to the CPU 31 times! Looking at line 2 we see 30 copies from GPU to the CPU. </span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">The flop rate on the GPU is </span><span style="font-family: Monaco; font-size: 12px; white-space: pre-wrap;" class="">920,026 </span><span style="white-space: pre-wrap;" class="">  which is fine, but the flop rate for the entire multiply time is a terrible </span><span style="font-family: Monaco; font-size: 12px; white-space: pre-wrap;" class="">28,598,</span><span style="white-space: pre-wrap;" class=""> this is because this time includes all the copies between the GPU and CPU and CPU and GPU. </span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">So let's see if we can figure out why all these copies are taking place from the GPU to the CPU.</span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">But first please verify that if you run with one MPI rank the "on GPU" and the overall flop rates for the MatMatMult() are almost the same and there is no copy from the GPU for each multiply?</span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">I think the parallel multiply is done with MatMatMultNumeric_MPIAIJ_MPIDense(). This code has two problems</span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">1) It</span><span style="white-space: pre-wrap;" class=""> uses MatMPIDenseScatter() to move to the other ranks their needed rows of the C matrix. That function has the call MatDenseGetArrayRead() normally would trigger a copy of C up to the CPU each time. But since C is not changing in your test run I guess it only triggers one copy.</span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">2) If uses MatMatMultNumericAdd_SeqAIJ_SeqDense(aij->B,workB,cdense->A,PETSC_TRUE);CHKERRQ(ierr); to do the off diagonal part of the product but this triggers for each multiply a copy of the result matrix from the CPU to the GPU (hugely expensive)</span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">For performance there needs to be a new routine </span><span style="white-space: pre-wrap;" class="">MatMatMultNumeric_MPIAIJCUSPRSE_MPICUDADense() that is smarter about the needed MPI communication so it only moves exactly what it needs to the other ranks and it does the off-diagonal part of the product on the GPU so it does not need to copy the result up to the CPU. </span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class="">Barry</span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><div class=""><span style="white-space: pre-wrap;" class=""><br class=""></span></div><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class=""><pre style="overflow-wrap: break-word;" class="">---------------------------------------------- PETSc Performance Summary: ----------------------------------------------<br class=""><br class="">/g/g15/yadav2/taco/petsc/bin/benchmark on a  named lassen457 with 2 processors, by yadav2 Wed Feb  2 17:23:19 2022<br class="">Using Petsc Release Version 3.16.3, unknown<br class=""><br class="">                         Max       Max/Min     Avg       Total<br class="">Time (sec):           1.163e+02     1.000   1.163e+02<br class="">Objects:              4.800e+01     1.000   4.800e+01<br class="">Flop:                 6.338e+11     1.065   6.144e+11  1.229e+12<br class="">Flop/sec:             5.451e+09     1.065   5.284e+09  1.057e+10<br class="">MPI Messages:         3.500e+01     1.000   3.500e+01  7.000e+01<br class="">MPI Message Lengths:  2.544e+09     1.000   7.267e+07  5.087e+09<br class="">MPI Reductions:       8.100e+01     1.000<br class=""><br class="">Flop counting convention: 1 flop = 1 real number operation of type (multiply/divide/add/subtract)<br class="">                            e.g., VecAXPY() for real vectors of length N --> 2N flop<br class="">                            and VecAXPY() for complex vectors of length N --> 8N flop<br class=""><br class="">Summary of Stages:   ----- Time ------  ----- Flop ------  --- Messages ---  -- Message Lengths --  -- Reductions --<br class="">                        Avg     %Total     Avg     %Total    Count   %Total     Avg         %Total    Count   %Total<br class=""> 0:      Main Stage: 1.1628e+02 100.0%  1.2288e+12 100.0%  7.000e+01 100.0%  7.267e+07      100.0%  6.300e+01  77.8%<br class=""><br class="">------------------------------------------------------------------------------------------------------------------------<br class="">See the 'Profiling' chapter of the users' manual for details on interpreting output.<br class="">Phase summary info:<br class="">   Count: number of times phase was executed<br class="">   Time and Flop: Max - maximum over all processors<br class="">                  Ratio - ratio of maximum to minimum over all processors<br class="">   Mess: number of messages sent<br class="">   AvgLen: average message length (bytes)<br class="">   Reduct: number of global reductions<br class="">   Global: entire computation<br class="">   Stage: stages of a computation. Set stages with PetscLogStagePush() and PetscLogStagePop().<br class="">      %T - percent time in this phase         %F - percent flop in this phase<br class="">      %M - percent messages in this phase     %L - percent message lengths in this phase<br class="">      %R - percent reductions in this phase<br class="">   Total Mflop/s: 10e-6 * (sum of flop over all processors)/(max time over all processors)<br class="">   GPU Mflop/s: 10e-6 * (sum of flop on GPU over all processors)/(max GPU time over all processors)<br class="">   CpuToGpu Count: total number of CPU to GPU copies per processor<br class="">   CpuToGpu Size (Mbytes): 10e-6 * (total size of CPU to GPU copies per processor)<br class="">   GpuToCpu Count: total number of GPU to CPU copies per processor<br class="">   GpuToCpu Size (Mbytes): 10e-6 * (total size of GPU to CPU copies per processor)<br class="">   GPU %F: percent flops on GPU in this event<br class="">------------------------------------------------------------------------------------------------------------------------<br class="">Event                Count      Time (sec)     Flop                              --- Global ---  --- Stage ----  Total   GPU    - CpuToGpu -   - GpuToCpu - GPU<br class="">                   Max Ratio  Max     Ratio   Max  Ratio  Mess   AvgLen  Reduct  %T %F %M %L %R  %T %F %M %L %R Mflop/s Mflop/s Count   Size   Count   Size  %F<br class="">---------------------------------------------------------------------------------------------------------------------------------------------------------------<br class=""><br class="">--- Event Stage 0: Main Stage<br class=""><br class="">BuildTwoSided          2 1.0 4.4400e-01567.5 0.00e+00 0.0 2.0e+00 4.0e+00 2.0e+00  0  0  3  0  2   0  0  3  0  3     0       0      0 0.00e+00    0 0.00e+00  0<br class="">BuildTwoSidedF         1 1.0 4.4395e-0115659.1 0.00e+00 0.0 0.0e+00 0.0e+00 1.0e+00  0  0  0  0  1   0  0  0  0  2     0       0      0 0.00e+00    0 0.00e+00  0<br class="">MatAssemblyBegin      32 1.0 4.4400e-017378.9 0.00e+00 0.0 0.0e+00 0.0e+00 1.0e+00  0  0  0  0  1   0  0  0  0  2     0       0      0 0.00e+00    0 0.00e+00  0<br class="">MatAssemblyEnd        32 1.0 1.8511e+00 2.2 0.00e+00 0.0 0.0e+00 0.0e+00 6.0e+00  1  0  0  0  7   1  0  0  0 10     0       0      0 0.00e+00    0 0.00e+00  0<br class="">MatZeroEntries         1 1.0 3.3306e-03 1.0 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00  0  0  0  0  0   0  0  0  0  0     0       0      0 0.00e+00    0 0.00e+00  0<br class="">MatLoad                1 1.0 1.7220e+01 1.0 0.00e+00 0.0 6.0e+00 -8.8e+07 2.1e+01 15  0  9-10 26  15  0  9-10 33     0       0      0 0.00e+00    0 0.00e+00  0<br class="">MatMatMultSym         60 1.0 9.2215e-01 2.6 0.00e+00 0.0 4.0e+00 7.3e+05 3.2e+01  1  0  6  0 40   1  0  6  0 51     0       0      0 0.00e+00    0 0.00e+00  0<br class="">MatMatMultNum         30 1.0 4.2967e+01 1.0 6.34e+11 1.1 6.0e+01 9.4e+07 0.0e+00 37100 86110  0  37100 86110  0 28598   920026      2 6.71e+03   30 8.73e+04 98<br class="">MatCUSPARSCopyTo       1 1.0 4.4761e-01 1.0 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00  0  0  0  0  0   0  0  0  0  0     0       0      1 3.80e+03    0 0.00e+00  0<br class="">MatDenseCopyTo         1 1.0 2.2742e-01 1.0 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00  0  0  0  0  0   0  0  0  0  0     0       0      1 2.91e+03    0 0.00e+00  0<br class="">MatDenseCopyFrom      31 1.0 1.2006e+01 1.0 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00 10  0  0  0  0  10  0  0  0  0     0       0      0 0.00e+00   31 9.02e+04  0<br class="">VecSet                 3 1.0 4.1917e-04 1.1 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00  0  0  0  0  0   0  0  0  0  0     0       0      0 0.00e+00    0 0.00e+00  0<br class="">SFSetGraph             1 1.0 1.9180e-04 1.1 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00  0  0  0  0  0   0  0  0  0  0     0       0      0 0.00e+00    0 0.00e+00  0<br class="">SFSetUp                1 1.0 1.3672e-02 1.1 0.00e+00 0.0 4.0e+00 7.3e+05 1.0e+00  0  0  6  0  1   0  0  6  0  2     0       0      0 0.00e+00    0 0.00e+00  0<br class="">---------------------------------------------------------------------------------------------------------------------------------------------------------------<br class=""><br class="">Memory usage is given in bytes:<br class=""><br class="">Object Type          Creations   Destructions     Memory  Descendants' Mem.<br class="">Reports information only for process 0.<br class=""><br class="">--- Event Stage 0: Main Stage<br class=""><br class="">              Matrix    37             30   2867511840     0.<br class="">              Viewer     2              0            0     0.<br class="">              Vector     4              1         1792     0.<br class="">           Index Set     2              2      1495248     0.<br class="">   Star Forest Graph     3              0            0     0.<br class="">========================================================================================================================<br class="">Average time to get PetscTime(): 3.83e-08<br class="">Average time for MPI_Barrier(): 7.874e-07<br class="">Average time for zero size MPI_Send(): 3.4035e-06<br class="">#PETSc Option Table entries:<br class="">-bench spmm<br class="">-enable_gpu<br class="">-log_view<br class="">-mat_type aijcusparse<br class="">-matload_block_size 1<br class="">-matrix /p/gpfs1/yadav2/tensors/petsc/arabic-2005.petsc<br class="">-n 20<br class="">-vec_type cuda<br class="">-warmup 10
```</pre><pre style="overflow-wrap: break-word;" class=""><br class=""></pre><pre style="overflow-wrap: break-word;" class="">Thanks,</pre><pre style="overflow-wrap: break-word;" class=""><br class=""></pre><pre style="overflow-wrap: break-word;" class="">Rohan Yadav</pre><pre style="overflow-wrap: break-word;" class=""><font face="Arial, Helvetica, sans-serif" class=""><span style="white-space:normal" class=""><br class=""></span></font></pre></div></div>
</div></blockquote></div><br class=""></body></html>