Index: openmp/libomptarget/include/interop.h =================================================================== --- openmp/libomptarget/include/interop.h +++ openmp/libomptarget/include/interop.h @@ -116,30 +116,6 @@ extern const char *__KAI_KMPC_CONVENTION omp_get_interop_rc_desc(const omp_interop_t, omp_interop_rc_t); -typedef struct kmp_tasking_flags { /* Total struct must be exactly 32 bits */ - /* Compiler flags */ /* Total compiler flags must be 16 bits */ - unsigned tiedness : 1; /* task is either tied (1) or untied (0) */ - unsigned final : 1; /* task is final(1) so execute immediately */ - unsigned merged_if0 : 1; // no __kmpc_task_{begin/complete}_if0 calls in if0 - unsigned destructors_thunk : 1; // set if the compiler creates a thunk to - unsigned proxy : 1; // task is a proxy task (it will be executed outside the - unsigned priority_specified : 1; // set if the compiler provides priority - unsigned detachable : 1; // 1 == can detach */ - unsigned unshackled : 1; /* 1 == unshackled task */ - unsigned target : 1; /* 1 == target task */ - unsigned reserved : 7; /* reserved for compiler use */ - unsigned tasktype : 1; /* task is either explicit(1) or implicit (0) */ - unsigned task_serial : 1; // task is executed immediately (1) or deferred (0) - unsigned tasking_ser : 1; // all tasks in team are either executed immediately - unsigned team_serial : 1; // entire team is serial (1) [1 thread] or parallel - unsigned started : 1; /* 1==started, 0==not started */ - unsigned executing : 1; /* 1==executing, 0==not executing */ - unsigned complete : 1; /* 1==complete, 0==not complete */ - unsigned freed : 1; /* 1==freed, 0==allocated */ - unsigned native : 1; /* 1==gcc-compiled task, 0==intel */ - unsigned reserved31 : 7; /* reserved for library use */ -} kmp_tasking_flags_t; - typedef enum omp_interop_backend_type_t { // reserve 0 omp_interop_backend_type_cuda_1 = 1, Index: openmp/libomptarget/src/api.cpp =================================================================== --- openmp/libomptarget/src/api.cpp +++ openmp/libomptarget/src/api.cpp @@ -15,6 +15,8 @@ #include "private.h" #include "rtl.h" +#include "llvm/ADT/SmallVector.h" + #include #include #include @@ -200,6 +202,105 @@ return Rc; } +// The helper function that calls omp_target_memcpy or omp_target_memcpy_rect +static int __kmpc_target_memcpy_async_helper(kmp_int32 Gtid, kmp_task_t *Task) { + if (Task == nullptr) + return OFFLOAD_FAIL; + + TargetMemcpyArgsTy *Args = (TargetMemcpyArgsTy *)Task->shareds; + + if (Args == nullptr) + return OFFLOAD_FAIL; + + // Call blocked version + int Rc = OFFLOAD_SUCCESS; + if (Args->IsRectMemcpy) { + Rc = omp_target_memcpy_rect( + Args->Dst, Args->Src, Args->ElementSize, Args->NumDims, Args->Volume, + Args->DstOffsets, Args->SrcOffsets, Args->DstDimensions, + Args->SrcDimensions, Args->DstDevice, Args->SrcDevice); + + DP("omp_target_memcpy_rect returns %d\n", Rc); + } else { + Rc = omp_target_memcpy(Args->Dst, Args->Src, Args->Length, Args->DstOffset, + Args->SrcOffset, Args->DstDevice, Args->SrcDevice); + + DP("omp_target_memcpy returns %d\n", Rc); + } + + // Release the arguments object + delete Args; + + return Rc; +} + +// Allocate and launch helper task +static int __kmpc_helper_task_creation(TargetMemcpyArgsTy *Args, + int DepObjCount, + omp_depend_t *DepObjList) { + // Create global thread ID + int Gtid = __kmpc_global_thread_num(nullptr); + int (*Fn)(kmp_int32, kmp_task_t *) = &__kmpc_target_memcpy_async_helper; + + // Setup the hidden helper flags; + kmp_int32 Flags = 0; + kmp_tasking_flags_t *InputFlags = (kmp_tasking_flags_t *)&Flags; + InputFlags->hidden_helper = 1; + + // Alloc helper task + kmp_task_t *Ptr = __kmpc_omp_target_task_alloc(nullptr, Gtid, Flags, + sizeof(kmp_task_t), 0, Fn, -1); + + if (Ptr == nullptr) { + // Task allocation failed, delete the argument object + delete Args; + + return OFFLOAD_FAIL; + } + + // Setup the arguments passed to helper task + Ptr->shareds = Args; + + // Convert the type of depend objects + llvm::SmallVector DepObjs; + for (int i = 0; i < DepObjCount; i++) { + omp_depend_t DepObj = DepObjList[i]; + DepObjs.push_back(*((kmp_depend_info_t *)DepObj)); + } + + // Launch the helper task + int Rc = __kmpc_omp_task_with_deps(nullptr, Gtid, Ptr, DepObjCount, + DepObjs.data(), 0, nullptr); + + return Rc; +} + +EXTERN int omp_target_memcpy_async(void *Dst, const void *Src, size_t Length, + size_t DstOffset, size_t SrcOffset, + int DstDevice, int SrcDevice, + int DepObjCount, omp_depend_t *DepObjList) { + TIMESCOPE(); + DP("Call to omp_target_memcpy_async, dst device %d, src device %d, " + "dst addr " DPxMOD ", src addr " DPxMOD ", dst offset %zu, " + "src offset %zu, length %zu\n", + DstDevice, SrcDevice, DPxPTR(Dst), DPxPTR(Src), DstOffset, SrcOffset, + Length); + + // Check the source and dest address + if (Dst == nullptr || Src == nullptr) + return OFFLOAD_FAIL; + + // Create task object + TargetMemcpyArgsTy *Args = new TargetMemcpyArgsTy( + Dst, Src, Length, DstOffset, SrcOffset, DstDevice, SrcDevice); + + // Create and launch helper task + int Rc = __kmpc_helper_task_creation(Args, DepObjCount, DepObjList); + + DP("omp_target_memcpy_async returns %d\n", Rc); + return Rc; +} + EXTERN int omp_target_memcpy_rect(void *Dst, const void *Src, size_t ElementSize, int NumDims, const size_t *Volume, @@ -260,6 +361,36 @@ return Rc; } +EXTERN int omp_target_memcpy_rect_async( + void *Dst, const void *Src, size_t ElementSize, int NumDims, + const size_t *Volume, const size_t *DstOffsets, const size_t *SrcOffsets, + const size_t *DstDimensions, const size_t *SrcDimensions, int DstDevice, + int SrcDevice, int DepObjCount, omp_depend_t *DepObjList) { + TIMESCOPE(); + DP("Call to omp_target_memcpy_rect_async, dst device %d, src device %d, " + "dst addr " DPxMOD ", src addr " DPxMOD ", dst offsets " DPxMOD ", " + "src offsets " DPxMOD ", dst dims " DPxMOD ", src dims " DPxMOD ", " + "volume " DPxMOD ", element size %zu, num_dims %d\n", + DstDevice, SrcDevice, DPxPTR(Dst), DPxPTR(Src), DPxPTR(DstOffsets), + DPxPTR(SrcOffsets), DPxPTR(DstDimensions), DPxPTR(SrcDimensions), + DPxPTR(Volume), ElementSize, NumDims); + + // Check the source and dest address + if (Dst == nullptr || Src == nullptr) + return OFFLOAD_FAIL; + + // Create task object + TargetMemcpyArgsTy *Args = new TargetMemcpyArgsTy( + Dst, Src, ElementSize, NumDims, Volume, DstOffsets, SrcOffsets, + DstDimensions, SrcDimensions, DstDevice, SrcDevice); + + // Create and launch helper task + int Rc = __kmpc_helper_task_creation(Args, DepObjCount, DepObjList); + + DP("omp_target_memcpy_rect_async returns %d\n", Rc); + return Rc; +} + EXTERN int omp_target_associate_ptr(const void *HostPtr, const void *DevicePtr, size_t Size, size_t DeviceOffset, int DeviceNum) { Index: openmp/libomptarget/src/exports =================================================================== --- openmp/libomptarget/src/exports +++ openmp/libomptarget/src/exports @@ -38,6 +38,8 @@ omp_target_is_present; omp_target_memcpy; omp_target_memcpy_rect; + omp_target_memcpy_async; + omp_target_memcpy_rect_async; omp_target_associate_ptr; omp_target_disassociate_ptr; llvm_omp_target_alloc_host; Index: openmp/libomptarget/src/private.h =================================================================== --- openmp/libomptarget/src/private.h +++ openmp/libomptarget/src/private.h @@ -98,7 +98,51 @@ * We maintain the same data structure for compatibility. */ typedef int kmp_int32; +typedef int64_t kmp_int64; typedef intptr_t kmp_intptr_t; + +typedef void *omp_depend_t; +struct kmp_task; +typedef kmp_int32 (*kmp_routine_entry_t)(kmp_int32, struct kmp_task *); +typedef struct kmp_task { + void *shareds; + kmp_routine_entry_t routine; + kmp_int32 part_id; +} kmp_task_t; + +typedef struct kmp_tasking_flags { /* Total struct must be exactly 32 bits */ + /* Compiler flags */ /* Total compiler flags must be 16 bits */ + unsigned tiedness : 1; /* task is either tied (1) or untied (0) */ + unsigned final : 1; /* task is final(1) so execute immediately */ + unsigned merged_if0 : 1; /* no __kmpc_task_{begin/complete}_if0 calls in if0 + code path */ + unsigned destructors_thunk : 1; /* set if the compiler creates a thunk to + invoke destructors from the runtime */ + unsigned proxy : 1; /* task is a proxy task (it will be executed outside the + context of the RTL) */ + unsigned priority_specified : 1; /* set if the compiler provides priority + setting for the task */ + unsigned detachable : 1; /* 1 == can detach */ + unsigned hidden_helper : 1; /* 1 == hidden helper task */ + unsigned reserved : 8; /* reserved for compiler use */ + + /* Library flags */ /* Total library flags must be 16 bits */ + unsigned tasktype : 1; /* task is either explicit(1) or implicit (0) */ + unsigned task_serial : 1; // task is executed immediately (1) or deferred (0) + unsigned tasking_ser : 1; // all tasks in team are either executed immediately + // (1) or may be deferred (0) + unsigned team_serial : 1; // entire team is serial (1) [1 thread] or parallel + // (0) [>= 2 threads] + /* If either team_serial or tasking_ser is set, task team may be NULL */ + /* Task State Flags: */ + unsigned started : 1; /* 1==started, 0==not started */ + unsigned executing : 1; /* 1==executing, 0==not executing */ + unsigned complete : 1; /* 1==complete, 0==not complete */ + unsigned freed : 1; /* 1==freed, 0==allocated */ + unsigned native : 1; /* 1==gcc-compiled task, 0==intel */ + unsigned reserved31 : 7; /* reserved for library use */ +} kmp_tasking_flags_t; + // Compiler sends us this info: typedef struct kmp_depend_info { kmp_intptr_t base_addr; @@ -117,6 +161,88 @@ kmp_depend_info_t *dep_list, kmp_int32 ndeps_noalias, kmp_depend_info_t *noalias_dep_list) __attribute__((weak)); + +kmp_task_t *__kmpc_omp_task_alloc(ident_t *loc_ref, kmp_int32 gtid, + kmp_int32 flags, size_t sizeof_kmp_task_t, + size_t sizeof_shareds, + kmp_routine_entry_t task_entry) + __attribute__((weak)); + +kmp_task_t * +__kmpc_omp_target_task_alloc(ident_t *loc_ref, kmp_int32 gtid, kmp_int32 flags, + size_t sizeof_kmp_task_t, size_t sizeof_shareds, + kmp_routine_entry_t task_entry, + kmp_int64 device_id) __attribute__((weak)); + +kmp_int32 __kmpc_omp_task_with_deps(ident_t *loc_ref, kmp_int32 gtid, + kmp_task_t *new_task, kmp_int32 ndeps, + kmp_depend_info_t *dep_list, + kmp_int32 ndeps_noalias, + kmp_depend_info_t *noalias_dep_list) + __attribute__((weak)); + +/** + * The argument set that is passed from asynchronous memory copy to block + * version of memory copy invoked in helper task + */ +struct TargetMemcpyArgsTy { + /** + * Common attribuutes + */ + void *Dst; + const void *Src; + int DstDevice; + int SrcDevice; + + /** + * The flag that denotes single dimensional or rectangle dimensional copy + */ + bool IsRectMemcpy; + + /** + * Arguments for single dimensional copy + */ + size_t Length; + size_t DstOffset; + size_t SrcOffset; + + /** + * Arguments for rectangle dimensional copy + */ + size_t ElementSize; + int NumDims; + const size_t *Volume; + const size_t *DstOffsets; + const size_t *SrcOffsets; + const size_t *DstDimensions; + const size_t *SrcDimensions; + + /** + * Constructor for single dimensional copy + */ + TargetMemcpyArgsTy(void *Dst, const void *Src, size_t Length, + size_t DstOffset, size_t SrcOffset, int DstDevice, + int SrcDevice) + : Dst(Dst), Src(Src), DstDevice(DstDevice), SrcDevice(SrcDevice), + IsRectMemcpy(false), Length(Length), DstOffset(DstOffset), + SrcOffset(SrcOffset), ElementSize(0), NumDims(0), Volume(0), + DstOffsets(0), SrcOffsets(0), DstDimensions(0), SrcDimensions(0){}; + + /** + * Constructor for rectangle dimensional copy + */ + TargetMemcpyArgsTy(void *Dst, const void *Src, size_t ElementSize, + int NumDims, const size_t *Volume, + const size_t *DstOffsets, const size_t *SrcOffsets, + const size_t *DstDimensions, const size_t *SrcDimensions, + int DstDevice, int SrcDevice) + : Dst(Dst), Src(Src), DstDevice(DstDevice), SrcDevice(SrcDevice), + IsRectMemcpy(true), Length(0), DstOffset(0), SrcOffset(0), + ElementSize(ElementSize), NumDims(NumDims), Volume(Volume), + DstOffsets(DstOffsets), SrcOffsets(SrcOffsets), + DstDimensions(DstDimensions), SrcDimensions(SrcDimensions){}; +}; + #ifdef __cplusplus } #endif Index: openmp/libomptarget/test/api/omp_target_memcpy_async1.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/api/omp_target_memcpy_async1.c @@ -0,0 +1,46 @@ +// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda + +// Test case for omp_target_memcpy_async, oringally from GCC + +#include +#include "stdio.h" +#include + +int main () { + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int q[128], i; + void *p; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + p = omp_target_alloc (130 * sizeof (int), d); + if (p == NULL) + return 0; + + for (i = 0; i < 128; i++) + q[i] = i; + + if (omp_target_memcpy_async(p, q, 128 * sizeof (int), sizeof (int), 0, d, id, 0, NULL)) { + abort(); + } + +#pragma omp taskwait + + int q2[128]; + for (i = 0; i < 128; ++i) + q2[i] = 0; + if (omp_target_memcpy_async (q2, p, 128 * sizeof(int), 0, sizeof (int), id, d, 0, NULL)) + abort (); + +#pragma omp taskwait + + for (i = 0; i < 128; ++i) + if (q2[i] != q[i]) + abort (); + + omp_target_free (p, d); + + return 0; +} Index: openmp/libomptarget/test/api/omp_target_memcpy_async2.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/api/omp_target_memcpy_async2.c @@ -0,0 +1,73 @@ +// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda + +#include +#include "stdio.h" +#include + +int main() { + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int a[128], b[64], c[32], e[16], q[128], i; + void *p; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + p = omp_target_alloc (130 * sizeof (int), d); + if (p == NULL) + return 0; + + for (i = 0; i < 128; ++i) + a[i] = i + 1; + for (i = 0; i < 64; ++i) + b[i] = i + 2; + for (i = 0; i < 32; i++) + c[i] = 0; + for (i = 0; i < 16; i++) + e[i] = i + 4; + + omp_depend_t obj[2]; + +#pragma omp parallel num_threads(5) +#pragma omp single + { +#pragma omp task depend(out: p) + omp_target_memcpy (p, a, 128 * sizeof (int), 0, 0, d, id); + +#pragma omp task depend(inout: p) + omp_target_memcpy (p, b, 64 * sizeof (int), 0, 0, d, id); + +#pragma omp task depend(out: c) + for (i = 0; i < 32; i++) + c[i] = i + 3; + +#pragma omp depobj(obj[0]) depend(inout: p) +#pragma omp depobj(obj[1]) depend(in: c) + omp_target_memcpy_async (p, c, 32 * sizeof (int), 0, 0, d, id, 2, obj); + +#pragma omp task depend(in: p) + omp_target_memcpy (p, e, 16 * sizeof (int), 0, 0, d, id); + } + +#pragma omp taskwait + + for (i = 0; i < 128; ++i) + q[i] = 0; + omp_target_memcpy(q, p, 128 * sizeof(int), 0, 0, id, d); + for (i = 0; i < 16; ++i) + if (q[i] != i + 4) + abort(); + for (i = 16; i < 32; ++i) + if (q[i] != i + 3) + abort(); + for (i = 32; i < 64; ++i) + if (q[i] != i + 2) + abort(); + for (i = 64; i < 128; ++i) + if (q[i] != i + 1) + abort(); + + omp_target_free (p, d); + + return 0; +} Index: openmp/libomptarget/test/api/omp_target_memcpy_rect_async1.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/api/omp_target_memcpy_rect_async1.c @@ -0,0 +1,66 @@ +// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda + +#include +#include +#include + +#define NUM_DIMS 3 + +int main() { + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int q[128], q2[128], i; + void *p; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + p = omp_target_alloc (130 * sizeof (int), d); + if (p == NULL) + return 0; + + if (omp_target_memcpy_rect_async (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, + NULL, d, id, 0, NULL) < 3 + || omp_target_memcpy_rect_async (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, + NULL, id, d, 0, NULL) < 3 + || omp_target_memcpy_rect_async (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, + NULL, id, id, 0, NULL) < 3) + abort (); + + for (i = 0; i < 128; i++) + q[i] = 0; + if (omp_target_memcpy (p, q, 128 * sizeof (int), 0, 0, d, id) != 0) + abort (); + + for (i = 0; i < 128; i++) + q[i] = i + 1; + + size_t volume[NUM_DIMS] = { 1, 2, 3 }; + size_t dst_offsets[NUM_DIMS] = { 0, 0, 0 }; + size_t src_offsets[NUM_DIMS] = { 0, 0, 0 }; + size_t dst_dimensions[NUM_DIMS] = { 3, 4, 5 }; + size_t src_dimensions[NUM_DIMS] = { 2, 3, 4 }; + + if (omp_target_memcpy_rect_async (p, q, sizeof (int), NUM_DIMS, volume, + dst_offsets, src_offsets, dst_dimensions, + src_dimensions, d, id, 0, NULL) != 0) + abort (); + +#pragma omp taskwait + + for (i = 0; i < 128; i++) + q2[i] = 0; + if (omp_target_memcpy (q2, p, 128 * sizeof (int), 0, 0, id, d) != 0) + abort (); + + /* q2 is expected to contain: 1 2 3 0 0 5 6 7 0 0 .. 0 */ + if (q2[0] != 1 || q2[1] != 2 || q2[2] !=3 || q2[3] != 0 || q2[4] != 0 + || q2[5] != 5 || q2[6] != 6 || q2[7] != 7) + abort (); + for (i = 8; i < 128; ++i) + if (q2[i] != 0) + abort (); + + omp_target_free (p, d); + return 0; +} Index: openmp/libomptarget/test/api/omp_target_memcpy_rect_async2.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/api/omp_target_memcpy_rect_async2.c @@ -0,0 +1,89 @@ +// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda + +#include +#include + +#define NUM_DIMS 3 + +int main () { + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int a[128], b[64], c[128], e[16], q[128], i; + void *p; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + p = omp_target_alloc (130 * sizeof (int), d); + if (p == NULL) + return 0; + + for (i = 0; i < 128; i++) + q[i] = 0; + if (omp_target_memcpy (p, q, 128 * sizeof (int), 0, 0, d, id) != 0) + abort (); + + size_t volume[NUM_DIMS] = { 2, 2, 3 }; + size_t dst_offsets[NUM_DIMS] = { 0, 0, 0 }; + size_t src_offsets[NUM_DIMS] = { 0, 0, 0 }; + size_t dst_dimensions[NUM_DIMS] = { 3, 4, 5 }; + size_t src_dimensions[NUM_DIMS] = { 2, 3, 4 }; + + for (i = 0; i < 128; i++) + a[i] = 42; + for (i = 0; i < 64; i++) + b[i] = 24; + for (i = 0; i < 128; i++) + c[i] = 0; + for (i = 0; i < 16; i++) + e[i] = 77; + + omp_depend_t obj[2]; + +#pragma omp parallel num_threads(5) +#pragma omp single + { +#pragma omp task depend (out: p) + omp_target_memcpy (p, a, 128 * sizeof (int), 0, 0, d, id); + +#pragma omp task depend(inout: p) + omp_target_memcpy (p, b, 64 * sizeof (int), 0, 0, d, id); + +#pragma omp task depend(out: c) + for (i = 0; i < 128; i++) + c[i] = i + 1; + +#pragma omp depobj(obj[0]) depend(inout: p) +#pragma omp depobj(obj[1]) depend(in: c) + + /* This produces: 1 2 3 - - 5 6 7 - - at positions 0..9 and + 13 14 15 - - 17 18 19 - - at positions 20..29. */ + omp_target_memcpy_rect_async (p, c, sizeof (int), NUM_DIMS, volume, + dst_offsets, src_offsets, dst_dimensions, + src_dimensions, d, id, 2, obj); + +#pragma omp task depend(in: p) + omp_target_memcpy (p, e, 16 * sizeof (int), 0, 0, d, id); + } + +#pragma omp taskwait + + if (omp_target_memcpy (q, p, 128 * sizeof(int), 0, 0, id, d) != 0) + abort (); + + for (i = 0; i < 16; ++i) + if (q[i] != 77) + abort (); + if (q[20] != 13 || q[21] != 14 || q[22] != 15 || q[25] != 17 || q[26] != 18 + || q[27] != 19) + abort (); + for (i = 28; i < 64; ++i) + if (q[i] != 24) + abort (); + for (i = 64; i < 128; ++i) + if (q[i] != 42) + abort (); + + omp_target_free (p, d); + return 0; +}