Improve gpuistl using cudaGraphs #5852
Conversation
multitalentloes
force-pushed
the
cudagraph_experiment
branch
from
d990b2e to
47726cc
Compare
| template void computeDiluDiagonal<T, blocksize>( \ | ||
| T*, int*, int*, int*, int*, const int, int, T*, int, cudaStream_t); \ | ||
| template void computeDiluDiagonalSplit<blocksize, T, double, MatrixStorageMPScheme::DOUBLE_DIAG_DOUBLE_OFFDIAG>( \ | ||
| const T*, \ |
There was a problem hiding this comment.
clang-format didnt make this bit prettier...
|
Speedup is measured to be 10% to 20% on the apply of both ILU0 and DILU. No speedup measured for the update so I am not introducing cudaGraphs there, no idea why we do not see the same runtime reduction when the update also has lots of short kernels... Still not mergable as the results are not verified on a consumer grade AMD-card to ensure there is no slowdown there (speedup is not expected with current rocm versions and current hardware generation) |
|
Now tested on AMD - 1.02 speedup on the ILU0 apply, though I am not sure this is significant/consistently reproducible. |
is also supported in HIP, though not does not seem to affect performance in any clear way. 1.1 to 1.2 speedup in Nvidia GPUs.
multitalentloes
force-pushed
the
cudagraph_experiment
branch
from
acd2c7a to
8ea1b29
Compare
| @@ -174,6 +174,19 @@ GpuBuffer<T>::copyFromHost(const T* dataPointer, size_t numberOfElements) | |||
| OPM_GPU_SAFE_CALL(cudaMemcpy(data(), dataPointer, numberOfElements * sizeof(T), cudaMemcpyHostToDevice)); | |||
| } | |||
|
|
|||
| template <class T> | |||
| void | |||
| GpuBuffer<T>::copyFromHost(const T* dataPointer, size_t numberOfElements, cudaStream_t stream) | |||
There was a problem hiding this comment.
various functions have been changed to include the cudastream, this is because to record a set of GPU activities in a cudaGraph, a created stream must be used.
Using cudagraphs reduces the overhead associated with many consecutive kernel launches in GPU ILU and DILU.
Code changes contain very few linechanges in the apply, though the stream has to be specified and many functions need an updated signature.
Speedups on spe1, spe11 and sleipner are typically around 1.1 to 1.2 in the preconditioners apply.
Also no speedup yet in the update, seems strange given the same kernel pattern...