Use a new file to store the outlined function for the device (CPU or GPU) about rexompiler

[ouankou]

Solution 2:

For outlining in OpenMP, ROSE enables useParameterWrapper and useStructureWrapper to include the variables into a data structure and pass it to the outlined function.

Another flag useNewFile is to store the outlined function in a new file. However, ROSE doesn't support using useNewFile and useStructureWrapper at the same time. Otherwise, it will complain that Outliner::generateParameterStructureDeclaration() separated file case is not yet handled.

The source code that throws the error message is: https://github.com/passlab/rexompiler/blob/master/src/midend/programTransformation/astOutlining/Transform.cc#L140

It means solution 2 requires us to:

1. either explore the details of the error and implements any missing feature in ROSE,
2. or significantly change the workflow in omplowering to not using structure wrapper.

The OpenMP transformation for FORTRAN doesn't use useStructureWrapper. It's possible that we borrow the same workflow to C/C++. However, there are much code in ROSE specifically for FORTRAN in omplowering, outliner, and so on. To do that, we have to change many original modules. It's probably not a good idea.

[ouankou]

Solution 1:

So far, a new file with a given name could be created for the outlined function. However, it seems only the function SgNode itself is copied but not its full subtree.
The new file and the original one belong to the same SgProject. The unparsing will performed to a project so that the two files will be both unparsed.

Given a simple input:

#include <stdio.h>

int main (int argc, char** argv) {
      int i = 1;
        #pragma omp parallel if(i == 1)
      {
          #pragma omp single
            printf("Test\n");
      }
}

The regular transformed code is:

#include <stdio.h>
#include "rex_kmp.h"
static void OUT__1__5599__(int *__global_tid,int *__bound_tid,void *__out_argv);

int main(int argc,char **argv)
{
  int __global_tid_value = __kmpc_global_thread_num(0);
  int *__global_tid = &__global_tid_value;
  int status = 0;
  int i = 1;
  if (i == 1)
    __kmpc_fork_call(0,1,OUT__1__5599__,0);
   else {
    OUT__1__5599__(&( *__global_tid),0,0);
  }
}

static void OUT__1__5599__(int *__global_tid,int *__bound_tid,void *__out_argv)
{
  if (__kmpc_single(0, *__global_tid)) {
    printf("Test\n");
    __kmpc_end_single(0, *__global_tid);
  }
  __kmpc_barrier(0, *__global_tid);
}

The current implementation of solution 1 produces a separate file. The printf line is missing here.

static void OUT__1__5599__(int *__global_tid,int *__bound_tid,void *__out_argv)
{
  if (__kmpc_single(0, *__global_tid)) {
    __kmpc_end_single(0, *__global_tid);
  }
  __kmpc_barrier(0, *__global_tid);
}

There are several problems:

1. The separated file also needs to include rex_kmp.h.
2. The modifier static in the separated file should be removed.
3. The modifier static in the main file should be changed to extern.
4. The separated file should have the complete outlined function definition.
5. The outlined function definition in the main file should be removed. Its SgNode may or may not be kept, but it shouldn't be unparsed for sure.
6. This simple case doesn't have a parameter structure, but we expect the structure will encounter a similar issue with the outlined function.
7. Should we store all the outlined functions in one file and multiple files (one function per file)?
8. Other headers and functions that the outlined function depends on should also be added to the new file.

There are several ROSE builder APIs that could be used to find a forward declaration.

Work-in-progress output

rose-compiler -rose:openmp:lowering -rose:skipfinalCompileStep foo.c

The main file rose_foo.c:

#include <stdio.h>
#define N 1
#include "rex_kmp.h"
extern void OUT__1__5599__(int *__global_tid,int *__bound_tid,void *__out_argv);

int main(int argc,char **argv)
{
  int __global_tid_value = __kmpc_global_thread_num(0);
  int *__global_tid = &__global_tid_value;
  int status = 0;
  int i = 1;
  if (i == 1)
    __kmpc_fork_call(0,1,OUT__1__5599__,0);
   else {
    OUT__1__5599__(&( *__global_tid),0,0);
  }
}

