How do you map nested structures to target memory in OpenMP?

user_9035791 0
Viewed 86

I have been trying to offload an existing OpenMP parallel loop to a GPU target.

The problem is, the to-be-offloaded calculation function is using a vector of objects, and each of these objects has vectors of other objects as members.

Mapping a simple vector to GPU is quite easy:

std::vector<double> v;
// ...
double* vd = v.data();

#pragma omp target data
    map(to: vd[:v.size()])
{
    // ...
}

However, I'm at loss when it comes to mapping something like this:

// Example 1
struct Matrix {
    std::vector<double> matrix_data;
    std::size_t rows = 0
    std::size_t cols = 0;
}

std::vector<Matrix> mv;  // How to map "mv" to target?

Or, even more complicated:

// Example 2
struct S {
    std::vector<Matrix> nodes;
    std::map<int, Matrix> mm;
}

std::vector<S> sv;  // How to map "sv" to target?

Unfortunately, the main object that the existing program (and its parallelized loops) is operating on is of a type with many such nested object members. The calculation loops read/write to deeply nested variables all over the object.

Compiler: Intel oneAPI DPC++/C++ Compiler 2022.0.0, Windows and Linux. GPU: Intel gen9 (i9-9900K).

P.S. I have not been able to find a good guide on OpenMP targets. It seems that there are either very simple beginner tutorials focusing only on saxpy, or the dry specification itself.

If it's not possible with OpenMP, is there another API that can be easily used with such C++ structures (ideally, without conflicting with other OpenMP code in the program)? Perhaps by defining mapper functions for each type?

Thanks

target openmp mapping gpu c++
1 Answers

AFAIK, you cannot map any advanced C++ data structure to a target device in OpenMP, especially std::map. std::map is generally implemented as a red-black tree and OpenMP has no (efficient) way to track and move the structure on the target device. In fact, this data structure is pretty inefficient on GPU and you should really avoid using it since it is not SIMD-friendly. Instead, you can use lookup table if numbers are small or a hashmap. Hashmaps are still not SIMD-friendly but there are ways to implement this on GPUs (though is a hard to do efficiently and even more tricky only with OpenMP). OpenMP is designed to map basic arrays and basic structure on the target device. In fact, the documentation states:

If a list item in a map clause is a variable of structure type then it is treated as if each structure element contained in the variable is a list item in the clause.

OpenMP does not know how to map std::vector since it cannot know that the array in the implementation point to a memory location that is an array and that one of the field is the size. You have to do this complex mapping yourself. I think you can use the directive declare mapper to do such a complex mapping. AFAIK, the idea is to define a mapper for a type and then specify it during the map so to use the user-defined mapping operations. Keep in mind that GPUs are efficient to compute SIMD-friendly numerical operations on large regular contiguous data structures. Transfer request have a pretty high latency and doing many small transfer is not efficient. Thus, it may be better to flatten/compactify your data structure so to send it on the target devices.

user_129395570
Related