diff --git a/include/camp/resource.hpp b/include/camp/resource.hpp
index 8a15c7cf..d6ddffc6 100644
--- a/include/camp/resource.hpp
+++ b/include/camp/resource.hpp
@@ -109,15 +109,15 @@ namespace resources
       }
 
       template <typename T>
-      T* allocate(size_t size, MemoryAccess ma = MemoryAccess::Device)
+      T* allocate(size_t n, MemoryAccess ma = MemoryAccess::Device)
       {
-        if (size == 0) {
+        if (n == 0) {
           return nullptr;
         }
         if (!m_value) {
           ::camp::throw_re("Empty Resource type allocate call.");
         }
-        return (T*)m_value->allocate(size * sizeof(T), ma);
+        return (T*)m_value->allocate(n * sizeof(T), ma);
       }
 
       void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device)
diff --git a/include/camp/resource/cuda.hpp b/include/camp/resource/cuda.hpp
index dbb9ebf8..9e7e20ce 100644
--- a/include/camp/resource/cuda.hpp
+++ b/include/camp/resource/cuda.hpp
@@ -184,14 +184,14 @@ namespace resources
         cudaError_t status = cudaPointerGetAttributes(&a, p);
         if (status == cudaSuccess) {
           switch (a.type) {
-            case cudaMemoryTypeUnregistered:
-              return MemoryAccess::Unknown;
             case cudaMemoryTypeHost:
               return MemoryAccess::Pinned;
             case cudaMemoryTypeDevice:
               return MemoryAccess::Device;
             case cudaMemoryTypeManaged:
               return MemoryAccess::Managed;
+            default:
+              return MemoryAccess::Unknown;
           }
         }
         ::camp::throw_re("invalid pointer detected");
@@ -272,44 +272,55 @@ namespace resources
 
       // Memory
       template <typename T>