The new file having the outlined function definition rose_foo_lib.c:

#include <stdio.h>
#define N 1
#include "rex_kmp.h"
int main(int ,char **);

void OUT__1__5599__(int *__global_tid,int *__bound_tid,void *__out_argv)
{
  if (__kmpc_single(0, *__global_tid)) {
    printf("Test\n");
    __kmpc_end_single(0, *__global_tid);
  }
  __kmpc_barrier(0, *__global_tid);
}

[yanyh15]

@ouankou Thank you for explore this in details. I could not find the code I created before, more than 10 years ago. I vaguely recall that it will need a deepcopy of AST in order to duplicate the AST. Did you try sageInterface::deepCopyNode (const SgNode *subtree). We can ask Liao for details.

I am thinking of two options for files vs regions: 1) a new file for each outlined function, or 2) one file for all the outlined function. It seems both are ok for me. Considering that we will support offloading (CUDA GPU) and vectorization (simd instruction or intrinsics), we can have one file for each architecture (CPU, CUDA, simd) that includes the outlined function for all the regions of a file, or one file per outlined region per architecture.

With either option, there are some details that need to be handled for the new file to be compilable, e.g. for the signatures of several things:

1. header include, we can copy the original include to the new file.
2. typedef, extern, #define can all be copied to the new file.
3. global variable needs to be added to the new file as extern-ed variable
4. all the function used by the outlined function body needs to have correct function signature in the new file
5. others (?). There might be other corner cases.

We can create a new header that include the above info and include this header in the file for the outline function, if we need to create several files for all the outline regions. This approach just simply copy all the signature definition of the original file. A better approach is just to extract the signature declaration of those symbols used in the outlined body, as well as the macro used.

[ouankou]

@yanyh15 Yes, I used the similar API deepCopy. I tried both and got the same full subtree. They are probably the same thing. Currently, the first 5 problems have been solved. I'm working on problem 7. For simple cases, it works fine. The function declarations are all copied to the new file and the two generated files can be compiled to get the correct executable.

However, I'm sure there will be issues for normal cases due to the points you listed. Additionally, since some ROSE APIs don't work in my code, there must be something wrong underneath. I'll wrap up the current implementation and push a commit. Then we can keep updating the design from there.

This is only an initial implementation to simply store the outlined function on CPU in a new file. In the meantime, I'll start to look at the target transformation to see how to create a CUDA kernel properly. If that works, we can take your hand-wrtitten AXPY offloading code and create a complete prototype in REX compiler.

[yanyh15]

This is a general problem: how to split a file into multiple files. It is not an easy job if we need to do perfect because of the situation we need to consider:
1.

[ouankou]

According to Dr. Liao, ROSE does keep all the definitions of global variables in the separate file having outlined functions. The compiler will complain that they are redeclared. We may need to manually scan those variables and mark them as extern declarations.

[ouankou]

Issues:

There could be multiple issues when we create a new file to hold the outlined functions.

1. Global variables (commit 3a7c6cf)
   The main file may have global variables. If they are static, they need to be copied. Otherwise, we have to set them to be extern without definition.
   For example, int a = 1; should be extern int a;, but static int b = 2; should be preserved as-is.

2. Preprocessing macros #define (commit 3a7c6cf)
   They should be copied exactly as-is because they are not shared across source files.

3. Function declaration
   Functions other than the outlined ones should only have the extern declaration in the new file if they are not static. Otherwise, they should be copied with the definition.
   The outline functions should have the forward declaration in the new file in case they are calling each other.
   In the main file, the outline functions only have the extern declaration there.

Solutions:

1. Generate a new header:
   We may need to analyze the source code and move the portion required by the outline functions into a new header. Then both the main file and new file include this header. The drawback of this solution is that the original source code could be modified heavily and additional compiler analysis is required.

2. Duplicate the required symbol signatures
   We cherry-pick the used symbol signatures into the new file and set up the storage modifier accordingly. The details are described in the issues above.

Implementation:

Solution 2 is preferred. For its implementation, we follow the procedure:

1. Duplicate the original source file.
2. Remove the unnecessary definitions.
3. Move the outlined functions and structures.