In a nutshell, this moves our libomptarget code to populate the offload subproject. With this commit, users need to enable the new LLVM/Offload subproject as a runtime in their cmake configuration. No further changes are expected for downstream code. Tests and other components still depend on OpenMP and have also not been renamed. The results below are for a build in which OpenMP and Offload are enabled runtimes. In addition to the pure `git mv`, we needed to adjust some CMake files. Nothing is intended to change semantics. ``` ninja check-offload ``` Works with the X86 and AMDGPU offload tests ``` ninja check-openmp ``` Still works but doesn't build offload tests anymore. ``` ls install/lib ``` Shows all expected libraries, incl. - `libomptarget.devicertl.a` - `libomptarget-nvptx-sm_90.bc` - `libomptarget.rtl.amdgpu.so` -> `libomptarget.rtl.amdgpu.so.18git` - `libomptarget.so` -> `libomptarget.so.18git` Fixes: https://github.com/llvm/llvm-project/issues/75124 --------- Co-authored-by: Saiyedul Islam <Saiyedul.Islam@amd.com>
121 lines
5.6 KiB
C++
121 lines
5.6 KiB
C++
// clang-format off
|
|
// RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic
|
|
// clang-format on
|
|
|
|
// REQUIRES: libomptarget-debug
|
|
|
|
// UNSUPPORTED: nvptx64-nvidia-cuda
|
|
// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
|
|
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
|
|
struct Descriptor {
|
|
int *datum;
|
|
long int x;
|
|
int *more_datum;
|
|
int xi;
|
|
int val_datum, val_more_datum;
|
|
long int arr[1][30];
|
|
int val_arr;
|
|
};
|
|
|
|
int main() {
|
|
Descriptor dat = Descriptor();
|
|
dat.datum = (int *)malloc(sizeof(int) * 10);
|
|
dat.more_datum = (int *)malloc(sizeof(int) * 20);
|
|
dat.xi = 3;
|
|
dat.arr[0][0] = 1;
|
|
|
|
dat.datum[7] = 7;
|
|
dat.more_datum[17] = 17;
|
|
dat.datum[dat.arr[0][0]] = 0;
|
|
|
|
/// The struct is mapped with type 0x0 when the pointer fields are mapped.
|
|
/// The struct is also map explicitely by the user. The second mapping by
|
|
/// the user must not overwrite the mapping set up for the pointer fields
|
|
/// when mapping the struct happens after the mapping of the pointers.
|
|
|
|
// clang-format off
|
|
// CHECK: omptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x0, Name=unknown
|
|
// CHECK: omptarget --> Entry 1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x1000000000001, Name=unknown
|
|
// CHECK: omptarget --> Entry 2: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=40, Type=0x1000000000011, Name=unknown
|
|
// CHECK: omptarget --> Entry 3: Base=[[MORE_DATUM_HST_PTR_BASE:0x.*]], Begin=[[MORE_DATUM_HST_PTR_BEGIN:0x.*]], Size=80, Type=0x1000000000011, Name=unknown
|
|
// clang-format on
|
|
|
|
/// The struct will be mapped in the same order as the above entries.
|
|
|
|
/// First argument is the struct itself and it will be mapped once.
|
|
|
|
// clang-format off
|
|
// CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[DAT_HST_PTR_BASE]], Size=288)...
|
|
// CHECK: PluginInterface --> MemoryManagerTy::allocate: size 288 with host pointer [[DAT_HST_PTR_BASE]].
|
|
// CHECK: omptarget --> Creating new map entry with HstPtrBase=[[DAT_HST_PTR_BASE]], HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtAllocBegin=[[DAT_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1, HoldRefCount=0, Name=unknown
|
|
// CHECK: omptarget --> Moving 288 bytes (hst:[[DAT_HST_PTR_BASE]]) -> (tgt:[[DAT_DEVICE_PTR_BASE]])
|
|
// clang-format on
|
|
|
|
/// Second argument is dat.datum:
|
|
// clang-format off
|
|
// CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[DATUM_HST_PTR_BASE]], Size=40)...
|
|
// CHECK: PluginInterface --> MemoryManagerTy::allocate: size 40 with host pointer [[DATUM_HST_PTR_BASE]].
|
|
// CHECK: omptarget --> Creating new map entry with HstPtrBase=[[DATUM_HST_PTR_BASE]], HstPtrBegin=[[DATUM_HST_PTR_BASE]], TgtAllocBegin=[[DATUM_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DATUM_DEVICE_PTR_BASE]], Size=40, DynRefCount=1, HoldRefCount=0, Name=unknown
|
|
// CHECK: omptarget --> Moving 40 bytes (hst:[[DATUM_HST_PTR_BASE]]) -> (tgt:[[DATUM_DEVICE_PTR_BASE]])
|
|
// clang-format on
|
|
|
|
/// Third argument is dat.more_datum:
|
|
// clang-format off
|
|
// CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], Size=80)...
|
|
// CHECK: PluginInterface --> MemoryManagerTy::allocate: size 80 with host pointer [[MORE_DATUM_HST_PTR_BEGIN]].
|
|
// CHECK: omptarget --> Creating new map entry with HstPtrBase=[[MORE_DATUM_HST_PTR_BEGIN]], HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], TgtAllocBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN:0x.*]], TgtPtrBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN]], Size=80, DynRefCount=1, HoldRefCount=0, Name=unknown
|
|
// CHECK: omptarget --> Moving 80 bytes (hst:[[MORE_DATUM_HST_PTR_BEGIN]]) -> (tgt:[[MORE_DATUM_DEVICE_PTR_BEGIN]])
|
|
// clang-format on
|
|
|
|
#pragma omp target enter data map(to : dat.datum[ : 10]) \
|
|
map(to : dat.more_datum[ : 20]) map(to : dat)
|
|
|
|
/// Checks induced by having a target region:
|
|
// clang-format off
|
|
// CHECK: omptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x223, Name=unknown
|
|
// CHECK: omptarget --> Mapping exists (implicit) with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=2 (incremented), HoldRefCount=0, Name=unknown
|
|
// CHECK: omptarget --> Obtained target argument [[DAT_DEVICE_PTR_BASE]] from host pointer [[DAT_HST_PTR_BASE]]
|
|
// clang-format on
|
|
|
|
#pragma omp target
|
|
{
|
|
dat.xi = 4;
|
|
dat.datum[7]++;
|
|
dat.more_datum[17]++;
|
|
dat.val_datum = dat.datum[7];
|
|
dat.val_more_datum = dat.more_datum[17];
|
|
dat.datum[dat.arr[0][0]] = dat.xi;
|
|
dat.val_arr = dat.datum[dat.arr[0][0]];
|
|
}
|
|
|
|
/// Post-target region checks:
|
|
// clang-format off
|
|
// CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1 (decremented), HoldRefCount=0
|
|
// clang-format on
|
|
|
|
#pragma omp target exit data map(from : dat)
|
|
|
|
/// Target data end checks:
|
|
// clang-format off
|
|
// CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
|
|
// CHECK: omptarget --> Moving 288 bytes (tgt:[[DAT_DEVICE_PTR_BASE]]) -> (hst:[[DAT_HST_PTR_BASE]])
|
|
// clang-format on
|
|
|
|
// CHECK: dat.xi = 4
|
|
// CHECK: dat.val_datum = 8
|
|
// CHECK: dat.val_more_datum = 18
|
|
// CHECK: dat.datum[dat.arr[0][0]] = 0
|
|
// CHECK: dat.val_arr = 4
|
|
|
|
printf("dat.xi = %d\n", dat.xi);
|
|
printf("dat.val_datum = %d\n", dat.val_datum);
|
|
printf("dat.val_more_datum = %d\n", dat.val_more_datum);
|
|
printf("dat.datum[dat.arr[0][0]] = %d\n", dat.datum[dat.arr[0][0]]);
|
|
printf("dat.val_arr = %d\n", dat.val_arr);
|
|
|
|
return 0;
|
|
}
|