From fd63cff2b4ed0df1d5182f6144ca3bcf6c772cea Mon Sep 17 00:00:00 2001 From: Robert Chen Date: Thu, 18 Jun 2020 19:57:21 -0700 Subject: [PATCH 1/3] Initial attempt to hostify omp resource. --- include/camp/resource/omp_target.hpp | 106 +++++++++++++++++++++++++-- 1 file changed, 100 insertions(+), 6 deletions(-) diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp index 4ef60ba6..34b6dc3f 100644 --- a/include/camp/resource/omp_target.hpp +++ b/include/camp/resource/omp_target.hpp @@ -122,9 +122,16 @@ namespace resources template T *allocate(size_t size) { - T *ret = static_cast(omp_target_alloc(sizeof(T) * size, dev)); - register_ptr_dev(ret, dev); - return ret; + int hdev = omp_get_initial_device(); + T *hostmem = new T[size]; + register_ptr_host(hostmem, hdev); + + T *devmem = static_cast(omp_target_alloc(sizeof(T) * size, dev)); + register_ptr_dev(devmem, dev); + + register_host_devptr(hostmem, devmem); + + return hostmem; } void *calloc(size_t size) { @@ -143,11 +150,41 @@ namespace resources void memcpy(void *dst, const void *src, size_t size) { // this is truly, insanely awful, need to think of something better - int dd = get_ptr_dev(dst); - int sd = get_ptr_dev(src); + //int dd = get_ptr_dev(dst); + //int sd = get_ptr_dev(src); // extra cast due to GCC openmp header bug - omp_target_memcpy(dst, (void *)src, size, 0, 0, dd, sd); + //omp_target_memcpy(dst, (void *)src, size, 0, 0, dd, sd); + + // witness the sadness . . . + // check for src = host, dst = device + auto ith = get_host_register().find(src); + auto itd = get_dev_register().find(dst); + if (ith != get_host_register().end() && itd != get_dev_register().end()) + { + int dd = get_ptr_dev(dst); + int sd = get_ptr_host(src); + // extra cast due to GCC openmp header bug + omp_target_memcpy(dst, (void *)src, size, 0, 0, dd, sd); + } + // check for src = device, dst = host + else + { + auto itth = get_host_register().find(dst); + auto ittd = get_dev_register().find(src); + if (itth != get_host_register().end() && ittd != get_dev_register().end()) + { + int dd = get_ptr_dev(src); + int sd = get_ptr_host(dst); + // extra cast due to GCC openmp header bug + omp_target_memcpy(dst, (void *)src, size, 0, 0, dd, sd); + } + else + { + printf( "TROUBLE TROUBLE TROUBLE TROUBLE TROUBLE\n" ); + } + } } + void memset(void *p, int val, size_t size) { char *local_addr = addr; @@ -161,6 +198,46 @@ namespace resources } } + void register_ptr_host(void *p, int device) + { + get_host_register()[p] = device; + } + int get_ptr_host(void const *p) + { + int ret = omp_get_initial_device(); + auto it = get_host_register().find(p); + if (it != get_host_register().end()) + { + ret = it->second; + } + return ret; + } + + void register_host_devptr(void *h, void *d) + { + get_host_devptr()[h] = d; + } + const void * obtain_devptr_withhost(void *h) + { + const void * ret = nullptr; + auto it = get_host_devptr().find(h); + if (it != get_host_devptr().end()) + { + ret = it->second; + } + + if (ret != nullptr) + { + return ret; + } + else + { + //fflush(); + abort(); + return nullptr; + } + } + void register_ptr_dev(void *p, int device) { #pragma omp critical(camp_register_ptr) @@ -187,9 +264,26 @@ namespace resources template std::map &get_dev_register() { + // key: device ptr, value: device num static std::map dev_register; return dev_register; } + + template + std::map &get_host_register() + { + // key: host ptr, value: device num + static std::map host_register; + return host_register; + } + + template + std::map &get_host_devptr() + { + // key: host ptr, value: device ptr + static std::map host_devptr; + return host_devptr; + } }; } // namespace v1 From ffbada76754825a4ba41967db098736b50b8b011 Mon Sep 17 00:00:00 2001 From: Robert Chen Date: Tue, 7 Jul 2020 13:38:25 -0700 Subject: [PATCH 2/3] Host allocation and map design. XL memcpy does not work, clang-ibm segfault in kernel. --- include/camp/resource/omp_target.hpp | 145 ++++----------------------- 1 file changed, 18 insertions(+), 127 deletions(-) diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp index 34b6dc3f..410ca32c 100644 --- a/include/camp/resource/omp_target.hpp +++ b/include/camp/resource/omp_target.hpp @@ -27,6 +27,11 @@ namespace resources inline namespace v1 { + struct hdmem + { + void * host_arr[10]; + }; + class OmpEvent { public: @@ -122,66 +127,34 @@ namespace resources template T *allocate(size_t size) { - int hdev = omp_get_initial_device(); - T *hostmem = new T[size]; - register_ptr_host(hostmem, hdev); - - T *devmem = static_cast(omp_target_alloc(sizeof(T) * size, dev)); - register_ptr_dev(devmem, dev); + T * hostmem = (T*)malloc(sizeof(T)*sizehost); - register_host_devptr(hostmem, devmem); +#pragma omp target enter data map( to: hostmem[0:size] ) return hostmem; } + void *calloc(size_t size) { void *p = allocate(size); this->memset(p, 0, size); return p; } + void deallocate(void *p) { -#pragma omp critical(camp_register_ptr) - { - get_dev_register().erase(p); - } - omp_target_free(p, dev); + free(p); } + void memcpy(void *dst, const void *src, size_t size) { - // this is truly, insanely awful, need to think of something better - //int dd = get_ptr_dev(dst); - //int sd = get_ptr_dev(src); - // extra cast due to GCC openmp header bug - //omp_target_memcpy(dst, (void *)src, size, 0, 0, dd, sd); - - // witness the sadness . . . - // check for src = host, dst = device - auto ith = get_host_register().find(src); - auto itd = get_dev_register().find(dst); - if (ith != get_host_register().end() && itd != get_dev_register().end()) + int initdev = omp_get_initial_device(); + int sdevice = omp_target_is_present( (void *)src, dev ) ? dev : initdev; + int ddevice = omp_target_is_present( (void *)dst, dev ) ? dev : initdev; + #pragma omp target data if(sdevice != initdev) device(sdevice) use_device_ptr(src) + #pragma omp target data if(ddevice != initdev) device(ddevice) use_device_ptr(dst) { - int dd = get_ptr_dev(dst); - int sd = get_ptr_host(src); - // extra cast due to GCC openmp header bug - omp_target_memcpy(dst, (void *)src, size, 0, 0, dd, sd); - } - // check for src = device, dst = host - else - { - auto itth = get_host_register().find(dst); - auto ittd = get_dev_register().find(src); - if (itth != get_host_register().end() && ittd != get_dev_register().end()) - { - int dd = get_ptr_dev(src); - int sd = get_ptr_host(dst); - // extra cast due to GCC openmp header bug - omp_target_memcpy(dst, (void *)src, size, 0, 0, dd, sd); - } - else - { - printf( "TROUBLE TROUBLE TROUBLE TROUBLE TROUBLE\n" ); - } + omp_target_memcpy(dst, (void *)src, size, 0, 0, ddevice, sdevice); } } @@ -190,6 +163,7 @@ namespace resources char *local_addr = addr; CAMP_ALLOW_UNUSED_LOCAL(local_addr); char *pc = (char *)p; +#pragma omp target data use_device_ptr(pc) #pragma omp target teams distribute parallel for device(dev) \ depend(inout \ : local_addr[0]) is_device_ptr(pc) nowait @@ -198,92 +172,9 @@ namespace resources } } - void register_ptr_host(void *p, int device) - { - get_host_register()[p] = device; - } - int get_ptr_host(void const *p) - { - int ret = omp_get_initial_device(); - auto it = get_host_register().find(p); - if (it != get_host_register().end()) - { - ret = it->second; - } - return ret; - } - - void register_host_devptr(void *h, void *d) - { - get_host_devptr()[h] = d; - } - const void * obtain_devptr_withhost(void *h) - { - const void * ret = nullptr; - auto it = get_host_devptr().find(h); - if (it != get_host_devptr().end()) - { - ret = it->second; - } - - if (ret != nullptr) - { - return ret; - } - else - { - //fflush(); - abort(); - return nullptr; - } - } - - void register_ptr_dev(void *p, int device) - { -#pragma omp critical(camp_register_ptr) - { - get_dev_register()[p] = device; - } - } - int get_ptr_dev(void const *p) - { - int ret = omp_get_initial_device(); -#pragma omp critical(camp_register_ptr) - { - auto it = get_dev_register().find(p); - if (it != get_dev_register().end()) { - ret = it->second; - } - } - return ret; - } - private: char *addr; int dev; - template - std::map &get_dev_register() - { - // key: device ptr, value: device num - static std::map dev_register; - return dev_register; - } - - template - std::map &get_host_register() - { - // key: host ptr, value: device num - static std::map host_register; - return host_register; - } - - template - std::map &get_host_devptr() - { - // key: host ptr, value: device ptr - static std::map host_devptr; - return host_devptr; - } }; } // namespace v1 From 657b713fb512f03d8ecb2805819281c3b71a814e Mon Sep 17 00:00:00 2001 From: Robert Chen Date: Tue, 7 Jul 2020 14:22:18 -0700 Subject: [PATCH 3/3] Add exit data to deallocate. --- include/camp/resource/omp_target.hpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp index 410ca32c..9d862e10 100644 --- a/include/camp/resource/omp_target.hpp +++ b/include/camp/resource/omp_target.hpp @@ -143,6 +143,9 @@ namespace resources void deallocate(void *p) { + char * pp = (char *)p; + CAMP_ALLOW_UNUSED_LOCAL(pp); +#pragma omp target exit data map( release: pp[:0] ) free(p); }