diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp
index 4ef60ba6..9d862e10 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,37 +127,46 @@ namespace resources
       template <typename T>
       T *allocate(size_t size)
       {
-        T *ret = static_cast<T *>(omp_target_alloc(sizeof(T) * size, dev));
-        register_ptr_dev(ret, dev);
-        return ret;
+        T * hostmem = (T*)malloc(sizeof(T)*sizehost);
+
+#pragma omp target enter data map( to: hostmem[0:size] )
+
+        return hostmem;
       }
+
       void *calloc(size_t size)
       {
         void *p = allocate<char>(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);
+        char * pp = (char *)p;
+        CAMP_ALLOW_UNUSED_LOCAL(pp);
+#pragma omp target exit data map( release: pp[:0] )
+        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);
+        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)
+        {
+          omp_target_memcpy(dst, (void *)src, size, 0, 0, ddevice, sdevice);
+        }
       }
+
       void memset(void *p, int val, size_t size)
       {
         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
@@ -161,35 +175,9 @@ namespace resources
         }
       }
 
-      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 <typename always_void_odr_helper = void>
-      std::map<const void *, int> &get_dev_register()
-      {
-        static std::map<const void *, int> dev_register;
-        return dev_register;
-      }
     };
 
   }  // namespace v1