Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Generate similiar *.ll code as Clang does for #pragma omp target parallel #49

Open
wants to merge 15 commits into
base: aomp-dev
Choose a base branch
from
Open

Generate similiar *.ll code as Clang does for #pragma omp target parallel #49

wants to merge 15 commits into from

Commits on May 10, 2022

  1. Start new control flow for SPMD

    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    62f21e9 View commit details
    Browse the repository at this point in the history

  2. Modify deinit calls

    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    f5e8e7d View commit details
    Browse the repository at this point in the history

  3. Added first version of calling kmpc_parallel_51 

    Done:
   Added declaration of kmpc_parallel_51
   Set constant args
   Added logic for setting up size of array which will
     contain target symbols
   Calling kmpc_parallel_51
   Moved deinit function just after kmpc_parallel_51

Not done:
   Passing symbols to target array
   Creating separate function which will reflect kernel code

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    bf15996 View commit details
    Browse the repository at this point in the history

  4. Add function which initializes mapped symbols. 

    kmpc_parallel_51 requires that offloaded symbols are passed
as addresses inside pointer array

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    fce4694 View commit details
    Browse the repository at this point in the history

  5. Add support for SPMD kernels with new OpenMP API. 

    Flang generates new SPMD kernels which use kmpc_parallel_51
function.

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    063001a View commit details
    Browse the repository at this point in the history

  6. Fixed handling of last argument. 

    Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    fe8d5ee View commit details
    Browse the repository at this point in the history

  7. Fixed passing arrays for spmd kernels. 

    Do not assign new dtype value for device symbols via
get_type function. Use macros PASSBYVALP PASSBYREFP instead.

Remove hack for load/store of the last symbol.

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    3ee5166 View commit details
    Browse the repository at this point in the history

  8. Fix passing args to helper function 

    Fixed passing scalars which type is different than int64.

Fixed passing allocatable arrays.

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    253c35e View commit details
    Browse the repository at this point in the history

  9. Pass complex arg as pointer to pair of floats

    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    e84c6e5 View commit details
    Browse the repository at this point in the history

  10. Do not modify the LLType of the argument 

    If we generate initialization function for SPMD
kernels we need to store addresses of the arguments before
we call kmpc_parallel_51 function. We use ptrtoint instruction
for scalar variables. Before this patch the LLVM IR code was generated
wrongly for complex variables:

void kernel_func(<float, float> *Arg_c)
//some code
ptrtoint i64* %Arg_c //error Arg_c was declared as pair of floats
//some code
call kmpc_parallel_51()
//some code

This patch causes that LLVM IR contains correct ptrtoint instruction:
void kernel_func(<float, float> *Arg_c)
//some code
ptrtoint <float, float>* %Arg_c //ok, Arg_c was declared as pair of floats
//some code
call kmpc_parallel_51()
//some code

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    b713196 View commit details
    Browse the repository at this point in the history

  11. Fix passing Fortran objects 

    Fortran objects should be passed as i64*

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    ce4a1b4 View commit details
    Browse the repository at this point in the history

  12. Fixed passing complex numbers. 

    Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    c9307ca View commit details
    Browse the repository at this point in the history

  13. Skip symbols which are not initialized 

    Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    8ed9589 View commit details
    Browse the repository at this point in the history

  14. Add support for passing integers by value 

    Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    97bd64e View commit details
    Browse the repository at this point in the history

  15. Use the same OpenMP API as Clang for target parallel for pragma 

    Clang uses kmpc_parallel_51 function for handlig target parallel for
pragma.

Flang should use the same functions as Clang for pragma
target parallel

Signed-off-by: Dominik Adamski <Dominik.Adamski@amd.com>
    @DominikAdamski
    DominikAdamski committed May 10, 2022
    Configuration menu
    Copy the full SHA
    3fb2323 View commit details
    Browse the repository at this point in the history