Fixed #291: Added a unique_span class - a uniquely-owned span. It s…

…hould provide the most flexibility in our example programs (and typical user programs) for direct use without explicit casting, re-specifying types and sizes, etc.

examples/by_api_module/event_management.cu

@@ -94,13 +94,12 @@ int main(int argc, char **argv)
 		cuda::event::do_record_timings,
 		cuda::event::not_interprocess);
 
-	constexpr size_t buffer_size = 12345678;
-	auto buffer = cuda::memory::managed::make_unique<char[]>(
-		device, buffer_size, cuda::memory::managed::initial_visibility_t::to_all_devices);
+	auto buffer = cuda::memory::managed::make_unique_span<char>(
+		device, 12345678, cuda::memory::managed::initial_visibility_t::to_all_devices);
 	auto wrapped_kernel = cuda::kernel::get(device, increment);
 	auto launch_config = cuda::launch_config_builder()
 		.kernel(&wrapped_kernel)
-		.overall_size(buffer_size)
+		.overall_size(buffer.size())
 		.use_maximum_linear_block()
 		.build();
 
@@ -110,7 +109,7 @@ int main(int argc, char **argv)
 		report_occurrence("In first callback (enqueued after first event but before first kernel)", event_1, event_2);
 	};
 	stream.enqueue.host_invokable(first_callback);
