# ROSE Compiler Framework/OpenMP Acclerator Model Implementation

ROSE Compiler Framework/OpenMP Support>

## Overview

We are experimenting a trial implementation of the OpenMP accelerator extension in OpenMP 4.0 specification. It is not meant to be the official or complete implementation due to the fast changing nature of the specification and our resource limitations.
The implementation is based on the existing OpenMP implementation in ROSE.

## Implemented Features

The implementation loosely follows the technical report http://www.openmp.org/mp-documents/TR1_167.pdf and later http://www.openmp.org/mp-documents/OpenMP4.0.0.pdf .
* pragma omp target + pragma omp parallel for
* map clauses
* reduction on GPUs
* pragma omp target data
* collapse clause

## Compile and Run

We incrementally release the work as part of the ROSE (EDG 4.x based version) compiler's hosted at https://github.com/rose-compiler/rose-develop
* clone it
* configure and build librose.so under buildtree/src
* build the translator roseompacc under tests/nonsmoke/functional/roseTests/ompLoweringTests/ by typing "make roseompacc"
The translator source is rose-omp-accelerator/tests/nonsmoke/functional/roseTests/ompLoweringTests/roseompacc.C
The command lines (two steps for now) to use roseompacc is exemplified in rose-omp-accelerator/tests/nonsmoke/functional/roseTests/ompLoweringTests/Makefile.am

## Implementation Details
### Parsing

flex/bison
* https://github.com/rose-compiler/rose-develop/blob/master/src/frontend/SageIII/omplexer.ll
* https://github.com/rose-compiler/rose-develop/blob/master/src/frontend/SageIII/ompparser.yy
AST attribute storing parsing results
* https://github.com/rose-compiler/rose-develop/blob/master/src/frontend/SageIII/OmpAttribute.h
* https://github.com/rose-compiler/rose-develop/blob/master/src/frontend/SageIII/OmpAttribute.C
IR generation: ROSETTA, convert attribute to dedicated AST
* https://github.com/rose-compiler/rose-develop/blob/master/src/frontend/SageIII/ompAstConstruction.cpp

### Translation

translation:
* https://github.com/rose-compiler/rose-develop/blob/master/src/midend/programTransformation/ompLowering/omp_lowering.cpp
** void transOmpTargetParallel translate a parallel region associated with a target directive
** void transOmpTargetLoop translate an omp for loop associated with a target directive
** void transOmpMapVariables translate data mapping clauses
Internally, it invokes the outliner acting in mode of classic (passing individual parameters without any wrapping array or structure)
* Outliner::enable_classic = true;

### Runtime Support

runtime support
* https://github.com/rose-compiler/rose-develop/blob/master/src/midend/programTransformation/ompLowering/libxomp.h
* https://github.com/rose-compiler/rose-develop/blob/master/src/midend/programTransformation/ompLowering/xomp.c
* https://github.com/rose-compiler/rose-develop/blob/master/src/midend/programTransformation/ompLowering/xomp_cuda_lib.cu
* https://github.com/rose-compiler/rose-develop/blob/master/src/midend/programTransformation/ompLowering/xomp_cuda_lib_inlined.cu
Rules about what to put into each file:
* libxomp.h runtime interface to compiler writers
* xomp.c runtime support for CPU threading
* xomp_cuda_lib.cu provides the regular C functions wrapping CUDA code (without the __device__ keyword) called by the host code.
* xomp_cuda_lib_inlined.cu provides the __device__ functions called by the CUDA kernel functions (with __global__ keyword).
At the time the work being done, CUDA and/or nvcc does not have linker for device code (linking __global__ and _device_ functions from different files). We have to put some common device functions into this file. So they can be included in the same file in which the generated CUDA kernels are generated (__global__ functions can only call __device__ functions from the same file).

### Device Data Management

The runtime provides two ways to manage device data
* explicit manage of data allocation, copy-to, copy-back, and deallocation
* automatic management using device data environment (DDE) functions
** xomp_deviceDataEnvironmentEnter; // enter an environment , pushed to an environment stack
** xomp_deviceDataEnvironmentPrepareVariable; // reuse data if already in parent environments, or allocate, register with the current environment stack
** xomp_deviceDataEnvironmentExit; // auto deallocation, copy back data. pop environment stack
Examples explicit data management
Automated data management supporting reusing nested device data environments

### Environment variables

Control the number of devices being used in the multi-GPU support by default
The runtime will automatically detect this environment variable and use it to control the GPU count.
You can retrieve this count n the code, like
If OMP_NUM_DEVICES is specified, xomp_get_num_devices will internally call int omp_get_max_devices(void) to obtain the max number of devices supported by the hardware.
Relevant commit

### Loop scheduling

Runtime loop scheduling must be provided since the naive 1-iteration-to-1-thread mapping in CUDA tutorials won't work for large iteration space. (iterations count > total GPU threads).
We provide two loop schedulers in the runtime so each CUDA thread may get more than 1 iterations to handle large iteration space.
* a static even scheduler: this scheduler evenly divides up the iteration space into roughly equal sized chunks. It then assigns each chunk to a CUDA thread. This scheduling policy may stress the memory too much since each thread will touch a large range of data brought in by the iteration chunk
* a round-robin scheduler: Each thread grabs one iteration (or more iterations) at a time. It is identical to OpenMP's schedule(static, chunk) policy.
Testing use of the schedulers (using corresponding CPU versions) are given at:
* https://github.com/rose-compiler/rose-develop/blob/master/src/midend/programTransformation/ompLowering/xomp_accelerator_sched_test_v2.c
GPU version use example: the round-robin scheduler (now the new default scheduler)
Example of how to use the static-even scheduler (no longer the default scheduler in recent release due to performance drawbacks.):

### Testing

Input files under https://github.com/rose-compiler/rose-develop/tree/master/tests/nonsmoke/functional/CompileTests/OpenMP_tests
* axpy_ompacc.c
* axpy_ompacc2.c
* matrixmultiply-ompacc.c
* jacobi-ompacc.c
* jacobi-ompacc-opt1.c // use "target data" region
Test directory https://github.com/rose-compiler/rose-develop/tree/master/tests/nonsmoke/functional/roseTests/ompLoweringTests
* Test translator: roseompacc.C
* Test driver: Makefile.am
* make ompacc_test within the build tree will trigger the associated testing
** Make sure you have nvidia nvcc compiler installed to generate the executables. Otherwise only cuda file will be generated and the final compilation will fail.
Some sample generated CUDA files can be found at:
* https://github.com/chunhualiao/homp-demo

## publication

* C. Liao, Y. Yan, B. de Supinsky, D. Quinlan, B. Chapman, Early Experiences With The OpenMP Accelerator Model, IWOMP 2013, https://e-reports-ext.llnl.gov/pdf/755563.pdf