From 62596bc0f9a527b43f64e4575955158ff7c94104 Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Tue, 14 Jul 2026 10:50:54 -0700 Subject: [PATCH 1/9] Fix resource allocators with 0 size --- include/camp/resource/cuda.hpp | 12 +++++-- include/camp/resource/hip.hpp | 12 +++++-- include/camp/resource/omp_target.hpp | 47 ++++++++++++++++++---------- include/camp/resource/sycl.hpp | 12 +++++-- 4 files changed, 57 insertions(+), 26 deletions(-) diff --git a/include/camp/resource/cuda.hpp b/include/camp/resource/cuda.hpp index dbb9ebf8..22ac52f6 100644 --- a/include/camp/resource/cuda.hpp +++ b/include/camp/resource/cuda.hpp @@ -303,13 +303,19 @@ namespace resources void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device) { - void* p = allocate(size, ma); - this->memset(p, 0, size); - return p; + T* ret = nullptr; + if (size > 0) { + ret = allocate(size, ma); + 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); diff --git a/include/camp/resource/hip.hpp b/include/camp/resource/hip.hpp index 01f0925c..4398cec7 100644 --- a/include/camp/resource/hip.hpp +++ b/include/camp/resource/hip.hpp @@ -305,13 +305,19 @@ namespace resources void* calloc(size_t size, MemoryAccess ma) { - void* p = allocate(size, ma); - this->memset(p, 0, size); - return p; + T* ret = nullptr; + if (size > 0) { + ret = allocate(size, ma); + 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); diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp index 2d2ff4b0..9316e138 100644 --- a/include/camp/resource/omp_target.hpp +++ b/include/camp/resource/omp_target.hpp @@ -219,22 +219,31 @@ namespace resources template T* allocate(size_t size, MemoryAccess ma = MemoryAccess::Device) { - check_ma(ma); - T* ret = static_cast(omp_target_alloc(sizeof(T) * size, dev)); - register_ptr_dev(ret, dev); + T* ret = nullptr; + if (size > 0) { + check_ma(ma); + ret = static_cast(omp_target_alloc(sizeof(T) * size, dev)); + register_ptr_dev(ret, dev); + } return ret; } void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device) { - check_ma(ma); - void* p = allocate(size); - this->memset(p, 0, size); - return p; + void* ret = nullptr; + if (size > 0) { + check_ma(ma); + void* ret = allocate(size); + 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,22 +251,26 @@ 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); - // extra cast due to GCC openmp header bug - omp_target_memcpy(dst, (void*)src, size, 0, 0, dd, sd); + if (size > 0) { + // 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); + } } void memset(void* p, int val, size_t size) { - char* local_addr = addr; - CAMP_ALLOW_UNUSED_LOCAL(local_addr); - char* pc = (char*)p; + if (size > 0) { + char* local_addr = addr; + CAMP_ALLOW_UNUSED_LOCAL(local_addr); + char* pc = (char*)p; #pragma omp target teams distribute parallel for device(dev) \ depend(inout : local_addr[0]) is_device_ptr(pc) nowait - for (size_t i = 0; i < size; ++i) { - pc[i] = val; + for (size_t i = 0; i < size; ++i) { + pc[i] = val; + } } } diff --git a/include/camp/resource/sycl.hpp b/include/camp/resource/sycl.hpp index 75ee05df..4e81fd71 100644 --- a/include/camp/resource/sycl.hpp +++ b/include/camp/resource/sycl.hpp @@ -358,13 +358,19 @@ namespace resources void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device) { - void* p = allocate(size, ma); - this->memset(p, 0, size); - return p; + T* ret = nullptr; + if (size > 0) { + ret = allocate(size, ma); + 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); } From 485b3e4299e3c0d5d31c6efe5af88bcf46ffd8c4 Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Tue, 14 Jul 2026 18:00:01 -0700 Subject: [PATCH 2/9] Remove MemoryAccess::Unknown tests They are failing as we don't allow deallocating unknown memory accesses. --- test/resource.cpp | 3 --- 1 file changed, 3 deletions(-) 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(MemoryAccess::Device); } #ifdef CAMP_HAVE_CUDA TEST(CampResource, MemoryCuda) { - test_memory_ops(MemoryAccess::Unknown); test_memory_ops(MemoryAccess::Device); test_memory_ops(MemoryAccess::Pinned); test_memory_ops(MemoryAccess::Managed); @@ -1743,7 +1742,6 @@ TEST(CampResource, MemoryCuda) #ifdef CAMP_HAVE_HIP TEST(CampResource, MemoryHip) { - test_memory_ops(MemoryAccess::Unknown); test_memory_ops(MemoryAccess::Device); test_memory_ops(MemoryAccess::Pinned); test_memory_ops(MemoryAccess::Managed); @@ -1761,7 +1759,6 @@ TEST(CampResource, MemoryOmp) { test_memory_ops(MemoryAccess::Device); } #ifdef CAMP_HAVE_SYCL TEST(CampResource, MemorySycl) { - test_memory_ops(MemoryAccess::Unknown); test_memory_ops(MemoryAccess::Device); test_memory_ops(MemoryAccess::Pinned); test_memory_ops(MemoryAccess::Managed); From 7f8689f8ed79e445f4845dca15e95fa74523c8f1 Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Tue, 14 Jul 2026 18:00:47 -0700 Subject: [PATCH 3/9] Unify resource memory 0 size handling Also fix leak in Sycl::allocate --- include/camp/resource.hpp | 6 +-- include/camp/resource/cuda.hpp | 71 +++++++++++++++------------- include/camp/resource/hip.hpp | 71 +++++++++++++++------------- include/camp/resource/host.hpp | 20 +++++++- include/camp/resource/omp_target.hpp | 50 ++++++++++---------- include/camp/resource/sycl.hpp | 48 ++++++++++--------- 6 files changed, 147 insertions(+), 119 deletions(-) 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 - 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 22ac52f6..f62eff2b 100644 --- a/include/camp/resource/cuda.hpp +++ b/include/camp/resource/cuda.hpp @@ -272,42 +272,43 @@ namespace resources // Memory template - 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::Unknown: + 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; } return ret; } void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device) { - T* ret = nullptr; - if (size > 0) { - ret = allocate(size, ma); - this->memset(ret, 0, size); + if (size == 0) { + return nullptr; } + void* ret = allocate(size, ma); + this->memset(ret, 0, size); return ret; } @@ -340,19 +341,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 4398cec7..570f2c83 100644 --- a/include/camp/resource/hip.hpp +++ b/include/camp/resource/hip.hpp @@ -274,42 +274,43 @@ namespace resources // Memory template - 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::Unknown: + 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; } return ret; } void* calloc(size_t size, MemoryAccess ma) { - T* ret = nullptr; - if (size > 0) { - ret = allocate(size, ma); - this->memset(ret, 0, size); + if (size == 0) { + return nullptr; } + void* ret = allocate(size, ma); + this->memset(ret, 0, size); return ret; } @@ -342,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..26230486 100644 --- a/include/camp/resource/host.hpp +++ b/include/camp/resource/host.hpp @@ -114,11 +114,17 @@ namespace resources template 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) { + if (size == 0) { + return nullptr; + } void* p = allocate(size); this->memset(p, 0, size); return p; @@ -126,15 +132,27 @@ namespace resources 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 9316e138..addfb198 100644 --- a/include/camp/resource/omp_target.hpp +++ b/include/camp/resource/omp_target.hpp @@ -217,25 +217,25 @@ namespace resources // Memory template - T* allocate(size_t size, MemoryAccess ma = MemoryAccess::Device) + T* allocate(size_t n, MemoryAccess ma = MemoryAccess::Device) { - T* ret = nullptr; - if (size > 0) { - check_ma(ma); - ret = static_cast(omp_target_alloc(sizeof(T) * size, dev)); - register_ptr_dev(ret, dev); + if (n == 0) { + return nullptr; } + check_ma(ma); + T* ret = static_cast(omp_target_alloc(sizeof(T) * n, dev)); + register_ptr_dev(ret, dev); return ret; } void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device) { - void* ret = nullptr; - if (size > 0) { - check_ma(ma); - void* ret = allocate(size); - this->memset(ret, 0, size); + if (size == 0) { + return nullptr; } + check_ma(ma); + void* ret = allocate(size); + this->memset(ret, 0, size); return ret; } @@ -251,26 +251,28 @@ namespace resources void memcpy(void* dst, const void* src, size_t size) { - if (size > 0) { - // 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); + 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); + // extra cast due to GCC openmp header bug + omp_target_memcpy(dst, (void*)src, size, 0, 0, dd, sd); } void memset(void* p, int val, size_t size) { - if (size > 0) { - char* local_addr = addr; - CAMP_ALLOW_UNUSED_LOCAL(local_addr); - char* pc = (char*)p; + if (size == 0) { + return; + } + char* local_addr = addr; + CAMP_ALLOW_UNUSED_LOCAL(local_addr); + char* pc = (char*)p; #pragma omp target teams distribute parallel for device(dev) \ depend(inout : local_addr[0]) is_device_ptr(pc) nowait - for (size_t i = 0; i < size; ++i) { - pc[i] = val; - } + for (size_t i = 0; i < size; ++i) { + pc[i] = val; } } diff --git a/include/camp/resource/sycl.hpp b/include/camp/resource/sycl.hpp index 4e81fd71..ab53457b 100644 --- a/include/camp/resource/sycl.hpp +++ b/include/camp/resource/sycl.hpp @@ -335,34 +335,34 @@ namespace resources // Memory template - 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(size, qu); - switch (ma) { - case MemoryAccess::Unknown: - case MemoryAccess::Device: - ret = sycl::malloc_device(size, qu); - break; - case MemoryAccess::Pinned: - ret = sycl::malloc_host(size, qu); - break; - case MemoryAccess::Managed: - ret = sycl::malloc_shared(size, qu); - break; - } + switch (ma) { + case MemoryAccess::Unknown: + case MemoryAccess::Device: + ret = sycl::malloc_device(n, qu); + break; + case MemoryAccess::Pinned: + ret = sycl::malloc_host(n, qu); + break; + case MemoryAccess::Managed: + ret = sycl::malloc_shared(n, qu); + break; } return ret; } void* calloc(size_t size, MemoryAccess ma = MemoryAccess::Device) { - T* ret = nullptr; - if (size > 0) { - ret = allocate(size, ma); - this->memset(ret, 0, size); + if (size == 0) { + return nullptr; } + void* ret = allocate(size, ma); + this->memset(ret, 0, size); return ret; } @@ -377,16 +377,18 @@ namespace resources 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 From 6d0ab822e987c76fdcbf480ae651b2614cde05fc Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Tue, 14 Jul 2026 18:03:10 -0700 Subject: [PATCH 4/9] Unify MemoryAccess::Unknown behavior Throw on attempt to allocate unknown with gpu backends, unless size is 0, previously unknown mapped to device. Unknown is allowed on the deallcoate side as it will cause a runtime lookup of the pointer's attributes. Note that cpu backends ignore MemoryAccess entirely and Omp only supports device. --- include/camp/resource/cuda.hpp | 8 +++++--- include/camp/resource/hip.hpp | 4 +++- include/camp/resource/sycl.hpp | 4 +++- 3 files changed, 11 insertions(+), 5 deletions(-) diff --git a/include/camp/resource/cuda.hpp b/include/camp/resource/cuda.hpp index f62eff2b..da9cafe6 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"); @@ -280,7 +280,6 @@ namespace resources T* ret = nullptr; auto d{device_guard(device)}; switch (ma) { - case MemoryAccess::Unknown: case MemoryAccess::Device: CAMP_CUDA_API_INVOKE_AND_CHECK(cudaMalloc, (void**)&ret, @@ -298,6 +297,9 @@ namespace resources (void**)&ret, sizeof(T) * n); break; + case MemoryAccess::Unknown: + ::camp::throw_re("Unknown memory access type, cannot allocate"); + break; } return ret; } diff --git a/include/camp/resource/hip.hpp b/include/camp/resource/hip.hpp index 570f2c83..2d7236ff 100644 --- a/include/camp/resource/hip.hpp +++ b/include/camp/resource/hip.hpp @@ -282,7 +282,6 @@ namespace resources T* ret = nullptr; auto d{device_guard(device)}; switch (ma) { - case MemoryAccess::Unknown: case MemoryAccess::Device: CAMP_HIP_API_INVOKE_AND_CHECK(hipMalloc, (void**)&ret, @@ -300,6 +299,9 @@ namespace resources (void**)&ret, sizeof(T) * n); break; + case MemoryAccess::Unknown: + ::camp::throw_re("Unknown memory access type, cannot allocate"); + break; } return ret; } diff --git a/include/camp/resource/sycl.hpp b/include/camp/resource/sycl.hpp index ab53457b..9df170cc 100644 --- a/include/camp/resource/sycl.hpp +++ b/include/camp/resource/sycl.hpp @@ -342,7 +342,6 @@ namespace resources } T* ret = nullptr; switch (ma) { - case MemoryAccess::Unknown: case MemoryAccess::Device: ret = sycl::malloc_device(n, qu); break; @@ -352,6 +351,9 @@ namespace resources case MemoryAccess::Managed: ret = sycl::malloc_shared(n, qu); break; + case MemoryAccess::Unknown: + ::camp::throw_re("Unknown memory access type, cannot allocate"); + break; } return ret; } From 0d73df91a65129ab7b5aa4e26e9072ac676750d3 Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Wed, 15 Jul 2026 08:07:45 -0700 Subject: [PATCH 5/9] Fix Hip Managed lookup deallocation failure Look for the right enum value when checking for managed pointers in Hip that caused deallocations of managed pointers to fail if the attributes had to be looked up. --- include/camp/resource/hip.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/camp/resource/hip.hpp b/include/camp/resource/hip.hpp index 2d7236ff..0b5ab543 100644 --- a/include/camp/resource/hip.hpp +++ b/include/camp/resource/hip.hpp @@ -193,7 +193,7 @@ namespace resources return MemoryAccess::Pinned; case hipMemoryTypeDevice: return MemoryAccess::Device; - case hipMemoryTypeUnified: + case hipMemoryTypeManaged: return MemoryAccess::Managed; default: return MemoryAccess::Unknown; From a4f5c35bf2d698f728ebff6be95dd6171b02b4ac Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Wed, 15 Jul 2026 08:36:10 -0700 Subject: [PATCH 6/9] Remove ifdef for hip >= 6 --- include/camp/resource/hip.hpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/include/camp/resource/hip.hpp b/include/camp/resource/hip.hpp index 0b5ab543..ffc7fa3a 100644 --- a/include/camp/resource/hip.hpp +++ b/include/camp/resource/hip.hpp @@ -184,11 +184,7 @@ 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: From c83fe4b054315fa12f20060cd0d3b58aef8c0cfd Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Wed, 15 Jul 2026 08:45:20 -0700 Subject: [PATCH 7/9] Explicitly handle null pointers in Omp registration --- include/camp/resource/omp_target.hpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp index addfb198..d5d9485b 100644 --- a/include/camp/resource/omp_target.hpp +++ b/include/camp/resource/omp_target.hpp @@ -278,6 +278,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; @@ -286,6 +289,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); @@ -295,6 +301,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); From af0ee6077dd5f2a7c8d216651cb4db03d95cf3ca Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Wed, 15 Jul 2026 09:00:33 -0700 Subject: [PATCH 8/9] Alloc calloc to return nullptr previously it would unconditionally call memset which would presumably segfault --- include/camp/resource/cuda.hpp | 4 +++- include/camp/resource/hip.hpp | 4 +++- include/camp/resource/host.hpp | 8 +++++--- include/camp/resource/omp_target.hpp | 4 +++- include/camp/resource/sycl.hpp | 4 +++- 5 files changed, 17 insertions(+), 7 deletions(-) diff --git a/include/camp/resource/cuda.hpp b/include/camp/resource/cuda.hpp index da9cafe6..9e7e20ce 100644 --- a/include/camp/resource/cuda.hpp +++ b/include/camp/resource/cuda.hpp @@ -310,7 +310,9 @@ namespace resources return nullptr; } void* ret = allocate(size, ma); - this->memset(ret, 0, size); + if (ret != nullptr) { + this->memset(ret, 0, size); + } return ret; } diff --git a/include/camp/resource/hip.hpp b/include/camp/resource/hip.hpp index ffc7fa3a..5aa22261 100644 --- a/include/camp/resource/hip.hpp +++ b/include/camp/resource/hip.hpp @@ -308,7 +308,9 @@ namespace resources return nullptr; } void* ret = allocate(size, ma); - this->memset(ret, 0, size); + if (ret != nullptr) { + this->memset(ret, 0, size); + } return ret; } diff --git a/include/camp/resource/host.hpp b/include/camp/resource/host.hpp index 26230486..7ef3393d 100644 --- a/include/camp/resource/host.hpp +++ b/include/camp/resource/host.hpp @@ -125,9 +125,11 @@ namespace resources if (size == 0) { return nullptr; } - void* p = allocate(size); - this->memset(p, 0, size); - return p; + void* ret = allocate(size); + if (ret != nullptr) { + this->memset(ret, 0, size); + } + return ret; } void deallocate(void* p, MemoryAccess = MemoryAccess::Device) diff --git a/include/camp/resource/omp_target.hpp b/include/camp/resource/omp_target.hpp index d5d9485b..c5f8a199 100644 --- a/include/camp/resource/omp_target.hpp +++ b/include/camp/resource/omp_target.hpp @@ -235,7 +235,9 @@ namespace resources } check_ma(ma); void* ret = allocate(size); - this->memset(ret, 0, size); + if (ret != nullptr) { + this->memset(ret, 0, size); + } return ret; } diff --git a/include/camp/resource/sycl.hpp b/include/camp/resource/sycl.hpp index 9df170cc..6f32a0fc 100644 --- a/include/camp/resource/sycl.hpp +++ b/include/camp/resource/sycl.hpp @@ -364,7 +364,9 @@ namespace resources return nullptr; } void* ret = allocate(size, ma); - this->memset(ret, 0, size); + if (ret != nullptr) { + this->memset(ret, 0, size); + } return ret; } From c10081b882f03c5a9fe8c9ed9d9a60149de52434 Mon Sep 17 00:00:00 2001 From: Jason Burmark Date: Wed, 22 Jul 2026 11:48:57 -0700 Subject: [PATCH 9/9] Use std::calloc in Host::calloc This can be significantly faster than malloc+memset on some systems. --- include/camp/resource/host.hpp | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/include/camp/resource/host.hpp b/include/camp/resource/host.hpp index 7ef3393d..cfa1d10d 100644 --- a/include/camp/resource/host.hpp +++ b/include/camp/resource/host.hpp @@ -125,11 +125,7 @@ namespace resources if (size == 0) { return nullptr; } - void* ret = allocate(size); - if (ret != nullptr) { - this->memset(ret, 0, size); - } - return ret; + return std::calloc(size, 1); } void deallocate(void* p, MemoryAccess = MemoryAccess::Device)