QMCPACK / miniqmc

QMCPACK miniapp: a simplified real space QMC code for algorithm development, performance portability testing, and computer science experiments
Other
26 stars 34 forks source link

Explanation of OMP Code in `einspline_spo_omp.cpp` #266

Closed cabreraam closed 3 years ago

cabreraam commented 3 years ago

Admittedly, I am new to OpenMP GPU offload (and I am even further from understanding the mathematical methods used to calculate the total energy of the system).

When looking at the output of nvsys and nvprof when running miniqmc, I see that most of the offloaded GPU functionality is contained in einspline_spo_omp.cpp. My first question is, in the set function starting on line 88, the offload section references the array einsplines_ptr. I see that in line 119, this function is just a pointer. However, it is not specified in any of the #pragma omp target enter data map calls. I would believe that there is some kind of implicit assignment to einsplines_ptr, but what would it be? Maybe map(tofrom : einsplines_ptr[0:1])? Additionally, does this means that einsplines_ptr gets allocated into device memory? If so, how long does it persist? Until the entire program is finished? In evaluate_vgh, I see that einsplines[i] gets assigned to spline_m which gets used in a target offload section, yet it is not mapped. Has it already been mapped in the set or copy constructor functions?

Any help is appreciated!

Anthony

cabreraam commented 3 years ago

A colleague, @jdenny-ornl, was able to help me track down what I was interested in.

einsplines is a member variable of einspline_spo_omp, which is declared with OMPallocator--defined in OMPallocator.hpp--defines allocate and deallocate functions that omp target enter data map and omp target exit data map an allocated region, which in this case is einsplines. So, einsplines is not implicitly mapped, but rather explicitly mapped. I verified this by including this line of code:

std::cout << "einsplines.data() mapped already? " << omp_target_is_present(einsplines.data(),0) << std::endl;

before line 117 in einspline_spo_omp.cpp, and verified that einsplines.data() (and subsequently einsplines_ptr) has already been mapped, so the reference count to einsplines.data() is unaffected in the set function.