Skip to content

Target Code Generation

Attention

This page is for internal and development use only.

The general process to the target code generation is as follows:

  1. The code's AST is processed with RecursiveASTVisitors (clang/tools/sotoc/Visitors.{h,cpp})
  2. The found target code is recorded in TargetCodeFragment (clang/tools/sotoc/TargetCode.{h,pp} and clang/tools/sotoc/TargetCodeFragment.{h,cpp})
  3. The fragments are then 'reconstituted' into valid target code.

This transforms C code like 

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
int main(void) {
  int x = 0;
  int y = 1;
  int z = 2;

  #pragma omp target map(tofrom:x)
  {
    x = x + y + z;
  }
}
into 
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
void __omp_offloading_38_272a6c6b_main_l7(int *__sotoc_var_x, int y, int z)
{
   int x = *__sotoc_var_x;

  {
    x = x + y + z;
  }

  *__sotoc_var_x = x;
}

The target code is thereby transformed into a function with the mapped variables as arguments.

Extracting Target Code from the AST

The Clang AST of a source file can be viewed using the following command:

clang -Xclang -ast-dump -fsyntax-only -fopenmp code.c

Info

Without the -fopenmp flag Clang ignores #pragma omp.

Target regions are separate statements in the AST and functions and variables declared with #pragma omp declare target get a special attribute added to their AST node. The FindTargetCodeVisitor in Visitors.{h,cpp} is used for that.

Warning

This does not apply the type declarations and there is no guaranty that all functions used in the target region are declared in #pragma omp declare target.

To solve the issue we search all target regions, declared functions and variables to additional types, functions, and variables. This is done by the DiscoverTypesInDeclVisitor and DiscoverFunctionsInDeclVisitor. As types can be derived from other types (e.g., structs or typedef-chains) and functions can depend on other functions and additional types, those dependencies have to be resolved using the DeclResolver (in clang/tools/sotoc/DeclResolver.{h,cpp}).

Analysing and Recording the Target Code

Before target code can be generated, additional information about variable mappings have to be collected.

  • Private variables are not mapped. For those we have to create a declaration in the target region using the FindPrivateVariablesVisitor.

  • First-private variables (which is the default) are passed by value into the target region.

  • Variables mapped using from or tofrom, are passed by pointer. For ncc to process them, we bring the variable into scope by copying them into a local variable and afterwards back.

  • Arrays are mapped as pointer. If only an array slice (e.g., map(to:array[10:])) is mapped, we have to detect the bounds with the FindArraySectionVisitor and move the returned pointer (array[10]) in the target region. Multi-dimensional arrays have to be cast, so they can be accessed using the array[x] notation.

To recognize the variable form we use the TargetRegionVariable class in clang/tools/sotoc/TargetRegionVariable.{h,cpp}.

In addition, we have to rewrite the OpenMP pragma directive and add the appropriate clauses (OmpPragma.{h,cpp}). For example if a #pragma omp target parallel is encountered, an additional parallel region has to be added. Some clauses are not supported by the reduced #pragma. Those are filtered out or replaced. Parameters to clauses are also transferred as function arguments of the target function.

Code Generation

An Object of the TargetCode class, serves as a collection of TargetCodeFragments and generates the code. When adding code fragments, TargetCode sorts them according to their position in the original source file. We try to generate as little code as possible ourselves and let Clangs PrettyPrinter do most of the work, like functions and global variables.

Bug

Clangs PrettyPrinter is designed for compiler diagnostics and can not do e.g., anonymous structs and enums.

Note

We extended the PrettyPrinter to handle Decls.

We still have to generate functions for regions with mapped variables and clause parameters as arguments and local declarations to bring variables into scope as mentioned above.

Bug

Mappings with missing elements in the middle are currently not completely mapped by libomptarget/the OpenMP runtime.

Last update: 2021-11-24
Back to top