Regards #291: Added support for untyped, owned, memory regions:

* Added a `unique_region` class - non-templated
* `unique_region`'s "decay" into `cuda::memory::region_t`'s - so there's no need to reimplement half the world for them
* Added a missing variant of `cuda::memory::copy()`
* Added an example program based on vectorAdd, which uses unique_regions (I would have replaced vectorAdd with it - had it not requirted to lines of source more than before...)
* `region_t`'s now have an `as_span` method, as well as an `operator span()`
* Removed some commented-out unique-pointer-related code

examples/CMakeLists.txt

@@ -56,6 +56,7 @@ link_libraries(cuda-api-wrappers::runtime-and-driver)
 
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "bin")
 add_executable(vectorAdd modified_cuda_samples/vectorAdd/vectorAdd.cu)
+add_executable(vectorAdd_unique_regions modified_cuda_samples/vectorAdd_unique_regions/vectorAdd_unique_regions.cu)
 add_executable(vectorAddMapped modified_cuda_samples/vectorAddMapped/vectorAddMapped.cu)
 add_executable(vectorAddManaged modified_cuda_samples/vectorAddManaged/vectorAddManaged.cu)
 add_executable(vectorAdd_nvrtc modified_cuda_samples/vectorAdd_nvrtc/vectorAdd_nvrtc.cpp)

examples/by_api_module/unified_addressing.cpp

@@ -28,13 +28,13 @@ void pointer_properties(const cuda::device_t& device)
 		cuda::context::create(device),
 		cuda::context::create(device)
 	};