-	stream.enqueue.kernel_launch(increment, launch_config, buffer.get(), buffer_size);
+	stream.enqueue.kernel_launch(increment, launch_config, buffer.data(), buffer.size());
 	auto second_callback = [&] {
 		report_occurrence("In second callback (enqueued after the first kernel but before the second event)",
 			event_1, event_2);
@@ -136,7 +135,7 @@ int main(int argc, char **argv)
 	report_occurrence("After synchronizing on event_2, but before synchronizing on the stream", event_1, event_2);
 	std::cout
 		<< cuda::event::time_elapsed_between(event_1, event_2).count() << " msec have elapsed, "
-		<< "executing the second kernel (\"increment\") on a buffer of " << buffer_size
+		<< "executing the second kernel (\"increment\") on a buffer of " << buffer.size()
 		<< " chars and triggering two callbacks.\n";
 	// ... and this should make the third kernel execute
 	stream.synchronize();

examples/by_api_module/stream_management.cu

@@ -154,21 +154,21 @@ int main(int argc, char **argv)
 #endif
 
 	constexpr auto buffer_size = 12345678;
-	auto buffer = cuda::memory::managed::make_unique<char[]>(
+	auto buffer = cuda::memory::managed::make_unique_span<char>(
 		buffer_size,
 		device.supports_concurrent_managed_access() ?
 			cuda::memory::managed::initial_visibility_t::to_supporters_of_concurrent_managed_access:
 			cuda::memory::managed::initial_visibility_t::to_all_devices);
-	print_first_char(buffer.get());
-	std::fill(buffer.get(), buffer.get() + buffer_size, 'a');
-	print_first_char(buffer.get());
+	print_first_char(buffer.data());
+	std::fill(buffer.begin(), buffer.end(), 'a');
+	print_first_char(buffer.data());
 
 	auto event_1 = cuda::event::create(device, cuda::event::sync_by_blocking);
 	stream_1.enqueue.kernel_launch(print_message<N,3>, single_thread_config, message<N>("I'm on stream 1"));
-	stream_1.enqueue.memset(buffer.get(), 'b', buffer_size);
+	stream_1.enqueue.memset(buffer, 'b');
 	auto callback = [&]() {
 		std::cout << "Callback from stream 1!... \n";
-		print_first_char(buffer.get());
+		print_first_char(buffer.data());
 	};
 	stream_1.enqueue.host_invokable(callback);
 	auto threads_per_block = cuda::kernel::get(device, increment).get_attribute(CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK);
@@ -177,15 +177,15 @@ int main(int argc, char **argv)
 	// TODO: The following doesn't have much of a meaningful effect; we should modify this example
 	// so that the attachment has some observable effect
 	stream_1.enqueue.attach_managed_region(buffer.get());
-	stream_1.enqueue.kernel_launch(increment, launch_config, buffer.get(), buffer_size);
+	stream_1.enqueue.kernel_launch(increment, launch_config, buffer.data(), buffer_size);
 	event_1.record(stream_1);
 	stream_1.enqueue.kernel_launch(print_message<N,4>, single_thread_config, message<N>("I'm on stream 1"));
 	stream_2.enqueue.wait(event_1);
-	stream_2.enqueue.kernel_launch(print_first_char_kernel, launch_config , buffer.get());
+	stream_2.enqueue.kernel_launch(print_first_char_kernel, launch_config , buffer.data());
 	stream_2.enqueue.kernel_launch(print_message<N,5>, single_thread_config, message<N>("I'm on stream 2"));
 	bool idleness_1 = stream_2.has_work_remaining();
 	device.synchronize();
-	print_first_char(buffer.get());
+	print_first_char(buffer.data());
 	// cuda::memory::managed::free(buffer);
 	bool idleness_2 = stream_2.has_work_remaining();
 	std::cout << std::boolalpha

examples/modified_cuda_samples/asyncAPI/asyncAPI.cu

@@ -29,12 +29,12 @@ __global__ void increment_kernel(datum*g_data, datum inc_value)
 	g_data[global_idx] = g_data[global_idx] + inc_value;
 }
 
-bool correct_output(int *data, const int n, const int x)
+bool correct_output(cuda::span<const int> data, const int x)
 {
-	for (int i = 0; i < n; i++)
+	for (size_t i = 0; i < data.size(); i++)
 		if (data[i] != x)
 		{
-			printf("Error! data[%d] = %d, ref = %d\n", i, data[i], x);
+			printf("Error! data[%lu] = %d, ref = %d\n", i, data[i], x);
 			return false;
 		}
 	return true;
@@ -51,15 +51,14 @@ int main(int, char **)
 
 	std::cout << "CUDA device [" <<  device.name() << "]\n";
 
-	int n = 16 * 1024 * 1024;
-	int num_bytes = n * sizeof(datum);
+	const int n = 16 * 1024 * 1024;
 	int value = 26;
 
 	// allocate host memory
-	auto a = cuda::memory::host::make_unique<datum[]>(n);
-	cuda::memory::host::zero(a.get(), num_bytes);
+	auto a = cuda::memory::host::make_unique_span<datum>(n);
+	cuda::memory::host::zero(a);
 
-	auto d_a = cuda::memory::make_unique<datum[]>(device, n);
+	auto d_a = cuda::memory::make_unique_span<datum>(device, n);
 
 	auto launch_config = cuda::launch_config_builder()
 		.overall_size(n)
@@ -80,9 +79,9 @@ int main(int, char **)
 	auto stream = device.default_stream(); // device.create_stream(cuda::stream::async);
 	auto cpu_time_start = std::chrono::high_resolution_clock::now();
 	stream.enqueue.event(start_event);
-	stream.enqueue.copy(d_a.get(), a.get(), num_bytes);
-	stream.enqueue.kernel_launch(increment_kernel, launch_config, d_a.get(), value);
-	stream.enqueue.copy(a.get(), d_a.get(), num_bytes);
+	stream.enqueue.copy(d_a, a);
+	stream.enqueue.kernel_launch(increment_kernel, launch_config, d_a.data(), value);
+	stream.enqueue.copy(a, d_a);
 	stream.enqueue.event(end_event);
 	auto cpu_time_end = std::chrono::high_resolution_clock::now();
 
@@ -99,7 +98,7 @@ int main(int, char **)
 	std::cout << "time spent by CPU in CUDA calls: " << std::setprecision(2)<< (cpu_time_end - cpu_time_start).count() << '\n';
 	std::cout << "CPU executed " << counter << " iterations while waiting for GPU to finish\n";
 
-	auto bFinalResults = correct_output(a.get(), n, value);
+	auto bFinalResults = correct_output(a, value);
 
 	std::cout << (bFinalResults ? "SUCCESS" : "FAILURE") << '\n';
 

examples/modified_cuda_samples/bandwidthtest/bandwidthtest.cpp

@@ -89,17 +89,17 @@ int main()
 		std::unique_ptr<float[]>(new float[nElements])
 	);
 
-	auto device_buffer = cuda::memory::device::make_unique<float[]>(nElements);
+	auto device_buffer = cuda::memory::device::make_unique_span<float>(nElements);
 
 	auto pinned_host_buffers = std::make_pair(
-		cuda::memory::host::make_unique<float[]>(nElements),
-		cuda::memory::host::make_unique<float[]>(nElements)
+		cuda::memory::host::make_unique_span<float>(nElements),
+		cuda::memory::host::make_unique_span<float>(nElements)
 	);
 
 	auto h_aPageable = pageable_host_buffers.first.get();
 	auto h_bPageable = pageable_host_buffers.second.get();
-	auto h_aPinned = pinned_host_buffers.first.get();
-	auto h_bPinned = pinned_host_buffers.second.get();
+	auto h_aPinned = pinned_host_buffers.first.data();
+	auto h_bPinned = pinned_host_buffers.second.data();
 
 	std::iota(h_aPageable, h_aPageable + nElements, 0.0);
 	cuda::memory::copy(h_aPinned, h_aPageable, bytes);
@@ -112,6 +112,6 @@ int main()
 	std::cout << "\nTransfer size (MB): " << (bytes / Mi) << "\n";
 
 	// perform copies and report bandwidth
-	profileCopies(h_aPageable, h_bPageable, device_buffer.get(), nElements, "Pageable");
-	profileCopies(h_aPinned, h_bPinned, device_buffer.get(), nElements, "Pinned");
+	profileCopies(h_aPageable, h_bPageable, device_buffer.data(), nElements, "Pageable");
+	profileCopies(h_aPinned, h_bPinned, device_buffer.data(), nElements, "Pinned");
 }

