[Openmp-commits] [PATCH] D81054: [OpenMP] Introduce target memory manager

[Shilei Tian via Phabricator via Openmp-commits]

tianshilei1992 added a comment.

Thank you Jon for the review! The comments are really precious.

> Memory allocators are notoriously difficult to implement efficiently in the face of unknown workload. Can you share benchmarks that lead to this design?

The benchmark is quite simple:

  #pragma omp parallel for
  for (int i = 0; i < 4096; ++i) {
  #pragma omp target team distribute map(...)
    { /* kernel here */ }
  }

We have 4096 tasks, and depending on the number of threads `N`, we have `N` target region offloaded almost at the same time such that `M` of them might be executing simultaneously. For each kernel, the RT will allocate memory for it before its execution and free the memory after the execution. From NVVP, we observed that `cuMemFree` is very expensive, especially the computation is light but depends on a large amount of memory. What's more, during the `cuMemFree`, there is no context switch on the device, even though there are actually multiple kernels executing at the same time. From the profiling result of IBM XL OpenMP, we found that they don't call `cuMemFree` after each execution, and that is why we're thinking to have a memory pool.

> Similarly they're really easy to get wrong. There's a lot of subtle arithmetic in this patch. It would be prudent to cover this with tests, e.g. stub out the RTL->data_alloc calls for malloc/free and run targeted and fuzz tests under valgrind.

That sounds reasonable. Will do it.

> A second part of this puzzle is device side memory pools, so that malloc/free from a target region doesn't (always) have to call into the host. That may end up being quite platform dependent. That seems orthogonal to this patch.

That part should be covered in the plugin which is currently not the focus of this patch. But maybe we could avoid doing memory allocation and free in plugin I guess.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D81054/new/

https://reviews.llvm.org/D81054