The purpose of __tgt_offload_entry and how compiler generates the content for the offloading entry table and omp_offloading_entries section

Runtimes OpenMP
1

Hi Friends, hope you all are cool and safe.

We are experimenting with the libomptarget library to directly use with source code. I read https://reviews.llvm.org/D64943 about how to create a fat binary using clang-offload-wrapper utility. Given a CUDA program, I am able to create a source file for what the clang-offload-wrapper can produce. I however am not sure what to initialize the __start_omp_offloading_entries and __stop_omp_offloading_entries symbols. They are implicitly created by the linker when linking and creating the omp_offloading_entries section. What I am not sure are what is the purpose of the offload entry table and what could be put in the entry table by the compiler/linker (the content of the omp_offloading_entries section)? From what I read of the libomptarget design document, it seems they are used to store global variables and functions that could be mapped to devices and they are used at the runtime to help the mapping. How far of my understanding is to what it really is?

For example, for the following axpy omp offloading, what could be put in the offloading entry table? I built the program and did an objdump to see the content of the section, but not much clue what they are.

void axpy(int N, float *Y, float *X, float a) {
int i,j;
#pragma omp target map(to:X[0:N], N) map(tofrom:Y[0:N])
#pragma omp parallel for
for (i = 0; i < N; ++i) {
Y[i] += a * X[i];
}
}

Thank you
Yonghong

2

The table is used for looking up the corresponding device function or variable for host function or variable

When you call a kernel, you are actually using the host address of the entry point. The runtime uses the table to look up the corresponding address in the device image to make the call.

The lookup is done positional entry in the table, so the table on the host must match the table on the device.

Ravi