/// creates host to the target data mapping, store it in the
/// libtarget.so internal structure (an entry in a stack of data maps)
/// and passes the data to the device;
EXTERN void __tgt_target_data_begin(int32_t device_id, int32_t arg_num,
void** args_base, void **args, int64_t *arg_sizes, int32_t *arg_types)
{
DP("Entering data begin region for device %d with %d mappings\n",
device_id, arg_num);
// No devices available?
if (device_id == OFFLOAD_DEVICE_DEFAULT) {
device_id = omp_get_default_device();
DP("Use default device id %d\n", device_id);
}
if (Devices.size() <= (size_t)device_id){
DP("Device ID %d does not have a matching RTL.\n", device_id);
return;
}
// Get device info
DeviceTy & Device = Devices[device_id];
// Init the device if not done before
if (!Device.IsInit) {
if (Device.init() != OFFLOAD_SUCCESS) {
DP("failed to init device %d\n", device_id);
return;
}
}
target_data_begin(Device, arg_num, args_base, args, arg_sizes, arg_types);
}
int
main ()
{
int d = omp_get_default_device ();
int id = omp_get_initial_device ();
if (d < 0 || d >= omp_get_num_devices ())
d = id;
int a[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
int *b = a;
int shared_mem = 0;
#pragma omp target map (alloc: shared_mem)
shared_mem = 1;
if (omp_target_is_present (b, d) != shared_mem)
abort ();
#pragma omp target enter data map (to: a)
if (omp_target_is_present (b, d) == 0)
abort ();
#pragma omp target enter data map (alloc: b[:0])
if (omp_target_is_present (b, d) == 0)
abort ();
#pragma omp target exit data map (release: b[:0])
if (omp_target_is_present (b, d) == 0)
abort ();
#pragma omp target exit data map (release: b[:0])
if (omp_target_is_present (b, d) != shared_mem)
abort ();
#pragma omp target enter data map (to: a)
if (omp_target_is_present (b, d) == 0)
abort ();
#pragma omp target enter data map (always, to: b[:0])
if (omp_target_is_present (b, d) == 0)
abort ();
#pragma omp target exit data map (delete: b[:0])
if (omp_target_is_present (b, d) != shared_mem)
abort ();
#pragma omp target exit data map (from: b[:0])
return 0;
}
/// passes data to/from the target
EXTERN void __tgt_target_data_update(int32_t device_id, int32_t arg_num,
void** args_base, void **args, int64_t *arg_sizes, int32_t *arg_types)
{
DP("Entering data update with %d mappings\n", arg_num);
// No devices available?
if (device_id == OFFLOAD_DEVICE_DEFAULT) {
device_id = omp_get_default_device();
}
if (Devices.size() <= (size_t)device_id){
DP("Device ID %d does not have a matching RTL.\n", device_id);
return;
}
// Get device info
DeviceTy & Device = Devices[device_id];
if (!Device.IsInit) {
DP("uninit device: ignore\n");
return;
}
// process each input
for(int32_t i=0; i<arg_num; ++i) {
void *HstPtrBegin = args[i];
//void *HstPtrBase = args_base[i];
long IsLast;
void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast, false);
if (arg_types[i] & tgt_map_from) {
DP("Moving %ld bytes (tgt:%016lx) -> (hst:%016lx)\n", (long)arg_sizes[i],
(long)TgtPtrBegin, (long)HstPtrBegin);
Device.data_retrieve(HstPtrBegin, TgtPtrBegin, arg_sizes[i]);
}
if (arg_types[i] & tgt_map_to) {
DP("Moving %ld bytes (hst:%016lx) -> (tgt:%016lx)\n", (long)arg_sizes[i],
(long)HstPtrBegin, (long)TgtPtrBegin);
Device.data_submit(TgtPtrBegin, HstPtrBegin, arg_sizes[i]);
}
}
}
/// passes data from the target, release target memory and destroys
/// the host-target mapping (top entry from the stack of data maps)
/// created by the last __tgt_target_data_begin
EXTERN void __tgt_target_data_end(int32_t device_id, int32_t arg_num,
void** args_base, void **args, int64_t *arg_sizes, int32_t *arg_types)
{
DP("Entering data end region with %d mappings\n", arg_num);
// No devices available?
if (device_id == OFFLOAD_DEVICE_DEFAULT) {
device_id = omp_get_default_device();
}
if (Devices.size() <= (size_t)device_id){
DP("Device ID %d does not have a matching RTL.\n", device_id);
return;
}
DeviceTy & Device = Devices[device_id];
if (!Device.IsInit) {
DP("uninit device: ignore\n");
return;
}
target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types);
}
int32_t
omp_get_default_device_ (void)
{
return omp_get_default_device ();
}
int main(int argc, const char * argv[]) {
int i, fn=0;
#pragma omp target
#pragma omp parallel
{
omp_set_default_device(4);
printf("Soy el thread %d y esta region tiene: %d devices\n",omp_get_thread_num(),omp_get_default_device());
}
#pragma omp target
#pragma omp parallel
{
printf("Soy el thread %d y esta region tiene: %d devices\n",omp_get_thread_num(),omp_get_default_device());
}
return 0;
}
int
main ()
{
int d = omp_get_default_device ();
int id = omp_get_initial_device ();
int err;
int q[128], i;
void *p;
if (d < 0 || d >= omp_get_num_devices ())
d = id;
for (i = 0; i < 128; i++)
q[i] = i;
p = omp_target_alloc (130 * sizeof (int), d);
if (p == NULL)
return 0;
if (omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL, NULL,
d, id) < 3
|| omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
NULL, id, d) < 3
|| omp_target_memcpy_rect (NULL, NULL, 0, 0, NULL, NULL, NULL, NULL,
NULL, id, id) < 3)
abort ();
if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) == 0)
{
size_t volume[3] = { 128, 0, 0 };
size_t dst_offsets[3] = { 0, 0, 0 };
size_t src_offsets[3] = { 1, 0, 0 };
size_t dst_dimensions[3] = { 128, 0, 0 };
size_t src_dimensions[3] = { 128, 0, 0 };
if (omp_target_associate_ptr (q, p, 128 * sizeof (int), sizeof (int), d) != 0)
abort ();
if (omp_target_is_present (q, d) != 1
|| omp_target_is_present (&q[32], d) != 1
|| omp_target_is_present (&q[127], d) != 1)
abort ();
if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0,
d, id) != 0)
abort ();
#pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
{
int j;
err = 0;
for (j = 0; j < 128; j++)
if (q[j] != j)
err = 1;
else
q[j] += 4;
}
if (err)
abort ();
if (omp_target_memcpy_rect (q, p, sizeof (int), 1, volume,
dst_offsets, src_offsets, dst_dimensions,
src_dimensions, id, d) != 0)
abort ();
for (i = 0; i < 128; i++)
if (q[i] != i + 4)
abort ();
volume[2] = 2;
volume[1] = 3;
volume[0] = 6;
dst_offsets[2] = 1;
dst_offsets[1] = 0;
dst_offsets[0] = 0;
src_offsets[2] = 1;
src_offsets[1] = 0;
src_offsets[0] = 3;
dst_dimensions[2] = 2;
dst_dimensions[1] = 3;
dst_dimensions[0] = 6;
src_dimensions[2] = 3;
src_dimensions[1] = 4;
src_dimensions[0] = 6;
if (omp_target_memcpy_rect (p, q, sizeof (int), 3, volume,
dst_offsets, src_offsets, dst_dimensions,
src_dimensions, d, id) != 0)
abort ();
#pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
{
int j, k, l;
err = 0;
for (j = 0; j < 6; j++)
for (k = 0; k < 3; k++)
for (l = 0; l < 2; l++)
if (q[j * 6 + k * 2 + l] != 3 * 12 + 4 + 1 + l + k * 3 + j * 12)
err = 1;
}
if (err)
abort ();
if (omp_target_memcpy (p, p, 10 * sizeof (int), 51 * sizeof (int),
111 * sizeof (int), d, d) != 0)
abort ();
#pragma omp target if (d >= 0) device (d >= 0 ? d : 0) map(alloc:q[0:32]) map(from:err)
{
int j;
err = 0;
for (j = 0; j < 10; j++)
if (q[50 + j] != q[110 + j])
err = 1;
}
if (err)
abort ();
if (omp_target_disassociate_ptr (q, d) != 0)
abort ();
}
omp_target_free (p, d);
return 0;
}
__attribute__((noinline, noclone)) void
foo (S<C, I, L, UCR, CAR, SH, IPR> s)
{
int d = omp_get_default_device ();
int id = omp_get_initial_device ();
int sep = 1;
if (d < 0 || d >= omp_get_num_devices ())
d = id;
int err;
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
s.a = 35; s.b[0] = 36; s.b[1] = 37;
s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
sep = 0;
}
if (err) abort ();
err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
if (err) abort ();
s.a = 50; s.b[0] = 49; s.b[1] = 48;
s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)
|| omp_target_is_present (&s.c[1], d)
|| omp_target_is_present (s.d, d)
|| omp_target_is_present (&s.d[-2], d)
|| omp_target_is_present (&s.e, d)
|| omp_target_is_present (s.f, d)
|| omp_target_is_present (&s.g[1], d)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d)
|| !omp_target_is_present (&s.e, d)
|| !omp_target_is_present (s.f, d)
|| !omp_target_is_present (&s.g[1], d)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
s.a = 17; s.b[0] = 18; s.b[1] = 19;
s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
}
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)
|| omp_target_is_present (&s.c[1], d)
|| omp_target_is_present (s.d, d)
|| omp_target_is_present (&s.d[-2], d)
|| omp_target_is_present (&s.e, d)
|| omp_target_is_present (s.f, d)
|| omp_target_is_present (&s.g[1], d)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
if (err) abort ();
err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
if (err) abort ();
s.a = 33; s.b[0] = 34; s.b[1] = 35;
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d)
|| !omp_target_is_present (&s.e, d)
|| !omp_target_is_present (s.f, d)
|| !omp_target_is_present (&s.g[1], d)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
s.a = 49; s.b[0] = 48; s.b[1] = 47;
s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
}
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d)
|| !omp_target_is_present (&s.e, d)
|| !omp_target_is_present (s.f, d)
|| !omp_target_is_present (&s.g[1], d)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)
|| omp_target_is_present (&s.c[1], d)
|| omp_target_is_present (s.d, d)
|| omp_target_is_present (&s.d[-2], d)
|| omp_target_is_present (&s.e, d)
|| omp_target_is_present (s.f, d)
|| omp_target_is_present (&s.g[1], d)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
if (err) abort ();
err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
if (err) abort ();
}
/// performs the same actions as data_begin in case arg_num is
/// non-zero and initiates run of offloaded region on target platform;
/// if arg_num is non-zero after the region execution is done it also
/// performs the same action as data_update and data_end aboveThe
/// following types are used; this function return 0 if it was able to
/// transfer the execution to a target and an int different from zero
/// otherwise
static int target(int32_t device_id, void *host_ptr, int32_t arg_num,
void** args_base, void **args, int64_t *arg_sizes, int32_t *arg_types,
int32_t team_num, int32_t thread_limit,
int IsTeamConstruct, int IsConstrDestrRecursiveCall)
{
DP("Entering target region with entry point %016lx and device Id %d\n",
(Elf64_Addr)host_ptr, device_id);
if (device_id == OFFLOAD_DEVICE_DEFAULT) {
device_id = omp_get_default_device();
}
// got a new constructor/destructor?
if (device_id == OFFLOAD_DEVICE_CONSTRUCTOR ||
device_id == OFFLOAD_DEVICE_DESTRUCTOR) {
DP("Got a constructor/destructor\n");
for(unsigned D=0; D<Devices.size(); D++) {
DeviceTy & Device = Devices[D];
DP("device %d: enqueue constr/destr\n", D);
Device.PendingConstrDestrHostPtrList.push_back(host_ptr);
}
DP("Done with constructor/destructor\n");
return OFFLOAD_SUCCESS;
}
// No devices available?
if (! (device_id>=0 && (size_t)device_id<Devices.size())) {
DP("Device ID %d does not have a matching RTL.\n", device_id);
return OFFLOAD_FAIL;
}
// Get device info
DeviceTy & Device = Devices[device_id];
DP("Is the device %d (local is %d) initialized? %d\n",
device_id, Device.RTLDeviceID, (int)Device.IsInit);
// Init the device if not done before
if (!Device.IsInit){
assert(! IsConstrDestrRecursiveCall && "constr & destr should not init RT");
if (Device.init() != OFFLOAD_SUCCESS) {
DP("failed to init device %d\n", device_id);
return OFFLOAD_FAIL;
}
}
if (! IsConstrDestrRecursiveCall && ! Device.PendingConstrDestrHostPtrList.empty()) {
DP("has pending constr/destr... call now\n");
for (std::list<void *>::iterator
ii=Device.PendingConstrDestrHostPtrList.begin(),
ie=Device.PendingConstrDestrHostPtrList.end(); ii!=ie ; ++ii) {
void *ConstrDestrHostPtr = *ii;
int rc = target(device_id, ConstrDestrHostPtr, 0, NULL, NULL, NULL, NULL, 1, 1,
true /*team*/, true /*recursive*/);
if (rc != OFFLOAD_SUCCESS) {
DP("failed to run constr/destr... enqueue it\n");
return OFFLOAD_FAIL;
}
}
DP("done with pending constr/destr\n");
Device.PendingConstrDestrHostPtrList.clear();
}
// Find the table information in the map or look it up in the translation tables
TableMap *TM = 0;
HostPtrToTableMapTy::iterator TableMapIt = HostPtrToTableMap.find(host_ptr);
if (TableMapIt == HostPtrToTableMap.end()){
// We don't have a map. So search all the registered libraries
for (HostEntriesBeginToTransTableTy::iterator
ii = HostEntriesBeginToTransTable.begin(),
ie = HostEntriesBeginToTransTable.end(); !TM && ii!=ie ; ++ii){
// get the translation table (which contains all the good info)
TranslationTable *TransTable = &ii->second;
// iterate over all the host table entries to see if we can locate the host_ptr
__tgt_offload_entry *begin = TransTable->HostTable.EntriesBegin;
__tgt_offload_entry *end = TransTable->HostTable.EntriesEnd;
__tgt_offload_entry *cur = begin;
for (uint32_t i=0; cur < end; ++cur, ++i) {
if (cur->addr != host_ptr)
continue;
// we got a match, now fill the HostPtrToTableMap so that we
// may avoid this search next time.
TM = &HostPtrToTableMap[host_ptr];
TM->Table = TransTable;
TM->Index = i;
break;
}
}
} else {
TM = &TableMapIt->second;
}
// No map for this host pointer found!
if (!TM){
DP("Host ptr %016lx does not have a matching target pointer.\n",
(Elf64_Addr)host_ptr);
return OFFLOAD_FAIL;
}
// get target table
assert(TM->Table->TargetsTable.size() > (size_t)device_id
&& "Not expecting a device ID outside the tables bounds!");
__tgt_target_table *TargetTable = TM->Table->TargetsTable[device_id];
// if first call, need to move the data
if (!TargetTable)
{
// 1) get image
assert(TM->Table->TargetsImages.size() > (size_t)device_id
&& "Not expecting a device ID outside the tables bounds!");
__tgt_device_image *img = TM->Table->TargetsImages[device_id];
if (!img){
DP("No image loaded for device id %d.\n", device_id);
return OFFLOAD_FAIL;
}
// 2) load image into the target table
TargetTable = TM->Table->TargetsTable[device_id] = Device.load_binary(img);
// Unable to get table for this image: invalidate image and fail
if (!TargetTable){
DP("Unable to generate entries table for device id %d.\n", device_id);
TM->Table->TargetsImages[device_id] = 0;
return OFFLOAD_FAIL;
}
// Verify if the two tables sizes match
size_t hsize = TM->Table->HostTable.EntriesEnd -
TM->Table->HostTable.EntriesBegin;
size_t tsize = TargetTable->EntriesEnd - TargetTable->EntriesBegin;
// Invalid image for this host entries!
if (hsize != tsize){
DP("Host and Target tables mismatch for device id %d [%lx != %lx].\n",
device_id, hsize, tsize);
TM->Table->TargetsImages[device_id] = 0;
TM->Table->TargetsTable[device_id] = 0;
return OFFLOAD_FAIL;
}
assert(TM->Index < hsize &&
"Not expecting index greater than the table size");
// process global data that needs to be mapped
__tgt_target_table *HostTable = &TM->Table->HostTable;
for(__tgt_offload_entry *CurrDeviceEntry = TargetTable->EntriesBegin,
*CurrHostEntry = HostTable->EntriesBegin,
*EntryDeviceEnd = TargetTable->EntriesEnd;
CurrDeviceEntry != EntryDeviceEnd;
CurrDeviceEntry++, CurrHostEntry++) {
if (CurrDeviceEntry->size != 0) {
// has data
assert(CurrDeviceEntry->size == CurrHostEntry->size && "data size mismatch");
long IsLast;
assert(Device.getTgtPtrBegin(CurrHostEntry->addr, CurrHostEntry->size,
IsLast, false) == NULL && "data in declared target should not be already mapped");
// add entry to map
DP("add mapping from host 0x%llx to 0x%llx with size %lld\n\n",
(unsigned long long) CurrHostEntry->addr,
(unsigned long long)CurrDeviceEntry->addr,
(unsigned long long)CurrDeviceEntry->size);
Device.HostDataToTargetMap.push_front(HostDataToTargetTy(
(long)CurrHostEntry->addr, (long)CurrHostEntry->addr,
(long)CurrHostEntry->addr + CurrHostEntry->size,
(long)CurrDeviceEntry->addr,
(long)CurrDeviceEntry->addr+CurrDeviceEntry->size));
}
}
}
//Move data to device
target_data_begin(Device, arg_num, args_base, args, arg_sizes, arg_types);
std::vector<void*> tgt_args;
for(int32_t i=0; i<arg_num; ++i){
if (arg_types[i] & tgt_map_extra)
continue;
void * HstPtrBegin = args[i];
void * HstPtrBase = args_base[i];
void *TgtPtrBase;
long IsLast; // unused
if (arg_types[i] & tgt_map_pointer) {
DP("Obtaining target argument from host pointer %016lx to object %016lx \n",
(long)HstPtrBase, (long)HstPtrBegin);
void * TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast, false);
TgtPtrBase = TgtPtrBegin; // no offset for ptrs
} else {
DP("Obtaining target argument from host pointer %016lx\n",
(long)HstPtrBegin);
void * TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast, false);
assert(TgtPtrBegin && "NULL argument for hst ptr");
uint64_t PtrDelta = (uint64_t) HstPtrBegin - (uint64_t)HstPtrBase;
TgtPtrBase = (void *)((uint64_t) TgtPtrBegin - PtrDelta);
}
tgt_args.push_back(TgtPtrBase);
}
//Launch device execution
int rc;
DP("Launching target execution with pointer %016lx (index=%d).\n",
(Elf64_Addr)TargetTable->EntriesBegin[TM->Index].addr, TM->Index);
if (IsTeamConstruct) {
rc = Device.run_team_region(TargetTable->EntriesBegin[TM->Index].addr,
&tgt_args[0], tgt_args.size(), team_num, thread_limit);
} else {
rc = Device.run_region(TargetTable->EntriesBegin[TM->Index].addr,
&tgt_args[0], tgt_args.size());
}
if (rc)
return OFFLOAD_FAIL;
//Move data from device
target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types);
return OFFLOAD_SUCCESS;
}
int main(void){
#if CHECK
check_offloading();
#endif
/*
* Default device
*/
printf("Is%s initial device\n", omp_is_initial_device() ? "" : " not");
printf("Initial device: %d\n", omp_get_initial_device());
omp_set_default_device(1);
printf("Default device before task: %d\n", omp_get_default_device());
#pragma omp task
{
printf("Default device inside task: %d\n", omp_get_default_device());
omp_set_default_device(2);
printf("Default device inside task after resetting: %d\n",
omp_get_default_device());
}
#pragma omp taskwait
printf("Default device outside task: %d\n", omp_get_default_device());
// default device can set to whatever, if target fails, it goes to the host
const int default_device = 0;
omp_set_default_device(default_device);
// default device for omp target call MUST be >= 0 and <omp_get_num_devices() or
// the initial device. So when there are no devices, it must be the initial device
int default_device_omp_target_call = default_device;
if (omp_get_num_devices() == 0) {
default_device_omp_target_call = omp_get_initial_device();
}
#if DEBUG
printf("test on machine with %d devices\n", omp_get_num_devices());
#endif
/*
* Target alloc & target memcpy
*/
double A[N], B[N], C[N], D[N], E[N];
double *pA, *pB, *pC, *pD, *pE;
// map ptrs
pA = &A[0];
pB = &B[0];
pC = &C[0];
pD = &D[0];
pE = &E[0];
INIT();
pA = pA - 10;
pC = pC - 20;
pD = pD - 30;
void *device_A = omp_target_alloc(N*sizeof(double), default_device_omp_target_call);
void *device_C = omp_target_alloc(N*sizeof(double), default_device_omp_target_call);
void *device_D = omp_target_alloc(N*sizeof(double), default_device_omp_target_call);
double *dpA = (double *) device_A - 100;
double *dpC = (double *) device_C - 200;
double *dpD = (double *) device_D - 300;
printf("omp_target_alloc %s\n", device_A && device_C && device_D ?
"succeeded" : "failed");
omp_target_memcpy(dpC, pC, N*sizeof(double), 200*sizeof(double),
20*sizeof(double), default_device_omp_target_call, omp_get_initial_device());
omp_target_memcpy(dpD, pD, N*sizeof(double), 300*sizeof(double),
30*sizeof(double), default_device_omp_target_call, omp_get_initial_device());
#pragma omp target is_device_ptr(dpA, dpC, dpD) device(default_device)
{
#pragma omp parallel for schedule(static,1)
for (int i = 0; i < 992; i++)
dpA[i+100] = dpC[i+200] + dpD[i+300] + 1;
}
omp_target_memcpy(pA, dpA, N*sizeof(double), 10*sizeof(double),
100*sizeof(double), omp_get_initial_device(), default_device_omp_target_call);
int fail = 0;
VERIFY(0, N, A[i], (double)(i+2));
if (fail) {
printf ("Test omp_target_memcpy: Failed\n");
} else {
printf ("Test omp_target_memcpy: Succeeded\n");
}
/*
* target_is_present and target_associate/disassociate_ptr
*/
INIT();
if (offloading_disabled()) {
// If offloading is disabled just recreate the messages so that this can
// also be tested with no device.
printf("C is not present, associating it...\n");
printf("omp_target_associate_ptr C %s\n", 1 ? "succeeded" : "failed");
} else if (!omp_target_is_present(C, default_device_omp_target_call)) {
printf("C is not present, associating it...\n");
int rc = omp_target_associate_ptr(C, dpC, N*sizeof(double),
200*sizeof(double), default_device_omp_target_call);
printf("omp_target_associate_ptr C %s\n", !rc ? "succeeded" : "failed");
}
if (offloading_disabled()) {
// If offloading is disabled just recreate the messages so that this can
// also be tested with no device.
printf("D is not present, associating it...\n");
printf("omp_target_associate_ptr D %s\n", 1 ? "succeeded" : "failed");
} else if (!omp_target_is_present(D, default_device_omp_target_call)) {
printf("D is not present, associating it...\n");
int rc = omp_target_associate_ptr(D, dpD, N*sizeof(double),
300*sizeof(double), default_device_omp_target_call);
printf("omp_target_associate_ptr D %s\n", !rc ? "succeeded" : "failed");
}
#pragma omp target data map(from: C, D) device(default_device)
{
printf("Inside target data: A is%s present\n",
(omp_target_is_present(A, default_device_omp_target_call) && !offloading_disabled()) ? "" : " not");
printf("Inside target data: C is%s present\n",
omp_target_is_present(C, default_device_omp_target_call) ? "" : " not");
printf("Inside target data: D is%s present\n",
omp_target_is_present(D, default_device_omp_target_call) ? "" : " not");
// C and D are mapped "from", so there is no copy from host to device.
// If the association was successful, their corresponding device arrays
// are already populated from previous omp_target_memcpy with the correct
// values and the following target for-loop must yield the correct results.
#pragma omp target map(from: A) device(default_device)
{
#pragma omp parallel for schedule(static,1)
for (int i = 0; i < 992; i++)
A[i] = C[i] + D[i] + 1;
}
}
if (offloading_disabled()) {
printf("C is present, disassociating it...\n");
printf("omp_target_disassociate_ptr C %s\n", 1 ? "succeeded" : "failed");
} else if (omp_target_is_present(C, default_device_omp_target_call)) {
printf("C is present, disassociating it...\n");
int rc = omp_target_disassociate_ptr(C, default_device_omp_target_call);
printf("omp_target_disassociate_ptr C %s\n", !rc ? "succeeded" : "failed");
}
if (offloading_disabled()) {
printf("D is present, disassociating it...\n");
printf("omp_target_disassociate_ptr D %s\n", 1 ? "succeeded" : "failed");
} else if (omp_target_is_present(D, default_device_omp_target_call)) {
printf("D is present, disassociating it...\n");
int rc = omp_target_disassociate_ptr(D, default_device_omp_target_call);
printf("omp_target_disassociate_ptr D %s\n", !rc ? "succeeded" : "failed");
}
fail = 0;
VERIFY(0, N, A[i], (double)(i+2));
if (fail) {
printf ("Test omp_target_associate_ptr: Failed\n");
} else {
printf ("Test omp_target_associate_ptr: Succeeded\n");
}
omp_target_free(device_A, default_device_omp_target_call);
omp_target_free(device_C, default_device_omp_target_call);
omp_target_free(device_D, default_device_omp_target_call);
return 0;
}
void
foo (int f)
{
int d = f ? omp_get_num_devices () : omp_get_default_device ();
int h = 5;
#pragma omp target device (d)
if (omp_get_level () != 0)
abort ();
#pragma omp target if (v > 1)
if (omp_get_level () != 0 || !omp_is_initial_device ())
abort ();
#pragma omp target device (d) if (v > 1)
if (omp_get_level () != 0 || !omp_is_initial_device ())
abort ();
#pragma omp target if (v <= 1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
abort ();
#pragma omp target device (d) if (v <= 1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
abort ();
#pragma omp target if (0)
if (omp_get_level () != 0 || !omp_is_initial_device ())
abort ();
#pragma omp target device (d) if (0)
if (omp_get_level () != 0 || !omp_is_initial_device ())
abort ();
#pragma omp target if (1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
abort ();
#pragma omp target device (d) if (1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
abort ();
#pragma omp target data device (d) map (to: h)
{
#pragma omp target device (d)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5)
abort ();
#pragma omp target update device (d) from (h)
}
#pragma omp target data if (v > 1) map (to: h)
{
#pragma omp target if (v > 1)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6)
abort ();
#pragma omp target update if (v > 1) from (h)
}
#pragma omp target data device (d) if (v > 1) map (to: h)
{
#pragma omp target device (d) if (v > 1)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7)
abort ();
#pragma omp target update device (d) if (v > 1) from (h)
}
#pragma omp target data if (v <= 1) map (to: h)
{
#pragma omp target if (v <= 1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 8)
abort ();
#pragma omp target update if (v <= 1) from (h)
}
#pragma omp target data device (d) if (v <= 1) map (to: h)
{
#pragma omp target device (d) if (v <= 1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9)
abort ();
#pragma omp target update device (d) if (v <= 1) from (h)
}
#pragma omp target data if (0) map (to: h)
{
#pragma omp target if (0)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10)
abort ();
#pragma omp target update if (0) from (h)
}
#pragma omp target data device (d) if (0) map (to: h)
{
#pragma omp target device (d) if (0)
if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11)
abort ();
#pragma omp target update device (d) if (0) from (h)
}
#pragma omp target data if (1) map (to: h)
{
#pragma omp target if (1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 12)
abort ();
#pragma omp target update if (1) from (h)
}
#pragma omp target data device (d) if (1) map (to: h)
{
#pragma omp target device (d) if (1)
if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13)
abort ();
#pragma omp target update device (d) if (1) from (h)
}
if (h != 14)
abort ();
}