[libc-commits] [PATCH] D139839: [libc] Add a loader utility for AMDHSA architectures for testing

[Johannes Doerfert via Phabricator via libc-commits]

jdoerfert added a comment.

In D139839#3993272 <https://reviews.llvm.org/D139839#3993272>, @sivachandra wrote:

> On the general topic of adding dependencies on other projects like `openmp`, we should strive to keep the libc project as self-contained as possible.

The problem is that now we duplicate (and in the future extend) the "gpu launch" logic for each target.
We are effectively creating yet another offloading runtime. We will invent new abstraction layers, new features will be added, etc.

This patch has 400 lines of HSA magic copied probably from OpenMP. We will need the same for CUDA, OneAPI, maybe OpenCL, ...
Then we need extra logic in all of them to support allocators (via pre-allocation), in all their shapes (bump, free-lists, ...).
Then we need to put RPC logic in them, again per target, and all of it is then not yet tested with OpenMP offload.

A loader using the existing offloading runtime:

1. does not require clang changes
2. requires 2 minimal files (shown below)
3. will work on all supported GPU targets, the MPI target, virtual GPU target, ...



   // main.c
  #include <string.h>
  extern int __user_main(int, char *[]);
  
  int main(int argc, char *argv[]) {
    #pragma omp target enter data map(to: argv[:argc])
    for (int I = 0; I < argc; ++I) {
      size_t Len = strlen(argv[I]);
      #pragma omp target enter data map(to: argv[I][:Len])
    }
  
    int Ret;
    #pragma omp target teams num_teams(1) thread_limit(1024) map(from: Ret)
    { Ret = __user_main(argc, argv); }
    return Ret;
  }



   // header.h, could be replaced by a clang command line flag
  #pragma omp begin declare target device_type(nohost)
  int main(int, char *[]) asm("__user_main");


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D139839/new/

https://reviews.llvm.org/D139839