-	cuda::memory::device::unique_ptr<char[]> regions[2] = {
-		cuda::memory::make_unique<char[]>(contexts[0], fixed_size),
-		cuda::memory::make_unique<char[]>(contexts[1], fixed_size)
+	cuda::memory::device::unique_region regions[2] = {
+		cuda::memory::make_unique_region(contexts[0], fixed_size),
+		cuda::memory::make_unique_region(contexts[1], fixed_size)
 	};
 	void* raw_pointers[2] = {
-		regions[0].get(),
-		regions[1].get()
+		regions[0].data(),
+		regions[1].data()
 	};
 	cuda::memory::pointer_t<void> pointers[2] = {
 		cuda::memory::pointer::wrap(raw_pointers[0]),

examples/modified_cuda_samples/vectorAdd_unique_regions/vectorAdd_unique_regions.cu

@@ -0,0 +1,85 @@
+/**
+ * Derived from the nVIDIA CUDA 8.0 samples by
+ *
+ *   Eyal Rozenberg
+ *
+ * The derivation is specifically permitted in the nVIDIA CUDA Samples EULA
+ * and the deriver is the owner of this code according to the EULA.
+ *
+ * Use this reasonably. If you want to discuss licensing formalities, please
+ * contact the author.
+ */
+
+#include "../../common.hpp"
+
+#include <cuda/api.hpp>
+
+#include <iostream>
+#include <memory>
+#include <algorithm>
+
+__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
+{
+	int i = blockDim.x * blockIdx.x + threadIdx.x;
+	if (i < numElements) { C[i] = A[i] + B[i]; }
+}
+
+int main()
+{
+	if (cuda::device::count() == 0) {
+		std::cerr << "No CUDA devices on this system" << "\n";
+		exit(EXIT_FAILURE);
+	}
+
+	int numElements = 50000;
+	std::cout << "[Vector addition of " << numElements << " elements]\n";
+
+	// If we could rely on C++14, we would  use std::make_unique
+	auto h_A = std::unique_ptr<float[]>(new float[numElements]);
+	auto h_B = std::unique_ptr<float[]>(new float[numElements]);
+	auto h_C = std::unique_ptr<float[]>(new float[numElements]);
+
+	auto generator = []() { return rand() / (float) RAND_MAX; };
+	std::generate(h_A.get(), h_A.get() + numElements, generator);
+	std::generate(h_B.get(), h_B.get() + numElements, generator);
+
+	auto device = cuda::device::current::get();
+
+	auto d_A = cuda::memory::make_unique_region(device, numElements * sizeof(float));
+	auto d_B = cuda::memory::make_unique_region(device, numElements * sizeof(float));
+	auto d_C = cuda::memory::make_unique_region(device, numElements * sizeof(float));
+	auto sp_A = d_A.as_span<float>();
+	auto sp_B = d_B.as_span<float>();
+	auto sp_C = d_C.as_span<float>();
+
+	cuda::memory::copy(sp_A, h_A.get());
+	cuda::memory::copy(sp_B, h_B.get());
+
+	auto launch_config = cuda::launch_config_builder()
+		.overall_size(numElements)
+		.block_size(256)
+		.build();
+
+	std::cout
+		<< "CUDA kernel launch with " << launch_config.dimensions.grid.x
+		<< " blocks of " << launch_config.dimensions.block.x << " threads each\n";
+
+	cuda::launch(
+		vectorAdd, launch_config,
+		sp_A.data(), sp_B.data(), sp_C.data(), numElements
+	);
+
+	cuda::memory::copy(h_C.get(), sp_C);
+
+	// Verify that the result vector is correct
+	for (int i = 0; i < numElements; ++i) {
+		if (fabs(h_A.get()[i] + h_B.get()[i] - h_C.get()[i]) > 1e-5)  {
+			std::cerr << "Result verification failed at element " << i << "\n";
+			exit(EXIT_FAILURE);
+		}
+	}
+
+	std::cout << "Test PASSED\n";
+	std::cout << "SUCCESS\n";
+}
+

src/cuda/api.hpp

@@ -32,6 +32,7 @@
 #include "api/memory_pool.hpp"
 #endif
 #include "api/unique_ptr.hpp"
+#include "api/unique_region.hpp"
 #include "api/link_options.hpp"
 
 #include "api/device.hpp"
@@ -68,6 +69,7 @@
 #include "api/multi_wrapper_impls/context.hpp"
 #include "api/multi_wrapper_impls/stream.hpp"
 #include "api/multi_wrapper_impls/memory.hpp"
+#include "api/multi_wrapper_impls/unique_region.hpp"
 #include "api/multi_wrapper_impls/virtual_memory.hpp"
 #include "api/multi_wrapper_impls/kernel.hpp"
 #include "api/multi_wrapper_impls/kernel_launch.hpp"

src/cuda/api/memory.hpp

@@ -504,6 +504,16 @@ inline void copy(region_t destination, void* source, size_t num_bytes)
 	return copy(destination.start(), source, num_bytes);
 }
 
+inline void copy(void* destination, const_region_t source, size_t num_bytes)
+{
+#ifndef NDEBUG
+	if (source.size() < num_bytes) {
+		throw ::std::logic_error("Number of bytes to copy exceeds source size");
+	}
+#endif
+	return copy(destination, source.start(), num_bytes);
+}
+
 inline void copy(region_t destination, void* source)
 {
 	return copy(destination, source, destination.size());

src/cuda/api/multi_wrapper_impls/unique_region.hpp

@@ -0,0 +1,122 @@
+/**
+ * @file
+ *
+ * @brief Implementations of @ref `make_unique_region()` functions
+ */
+#pragma once
+#ifndef MULTI_WRAPPER_IMPLS_UNIQUE_REGION_HPP_
+#define MULTI_WRAPPER_IMPLS_UNIQUE_REGION_HPP_
+
+#include "../unique_region.hpp"
+#include "../types.hpp"
+
+namespace cuda {
+
+namespace memory {
+
+namespace device {
+
+inline unique_region make_unique_region(const context_t& context, cuda::size_t num_elements)
+{
+	return detail_::make_unique_region(context.handle(), num_elements);
+}
+
+/**
+ * @brief Create a variant of ::std::unique_pointer for an array in
+ * device-global memory
+ *
+ * @tparam T  an array type; _not_ the type of individual elements
+ *
+ * @param device        on which to construct the array of elements
+ * @param num_elements  the number of elements to allocate
+ * @return an ::std::unique_ptr pointing to the constructed T array
+ */
+inline unique_region make_unique_region(const device_t& device, size_t num_elements)
+{
+	auto pc = device.primary_context();
+	return make_unique_region(pc, num_elements);
+}
+
+/**
+ * @brief Create a variant of ::std::unique_pointer for an array in
+ * device-global memory on the current device.
+ *
+ * @note The allocation will be made in the device's primary context -
+ * which will be created if it has not yet been.
+ *
+ * @tparam T  an array type; _not_ the type of individual elements
+ *
+ * @param num_elements  the number of elements to allocate
+ *
+ * @return an ::std::unique_ptr pointing to the constructed T array
+ */
+inline unique_region make_unique_region(size_t num_elements)
+{
+	auto current_device_id = cuda::device::current::detail_::get_id();
+	auto pc = cuda::device::primary_context::detail_::leaky_get(current_device_id);
+	return make_unique_region(pc, num_elements);
+}
+
+} // namespace device
+
+namespace host {
+
+inline unique_region make_unique_region(
+	const context_t&    context,
+	size_t              num_bytes,
+	allocation_options  options)
+{
+	CAW_SET_SCOPE_CONTEXT(context.handle());
+	return unique_region{ allocate(num_bytes, options) };
+}
+
+inline unique_region make_unique_region(const device_t& device, size_t num_bytes)
+{
+	auto pc = device.primary_context();
+	return make_unique_region(pc, num_bytes);
+}
+
+inline unique_region make_unique_region(size_t num_bytes)
+{
+	return unique_region { allocate(num_bytes) };
+}
+
+} // namespace host
+
+namespace managed {
+
+inline unique_region make_unique_region(
+	const context_t&      context,
+	size_t                num_bytes,
+	initial_visibility_t  initial_visibility)
+{
+	CAW_SET_SCOPE_CONTEXT(context.handle());
+	return unique_region { detail_::allocate_in_current_context(num_bytes, initial_visibility) };
+}
+
+inline unique_region make_unique_region(
+	const device_t&       device,
+	size_t                num_bytes,
+	initial_visibility_t  initial_visibility)
+{
+	auto pc = device.primary_context();
+	return make_unique_region(pc, num_bytes, initial_visibility);
+}
+
+inline unique_region make_unique_region(
+	size_t                num_bytes,
+	initial_visibility_t  initial_visibility)
+{
+	auto current_device_id = cuda::device::current::detail_::get_id();
+	auto pc = cuda::device::primary_context::detail_::leaky_get(current_device_id);
+	return make_unique_region(pc, num_bytes, initial_visibility);
+}
+
+} // namespace managed
+
+} // namespace memory
+
+} // namespace cuda
+
+#endif // MULTI_WRAPPER_IMPLS_UNIQUE_REGION_HPP_
+