REX declare target

rex_declare_target.md

declare target creates two copies of declared symbols. Given the following code:

#include <stdio.h>

#pragma omp declare target
int x = 100;
#pragma omp end declare target

int main(int argc, char **argv) {

  printf("On host: x = %d\n", x); // output: x = 100

#pragma omp target
  {
    x = 200;
    printf("On device: x = %d\n", x); // output: x = 200
  }

  printf("On host: x = %d\n", x); // output: x = 100
  x = 300;

#pragma omp target
  { printf("On device: x = %d\n", x); } // output: x = 200

  return 0;
}

Transformation in REX:

During the analysis stage or the very beginning of lowering stage, after the REX AST has been created, REX collects all the omp declare target nodes. It takes each node as a starting point and gathers all subsequent sibling nodes in a list, excluding child nodes, until an omp end declare target node is encountered. All these sibling nodes are stored as pointers. They are then copied to a new file for GPU outlined functions and set as __device__ symbols.