-      T* allocate(size_t size, MemoryAccess ma = MemoryAccess::Device)
+      T* allocate(size_t n, MemoryAccess ma = MemoryAccess::Device)
       {
+        if (n == 0) {
+          return nullptr;
+        }
         T* ret = nullptr;
-        if (size > 0) {
-          auto d{device_guard(device)};
-          switch (ma) {
-            case MemoryAccess::Unknown:
-            case MemoryAccess::Device:
-              CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMalloc,
-                                             (void**)&ret,
-                                             sizeof(T) * size);
-              break;
-            case MemoryAccess::Pinned:
-              // TODO: do a test here for whether managed is *actually* shared
-              // so we can use the better performing memory
-              CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMallocHost,
-                                             (void**)&ret,
-                                             sizeof(T) * size);
-              break;
-            case MemoryAccess::Managed:
-              CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMallocManaged,
-                                             (void**)&ret,
-                                             sizeof(T) * size);
-              break;
-          }
+        auto d{device_guard(device)};
+        switch (ma) {
+          case MemoryAccess::Device:
+            CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMalloc,
+                                           (void**)&ret,
+                                           sizeof(T) * n);
+            break;
+          case MemoryAccess::Pinned:
+            // TODO: do a test here for whether managed is *actually* shared
+            // so we can use the better performing memory
+            CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMallocHost,
+                                           (void**)&ret,
+                                           sizeof(T) * n);
+            break;
+          case MemoryAccess::Managed:
+            CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMallocManaged,
+                                           (void**)&ret,
+                                           sizeof(T) * n);
+            break;
+          case MemoryAccess::Unknown:
+            ::camp::throw_re("Unknown memory access type, cannot allocate");
+            break;
         }
         return ret;
       }
 
       void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device)
       {
-        void* p = allocate<char>(size, ma);
-        this->memset(p, 0, size);
-        return p;
+        if (size == 0) {
+          return nullptr;
+        }
+        void* ret = allocate<char>(size, ma);
+        if (ret != nullptr) {
+          this->memset(ret, 0, size);
+        }
+        return ret;
       }
 
       void deallocate(void* p, MemoryAccess ma = MemoryAccess::Unknown)
       {
+        if (p == nullptr) {
+          return;
+        }
         auto d{device_guard(device)};
         if (ma == MemoryAccess::Unknown) {
           ma = get_access_type(p);
@@ -334,19 +345,21 @@ namespace resources
 
       void memcpy(void* dst, const void* src, size_t size)
       {
-        if (size > 0) {
-          auto d{device_guard(device)};
-          CAMP_CUDA_API_INVOKE_AND_CHECK(
-              cudaMemcpyAsync, dst, src, size, cudaMemcpyDefault, stream);
+        if (size == 0) {
+          return;
         }
+        auto d{device_guard(device)};
+        CAMP_CUDA_API_INVOKE_AND_CHECK(
+            cudaMemcpyAsync, dst, src, size, cudaMemcpyDefault, stream);
       }
 
       void memset(void* p, int val, size_t size)
       {
-        if (size > 0) {
-          auto d{device_guard(device)};
-          CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMemsetAsync, p, val, size, stream);
+        if (size == 0) {
+          return;
         }
+        auto d{device_guard(device)};
+        CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMemsetAsync, p, val, size, stream);
       }
 
       cudaStream_t get_stream() const { return stream; }
diff --git a/include/camp/resource/hip.hpp b/include/camp/resource/hip.hpp
index 01f0925c..5aa22261 100644
--- a/include/camp/resource/hip.hpp
+++ b/include/camp/resource/hip.hpp
@@ -184,16 +184,12 @@ namespace resources
         hipPointerAttribute_t a;
         hipError_t status = hipPointerGetAttributes(&a, p);
         if (status == hipSuccess) {
-#if (HIP_VERSION_MAJOR >= 6)
           switch (a.type) {
-#else
-          switch (a.memoryType) {
-#endif
             case hipMemoryTypeHost:
               return MemoryAccess::Pinned;
             case hipMemoryTypeDevice:
               return MemoryAccess::Device;
-            case hipMemoryTypeUnified:
+            case hipMemoryTypeManaged:
               return MemoryAccess::Managed;
             default:
               return MemoryAccess::Unknown;
@@ -274,44 +270,55 @@ namespace resources
 
       // Memory
       template <typename T>
-      T* allocate(size_t size, MemoryAccess ma = MemoryAccess::Device)
+      T* allocate(size_t n, MemoryAccess ma = MemoryAccess::Device)
       {
+        if (n == 0) {
+          return nullptr;
+        }
         T* ret = nullptr;
-        if (size > 0) {
-          auto d{device_guard(device)};
-          switch (ma) {
-            case MemoryAccess::Unknown:
-            case MemoryAccess::Device:
-              CAMP_HIP_API_INVOKE_AND_CHECK(hipMalloc,
-                                            (void**)&ret,
-                                            sizeof(T) * size);
-              break;
-            case MemoryAccess::Pinned:
-              // TODO: do a test here for whether managed is *actually* shared
-              // so we can use the better performing memory
-              CAMP_HIP_API_INVOKE_AND_CHECK(hipHostMalloc,
-                                            (void**)&ret,
-                                            sizeof(T) * size);
-              break;
-            case MemoryAccess::Managed:
-              CAMP_HIP_API_INVOKE_AND_CHECK(hipMallocManaged,
-                                            (void**)&ret,
-                                            sizeof(T) * size);
-              break;
-          }
+        auto d{device_guard(device)};
+        switch (ma) {
+          case MemoryAccess::Device:
+            CAMP_HIP_API_INVOKE_AND_CHECK(hipMalloc,
+                                          (void**)&ret,
+                                          sizeof(T) * n);
+            break;
+          case MemoryAccess::Pinned:
+            // TODO: do a test here for whether managed is *actually* shared
+            // so we can use the better performing memory
+            CAMP_HIP_API_INVOKE_AND_CHECK(hipHostMalloc,
+                                          (void**)&ret,
+                                          sizeof(T) * n);
+            break;
+          case MemoryAccess::Managed:
+            CAMP_HIP_API_INVOKE_AND_CHECK(hipMallocManaged,
+                                          (void**)&ret,
+                                          sizeof(T) * n);
+            break;
+          case MemoryAccess::Unknown:
+            ::camp::throw_re("Unknown memory access type, cannot allocate");
+            break;
         }
         return ret;
       }
 
       void* calloc(size_t size, MemoryAccess ma)
       {
-        void* p = allocate<char>(size, ma);
-        this->memset(p, 0, size);
-        return p;
+        if (size == 0) {
+          return nullptr;
+        }
+        void* ret = allocate<char>(size, ma);
+        if (ret != nullptr) {
+          this->memset(ret, 0, size);
+        }
+        return ret;
       }
 
       void deallocate(void* p, MemoryAccess ma = MemoryAccess::Unknown)
       {
+        if (p == nullptr) {
+          return;
+        }
         auto d{device_guard(device)};
         if (ma == MemoryAccess::Unknown) {
           ma = get_access_type(p);
@@ -336,19 +343,21 @@ namespace resources
 
       void memcpy(void* dst, const void* src, size_t size)
       {
-        if (size > 0) {
-          auto d{device_guard(device)};
-          CAMP_HIP_API_INVOKE_AND_CHECK(
-              hipMemcpyAsync, dst, src, size, hipMemcpyDefault, stream);
+        if (size == 0) {
+          return;
         }
+        auto d{device_guard(device)};
+        CAMP_HIP_API_INVOKE_AND_CHECK(
+            hipMemcpyAsync, dst, src, size, hipMemcpyDefault, stream);
       }
 
       void memset(void* p, int val, size_t size)
       {
-        if (size > 0) {
-          auto d{device_guard(device)};
-          CAMP_HIP_API_INVOKE_AND_CHECK(hipMemsetAsync, p, val, size, stream);
+        if (size == 0) {
+          return;
         }
+        auto d{device_guard(device)};
+        CAMP_HIP_API_INVOKE_AND_CHECK(hipMemsetAsync, p, val, size, stream);
       }
 
       hipStream_t get_stream() const { return stream; }
diff --git a/include/camp/resource/host.hpp b/include/camp/resource/host.hpp
index 83b597f6..7ef3393d 100644
--- a/include/camp/resource/host.hpp
+++ b/include/camp/resource/host.hpp
@@ -114,27 +114,47 @@ namespace resources
       template <typename T>
       T* allocate(size_t n, MemoryAccess = MemoryAccess::Device)
       {
+        if (n == 0) {
+          return nullptr;
+        }
         return (T*)std::malloc(sizeof(T) * n);
       }
 
       void* calloc(size_t size, MemoryAccess = MemoryAccess::Device)
       {
-        void* p = allocate<char>(size);
-        this->memset(p, 0, size);
-        return p;
+        if (size == 0) {
+          return nullptr;
+        }
+        void* ret = allocate<char>(size);
+        if (ret != nullptr) {
+          this->memset(ret, 0, size);
+        }
+        return ret;
       }
 
       void deallocate(void* p, MemoryAccess = MemoryAccess::Device)
       {
+        if (p == nullptr) {
+          return;
+        }
         std::free(p);
       }
 
       void memcpy(void* dst, const void* src, size_t size)
       {
+        if (size == 0) {
+          return;
+        }
         std::memcpy(dst, src, size);
       }
 
-      void memset(void* p, int val, size_t size) { std::memset(p, val, size); }
+      void memset(void* p, int val, size_t size)
+      {
+        if (size == 0) {
+          return;
+        }
+        std::memset(p, val, size);
+      }
 
       /*
        * \brief Compares two (Host) resources to see if they are equal
diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp
index 2d2ff4b0..c5f8a199 100644
--- a/include/camp/resource/omp_target.hpp
+++ b/include/camp/resource/omp_target.hpp
@@ -217,24 +217,35 @@ namespace resources
 
       // Memory
       template <typename T>
-      T* allocate(size_t size, MemoryAccess ma = MemoryAccess::Device)
+      T* allocate(size_t n, MemoryAccess ma = MemoryAccess::Device)
       {
+        if (n == 0) {
+          return nullptr;
+        }
         check_ma(ma);
-        T* ret = static_cast<T*>(omp_target_alloc(sizeof(T) * size, dev));
+        T* ret = static_cast<T*>(omp_target_alloc(sizeof(T) * n, dev));
         register_ptr_dev(ret, dev);
         return ret;
       }
 
       void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device)
       {
+        if (size == 0) {
+          return nullptr;
+        }
         check_ma(ma);
-        void* p = allocate<char>(size);
-        this->memset(p, 0, size);
-        return p;
+        void* ret = allocate<char>(size);
+        if (ret != nullptr) {
+          this->memset(ret, 0, size);
+        }
+        return ret;
       }
 
       void deallocate(void* p, MemoryAccess ma = MemoryAccess::Device)
       {
+        if (p == nullptr) {
+          return;
+        }
         check_ma(ma);
         deregister_ptr_dev(p);
         omp_target_free(p, dev);
@@ -242,6 +253,9 @@ namespace resources
 
       void memcpy(void* dst, const void* src, size_t size)
       {
+        if (size == 0) {
+          return;
+        }
         // this is truly, insanely awful, need to think of something better
         int dd = get_ptr_dev(dst);
         int sd = get_ptr_dev(src);
@@ -251,6 +265,9 @@ namespace resources
 
       void memset(void* p, int val, size_t size)
       {
+        if (size == 0) {
+          return;
+        }
         char* local_addr = addr;
         CAMP_ALLOW_UNUSED_LOCAL(local_addr);
         char* pc = (char*)p;
@@ -263,6 +280,9 @@ namespace resources
 
       void register_ptr_dev(void* p, int device)
       {
+        if (p == nullptr) {
+          return;
+        }
 #pragma omp critical(camp_register_ptr)
         {
           get_dev_register()[p] = device;
@@ -271,6 +291,9 @@ namespace resources
 
       void deregister_ptr_dev(void const* p)
       {
+        if (p == nullptr) {
+          return;
+        }
 #pragma omp critical(camp_register_ptr)
         {
           get_dev_register().erase(p);
@@ -280,6 +303,9 @@ namespace resources
       int get_ptr_dev(void const* p)
       {
         int ret = omp_get_initial_device();
+        if (p == nullptr) {
+          return ret;
+        }
 #pragma omp critical(camp_register_ptr)
         {
           auto it = get_dev_register().find(p);
diff --git a/include/camp/resource/sycl.hpp b/include/camp/resource/sycl.hpp
index 75ee05df..6f32a0fc 100644
--- a/include/camp/resource/sycl.hpp
+++ b/include/camp/resource/sycl.hpp
@@ -335,52 +335,64 @@ namespace resources
 
       // Memory
       template <typename T>
-      T* allocate(size_t size, MemoryAccess ma = MemoryAccess::Device)
+      T* allocate(size_t n, MemoryAccess ma = MemoryAccess::Device)
       {
+        if (n == 0) {
+          return nullptr;
+        }
         T* ret = nullptr;
-        if (size > 0) {
-          ret = sycl::malloc_shared<T>(size, qu);
-          switch (ma) {
-            case MemoryAccess::Unknown:
-            case MemoryAccess::Device:
-              ret = sycl::malloc_device<T>(size, qu);
-              break;
-            case MemoryAccess::Pinned:
-              ret = sycl::malloc_host<T>(size, qu);
-              break;
-            case MemoryAccess::Managed:
-              ret = sycl::malloc_shared<T>(size, qu);
-              break;
-          }
+        switch (ma) {
+          case MemoryAccess::Device:
+            ret = sycl::malloc_device<T>(n, qu);
+            break;
+          case MemoryAccess::Pinned:
+            ret = sycl::malloc_host<T>(n, qu);
+            break;
+          case MemoryAccess::Managed:
+            ret = sycl::malloc_shared<T>(n, qu);
+            break;
+          case MemoryAccess::Unknown:
+            ::camp::throw_re("Unknown memory access type, cannot allocate");
+            break;
         }
         return ret;
       }
 
       void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device)
       {
-        void* p = allocate<char>(size, ma);
-        this->memset(p, 0, size);
-        return p;
+        if (size == 0) {
+          return nullptr;
+        }
+        void* ret = allocate<char>(size, ma);
+        if (ret != nullptr) {
+          this->memset(ret, 0, size);
+        }
+        return ret;
       }
 
       void deallocate(void* p, MemoryAccess ma = MemoryAccess::Device)
       {
+        if (p == nullptr) {
+          return;
+        }
         CAMP_ALLOW_UNUSED_LOCAL(ma);
         sycl::free(p, qu);
       }
 
       void memcpy(void* dst, const void* src, size_t size)
       {
-        if (size > 0) {
-          qu.memcpy(dst, src, size).wait();
+        if (size == 0) {
+          return;
         }
+        qu.memcpy(dst, src, size).wait();
       }
 
       void memset(void* p, int val, size_t size)
       {
-        if (size > 0) {
-          qu.memset(p, val, size).wait();
+        if (size == 0) {
+          return;
         }
+        qu.memset(p, val, size).wait();
       }
 
       // implementation specific
diff --git a/test/resource.cpp b/test/resource.cpp
index a91ec2ad..b4ab0904 100644
--- a/test/resource.cpp
+++ b/test/resource.cpp
@@ -1729,7 +1729,6 @@ TEST(CampResource, MemoryHost) { test_memory_ops<Host>(MemoryAccess::Device); }
 #ifdef CAMP_HAVE_CUDA
 TEST(CampResource, MemoryCuda)
 {
-  test_memory_ops<Cuda>(MemoryAccess::Unknown);
   test_memory_ops<Cuda>(MemoryAccess::Device);
   test_memory_ops<Cuda>(MemoryAccess::Pinned);
   test_memory_ops<Cuda>(MemoryAccess::Managed);
@@ -1743,7 +1742,6 @@ TEST(CampResource, MemoryCuda)
 #ifdef CAMP_HAVE_HIP
 TEST(CampResource, MemoryHip)
 {
-  test_memory_ops<Hip>(MemoryAccess::Unknown);
   test_memory_ops<Hip>(MemoryAccess::Device);
   test_memory_ops<Hip>(MemoryAccess::Pinned);
   test_memory_ops<Hip>(MemoryAccess::Managed);
@@ -1761,7 +1759,6 @@ TEST(CampResource, MemoryOmp) { test_memory_ops<Omp>(MemoryAccess::Device); }
 #ifdef CAMP_HAVE_SYCL
 TEST(CampResource, MemorySycl)
 {
-  test_memory_ops<Sycl>(MemoryAccess::Unknown);
   test_memory_ops<Sycl>(MemoryAccess::Device);
   test_memory_ops<Sycl>(MemoryAccess::Pinned);
   test_memory_ops<Sycl>(MemoryAccess::Managed);