llvm dev - Jun 2020 - Is inlining of functions containing OpenMP offload regions possible?

Dear all,

we are trying OpenMP offloading functionality in LLVM for a target 
architecture we want to implement. To this end, we tried to compile a 
code segment such as the following one:

__attribute__((always_inline))
inline void function (int const range)
{
#pragma omp target
   {
     for (int i = 0; i < range; i++)
     {
	...device code
     }
   }
}

We expected that when calling this function, e.g. from main(), in the style

function(3);

that the function would be inlined and accordingly, the constant "3" 
would be known to the compiler as compile-time-constant, so loop 
unrolling would be possible on the device. However, as far as we can see 
(looking at the bitcode file that is passed to the device toolchain), 
this does not happen, and the "range" parameter stays a variable and
not
a constant in the created code.

So my question is: Is this type of inlining supported, i.e., in the case 
when the function to be inlined contains OpenMP offload regions?

I would appreciate an answer a lot!

Sincerely,

Kai Plociennik


-- 
Dr. Kai Plociennik
Fraunhofer-Institut für Techno- und Wirtschaftsmathematik ITWM
Competence Center High Performance Computing
Fraunhofer-Platz 1
67663 Kaiserslautern
Tel: +49 (0)631 31600 0
mail: kai.plociennik at itwm.fraunhofer.de
www.itwm.fraunhofer.de

Currently it is not supported as by default the outlining happens in the
frontend which is unaware of context.

Best regards,
Alexey Bataev
> 4 июня 2020 г., в 08:37, Kai Plociennik via llvm-dev <llvm-dev at
lists.llvm.org> написал(а):
> 
> Dear all,
> 
> we are trying OpenMP offloading functionality in LLVM for a target
architecture we want to implement. To this end, we tried to compile a code
segment such as the following one:
> 
> __attribute__((always_inline))
> inline void function (int const range)
> {
> #pragma omp target
>  {
>    for (int i = 0; i < range; i++)
>    {
>    ...device code
>    }
>  }
> }
> 
> We expected that when calling this function, e.g. from main(), in the style
> 
> function(3);
> 
> that the function would be inlined and accordingly, the constant
"3" would be known to the compiler as compile-time-constant, so loop
unrolling would be possible on the device. However, as far as we can see
(looking at the bitcode file that is passed to the device toolchain), this does
not happen, and the "range" parameter stays a variable and not a
constant in the created code.
> 
> So my question is: Is this type of inlining supported, i.e., in the case
when the function to be inlined contains OpenMP offload regions?
> 
> I would appreciate an answer a lot!
> 
> Sincerely,
> 
> Kai Plociennik
> 
> 
> -- 
> Dr. Kai Plociennik
> Fraunhofer-Institut für Techno- und Wirtschaftsmathematik ITWM
> Competence Center High Performance Computing
> Fraunhofer-Platz 1
> 67663 Kaiserslautern
> Tel: +49 (0)631 31600 0
> mail: kai.plociennik at itwm.fraunhofer.de
> www.itwm.fraunhofer.de
>