examples/modified_cuda_samples/clock_nvrtc/clock.cpp

@@ -154,20 +154,20 @@ int main()
 	{
 		const auto dynamic_shared_mem_size = sizeof(float) * 2 * num_threads_per_block;
 
-		auto d_input = cuda::memory::make_unique<float[]>(device, input_size);
-		auto d_output = cuda::memory::make_unique<float[]>(device, num_blocks);
+		auto d_input = cuda::memory::make_unique_span<float>(device, input_size);
+		auto d_output = cuda::memory::make_unique_span<float>(device, num_blocks);
 			// Note: We won't actually be checking the output...
-		auto d_timers = cuda::memory::make_unique<clock_t []>(device, num_timers);
-		cuda::memory::copy(d_input.get(), input.get(), input_size * sizeof(float));
+		auto d_timers = cuda::memory::make_unique_span<clock_t>(device, num_timers);
+		cuda::memory::copy(d_input, input.get());
 
 		auto launch_config = cuda::launch_config_builder()
 			.num_blocks(num_blocks)
 			.block_size(num_threads_per_block)
 			.dynamic_shared_memory_size(dynamic_shared_mem_size)
 			.build();
-		cuda::launch(kernel_in_module, launch_config, d_input.get(), d_output.get(), d_timers.get());
+		cuda::launch(kernel_in_module, launch_config, d_input.data(), d_output.data(), d_timers.data());
 		device.synchronize();
-		cuda::memory::copy(timers.get(), d_timers.get(), num_timers * sizeof(clock_t));
+		cuda::memory::copy(timers.get(), d_timers);
 	} // The allocated device buffers are released here
 	long double average_elapsed_clock_ticks_per_block = compute_average_elapsed_clocks(timers.get(), num_blocks);