330d8983d2
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>
94 lines
2.8 KiB
C++
94 lines
2.8 KiB
C++
// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=DEBUG
|
|
// REQUIRES: libomptarget-debug
|
|
|
|
#include <cassert>
|
|
#include <cstdio>
|
|
#include <cstdlib>
|
|
|
|
// Data structure definitions copied from OpenMP RTL.
|
|
struct __tgt_target_non_contig {
|
|
int64_t offset;
|
|
int64_t width;
|
|
int64_t stride;
|
|
};
|
|
|
|
enum tgt_map_type { OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000 };
|
|
|
|
// OpenMP RTL interfaces
|
|
#ifdef __cplusplus
|
|
extern "C" {
|
|
#endif
|
|
void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
|
|
void **args_base, void **args, int64_t *arg_sizes,
|
|
int64_t *arg_types);
|
|
#ifdef __cplusplus
|
|
}
|
|
#endif
|
|
|
|
int main() {
|
|
// case 1
|
|
// int arr[3][4][5][6];
|
|
// #pragma omp target update to(arr[0:2][1:3][1:2][:])
|
|
// set up descriptor
|
|
__tgt_target_non_contig non_contig[5] = {
|
|
{0, 2, 480}, {1, 3, 120}, {1, 2, 24}, {0, 6, 4}, {0, 1, 4}};
|
|
int64_t size = 4, type = OMP_TGT_MAPTYPE_NON_CONTIG;
|
|
|
|
void *base;
|
|
void *begin = &non_contig;
|
|
int64_t *sizes = &size;
|
|
int64_t *types = &type;
|
|
|
|
// The below diagram is the visualization of the non-contiguous transfer after
|
|
// optimization. Note that each element represent the innermost dimension
|
|
// (unit size = 24) since the stride * count of last dimension is equal to the
|
|
// stride of second last dimension.
|
|
//
|
|
// OOOOO OOOOO OOOOO
|
|
// OXXOO OXXOO OOOOO
|
|
// OXXOO OXXOO OOOOO
|
|
// OXXOO OXXOO OOOOO
|
|
__tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
|
|
sizes, types);
|
|
// DEBUG: offset 144
|
|
// DEBUG: offset 264
|
|
// DEBUG: offset 384
|
|
// DEBUG: offset 624
|
|
// DEBUG: offset 744
|
|
// DEBUG: offset 864
|
|
|
|
// case 2
|
|
// double darr[3][4][5];
|
|
// #pragma omp target update to(darr[0:2:2][2:2][:2:2])
|
|
// set up descriptor
|
|
__tgt_target_non_contig non_contig_2[4] = {
|
|
{0, 2, 320}, {2, 2, 40}, {0, 2, 16}, {0, 1, 8}};
|
|
int64_t size_2 = 4, type_2 = OMP_TGT_MAPTYPE_NON_CONTIG;
|
|
|
|
void *base_2;
|
|
void *begin_2 = &non_contig_2;
|
|
int64_t *sizes_2 = &size_2;
|
|
int64_t *types_2 = &type_2;
|
|
|
|
// The below diagram is the visualization of the non-contiguous transfer after
|
|
// optimization. Note that each element represent the innermost dimension
|
|
// (unit size = 24) since the stride * count of last dimension is equal to the
|
|
// stride of second last dimension.
|
|
//
|
|
// OOOOO OOOOO OOOOO
|
|
// OOOOO OOOOO OOOOO
|
|
// XOXOO OOOOO XOXOO
|
|
// XOXOO OOOOO XOXOO
|
|
__tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base_2, &begin_2,
|
|
sizes_2, types_2);
|
|
// DEBUG: offset 80
|
|
// DEBUG: offset 96
|
|
// DEBUG: offset 120
|
|
// DEBUG: offset 136
|
|
// DEBUG: offset 400
|
|
// DEBUG: offset 416
|
|
// DEBUG: offset 440
|
|
// DEBUG: offset 456
|
|
return 0;
|
|
}
|