Fixes #256: Replaced our library's detail namespaces with detail_…

…, to avoid libcu++ having ambiguity of some of its own `detail` namespaces.

examples/by_runtime_api_module/device_management.cpp

@@ -191,14 +191,14 @@ int main(int argc, char **argv)
 		auto device_1 = cuda::device::get(1);
 		cuda::device::current::set(device_0);
 		assert(cuda::device::current::get() == device_0);
-		assert(cuda::device::current::detail::get_id() == device_0.id());
+		assert(cuda::device::current::detail_::get_id() == device_0.id());
 		cuda::device::current::set(device_1);
 		assert(cuda::device::current::get() == device_1);
-		assert(cuda::device::current::detail::get_id() == device_1.id());
+		assert(cuda::device::current::detail_::get_id() == device_1.id());
 	}
 
 	try {
-		cuda::device::current::detail::set(device_count);
+		cuda::device::current::detail_::set(device_count);
 		die_("Should not have been able to set the current device to "
 			+ std::to_string(device_count) + " since that's the device count, and "
 			+ "the maximum valid ID should be " + std::to_string(device_count - 1)

examples/by_runtime_api_module/error_handling.cu

@@ -28,7 +28,7 @@ int main(int argc, char **argv)
 	}
 
 	try {
-		cuda::device::current::detail::set(device_count);
+		cuda::device::current::detail_::set(device_count);
 		die_("An exception should have be thrown");
 	}
 	catch(cuda::runtime_error& e) {
@@ -51,7 +51,7 @@ int main(int argc, char **argv)
 	// clearing the error
 
 	try {
-		cuda::device::current::detail::set(device_count);
+		cuda::device::current::detail_::set(device_count);
 		die_("An exception should have be thrown");
 	}
 	catch(cuda::runtime_error&) { }

src/cuda/api/array.hpp

@@ -20,7 +20,7 @@ class device_t;
 
 namespace array {
 
-namespace detail {
+namespace detail_ {
 
 template<typename T>
 cudaArray* allocate_on_current_device(array::dimensions_t<3> dimensions)
@@ -49,7 +49,7 @@ cudaArray* allocate(device_t& device, array::dimensions_t<3> dimensions);
 template<typename T>
 cudaArray* allocate(device_t& device, array::dimensions_t<2> dimensions);
 
-} // namespace detail
+} // namespace detail_
 
 } // namespace array
 
@@ -91,7 +91,7 @@ class array_t {
 	 * Creates and wraps a new CUDA array.
 	 */
 	array_t(device_t& device, array::dimensions_t<NumDimensions> dimensions)
-		: array_t(array::detail::allocate<T>(device, dimensions), dimensions) {}
+		: array_t(array::detail_::allocate<T>(device, dimensions), dimensions) {}
 	array_t(const array_t& other) = delete;
 	array_t(array_t&& other) noexcept : array_t(other.raw_array_, other.dimensions_)
 	{

src/cuda/api/current_device.hpp

@@ -37,7 +37,7 @@ namespace device {
 
 namespace current {
 
-namespace detail {
+namespace detail_ {
 
 /**
  * Obtains the numeric id of the device set as current for the CUDA Runtime API
@@ -79,14 +79,14 @@ inline void set(const id_t* device_ids, size_t num_devices)
 }
 
 /**
- * @note See the out-of-`detail::` version of this class.
+ * @note See the out-of-`detail_::` version of this class.
  */
 class scoped_override_t {
 protected:
 	static id_t  replace(id_t new_device_id)
 	{
-		id_t previous_device_id = device::current::detail::get_id();
-		device::current::detail::set(new_device_id);
+		id_t previous_device_id = device::current::detail_::get_id();
+		device::current::detail_::set(new_device_id);
 		return previous_device_id;
 	}
 
@@ -102,12 +102,12 @@ class scoped_override_t {
 };
 
 
-} // namespace detail
+} // namespace detail_
 
 /**
  * Reset the CUDA Runtime API's current device to its default value - the default device
  */
-inline void set_to_default() { return detail::set(device::default_device_id); }
+inline void set_to_default() { return detail_::set(device::default_device_id); }
 
 void set(device_t device);
 
@@ -116,9 +116,9 @@ void set(device_t device);
  * what remains of the current scope, and changing it back to its previous value
  * when exiting the scope.
  */
-class scoped_override_t : private detail::scoped_override_t {
+class scoped_override_t : private detail_::scoped_override_t {
 protected:
-	using parent = detail::scoped_override_t;
+	using parent = detail_::scoped_override_t;
 public:
 	scoped_override_t(device_t& device);
 	scoped_override_t(device_t&& device);

src/cuda/api/detail/device_properties.hpp

@@ -1,5 +1,5 @@
 /**
- * @file detail/device_properties.hpp
+ * @file detail_/device_properties.hpp
  *
  * @brief Implementation of methods and helper functions for device-property-related classes.
  *
@@ -47,15 +47,15 @@ inline constexpr bool operator >=(const compute_architecture_t& lhs, const compu
 	return lhs.major > rhs.major;
 }
 
-namespace detail {
+namespace detail_ {
 
 constexpr const int invalid_architecture_return { 0 };
 enum : memory::shared::size_t { KiB = 1024 };
 
 template <typename T>
 inline T ensure_arch_property_validity(T v, const compute_architecture_t& arch)
 {
-	if (v == detail::invalid_architecture_return) {
+	if (v == detail_::invalid_architecture_return) {
 		throw ::std::invalid_argument("No architecture numbered " + ::std::to_string(arch.major));
 	}
 	return v;
@@ -147,30 +147,30 @@ inline constexpr unsigned max_in_flight_threads_per_processor(const compute_arch
 		invalid_architecture_return;
 }
 
-} // namespace detail
+} // namespace detail_
 
 inline const char* compute_architecture_t::name() const {
-	return detail::ensure_arch_property_validity(detail::architecture_name(*this), *this);
+	return detail_::ensure_arch_property_validity(detail_::architecture_name(*this), *this);
 }
 
 inline unsigned compute_architecture_t::max_in_flight_threads_per_processor() const
 {
-	return detail::ensure_arch_property_validity(detail::max_in_flight_threads_per_processor(*this), *this);
+	return detail_::ensure_arch_property_validity(detail_::max_in_flight_threads_per_processor(*this), *this);
 }
 
 inline unsigned compute_architecture_t::max_shared_memory_per_block() const
 {
-	return detail::ensure_arch_property_validity(detail::max_shared_memory_per_block(*this), *this);
+	return detail_::ensure_arch_property_validity(detail_::max_shared_memory_per_block(*this), *this);
 }
 
 inline unsigned compute_architecture_t::max_resident_warps_per_processor() const
 {
-	return detail::ensure_arch_property_validity(detail::max_resident_warps_per_processor(*this), *this);
+	return detail_::ensure_arch_property_validity(detail_::max_resident_warps_per_processor(*this), *this);
 }
 
 inline unsigned compute_architecture_t::max_warp_schedulings_per_processor_cycle() const
 {
-	return detail::ensure_arch_property_validity(detail::max_warp_schedulings_per_processor_cycle(*this), *this);
+	return detail_::ensure_arch_property_validity(detail_::max_warp_schedulings_per_processor_cycle(*this), *this);
 }
 
 // compute_capability_t-related
@@ -224,7 +224,7 @@ inline constexpr compute_capability_t make_compute_capability(unsigned major, un
 	return { {major}, minor };
 }
 
-namespace detail {
+namespace detail_ {
 
 inline constexpr unsigned max_in_flight_threads_per_processor(const compute_capability_t& cc)
 {
@@ -262,26 +262,26 @@ inline constexpr unsigned max_resident_warps_per_processor(const compute_capabil
 		max_resident_warps_per_processor(cc.architecture);
 }
 
-} // namespace detail
+} // namespace detail_
 
 inline unsigned compute_capability_t::max_in_flight_threads_per_processor() const
 {
-	return detail::ensure_arch_property_validity(detail::max_in_flight_threads_per_processor(*this), architecture);
+	return detail_::ensure_arch_property_validity(detail_::max_in_flight_threads_per_processor(*this), architecture);
 }
 
 inline unsigned compute_capability_t::max_warp_schedulings_per_processor_cycle() const
 {
-	return detail::ensure_arch_property_validity(detail::max_warp_schedulings_per_processor_cycle(*this), architecture);
+	return detail_::ensure_arch_property_validity(detail_::max_warp_schedulings_per_processor_cycle(*this), architecture);
 }
 
 inline unsigned compute_capability_t::max_shared_memory_per_block() const
 {
-	return detail::ensure_arch_property_validity(detail::max_shared_memory_per_block(*this), architecture);
+	return detail_::ensure_arch_property_validity(detail_::max_shared_memory_per_block(*this), architecture);
 }
 
 inline unsigned compute_capability_t::max_resident_warps_per_processor() const
 {
-	return detail::ensure_arch_property_validity(detail::max_resident_warps_per_processor(*this), architecture);
+	return detail_::ensure_arch_property_validity(detail_::max_resident_warps_per_processor(*this), architecture);
 }
 
 // properties_t-related

src/cuda/api/device.hpp

@@ -155,7 +155,7 @@ class device_t {
 	using resource_id_t = cudaLimit;
 
 protected: // types
-	using scoped_setter_t = device::current::detail::scoped_override_t;
+	using scoped_setter_t = device::current::detail_::scoped_override_t;
 	using flags_t = unsigned;
 
 	///@cond
@@ -180,8 +180,8 @@ class device_t {
 	protected:
 		const device::id_t device_id_;
 
-		using deleter = memory::device::detail::deleter;
-		using allocator = memory::device::detail::allocator;
+		using deleter = memory::device::detail_::deleter;
+		using allocator = memory::device::detail_::allocator;
 
 	public:
 		///@cond
@@ -199,7 +199,7 @@ class device_t {
 		memory::region_t allocate(size_t size_in_bytes)
 		{
 			scoped_setter_t set_device_for_this_scope(device_id_);
-			return memory::device::detail::allocate(size_in_bytes);
+			return memory::device::detail_::allocate(size_in_bytes);
 		}
 
 		// Perhaps drop this? it should really go into a managed namespace
@@ -231,7 +231,7 @@ class device_t {
 				initial_visibility_t::to_supporters_of_concurrent_managed_access)
 		{
 			scoped_setter_t set_device_for_this_scope(device_id_);
-			return cuda::memory::managed::detail::allocate(size_in_bytes, initial_visibility);
+			return cuda::memory::managed::detail_::allocate(size_in_bytes, initial_visibility);
 		}
 
 		/**
@@ -688,7 +688,7 @@ class device_t {
 	 */
 	device_t& make_current()
 	{
-		device::current::detail::set(id());
+		device::current::detail_::set(id());
 		return *this;
 	}
 
@@ -754,12 +754,12 @@ namespace current {
 /**
  * Obtains (a proxy for) the device which the CUDA runtime API considers to be current.
  */
-inline device_t get() { return device::get(detail::get_id()); }
+inline device_t get() { return device::get(detail_::get_id()); }
 
 /**
  * Tells the CUDA runtime API to consider the specified device as the current one.
  */
-inline void set(device_t device) { detail::set(device.id()); }
+inline void set(device_t device) { detail_::set(device.id()); }
 
 } // namespace current
 
@@ -771,7 +771,7 @@ inline void set(device_t device) { detail::set(device.id()); }
  */
 inline device_t get(pci_location_t pci_id)
 {
-	auto resolved_id = device::detail::resolve_id(pci_id);
+	auto resolved_id = device::detail_::resolve_id(pci_id);
 	return get(resolved_id);
 }
 
@@ -797,7 +797,7 @@ inline device_t get(const ::std::string& pci_id_str)
 inline void synchronize(device_t& device)
 {
 	auto device_id = device.id();
-	device::current::detail::scoped_override_t set_device_for_this_scope(device_id);
+	device::current::detail_::scoped_override_t set_device_for_this_scope(device_id);
 	auto status = cudaDeviceSynchronize();
 	throw_if_error(status, "Failed synchronizing " + ::std::to_string(device_id));
 }

src/cuda/api/devices.hpp

@@ -11,7 +11,7 @@
 
 namespace cuda {
 
-namespace detail {
+namespace detail_ {
 
 // Note that while nothing constrains you from instantiating
 // this class many times, all instances are the same (as CUDA
@@ -203,11 +203,11 @@ inline bool operator!= (
 	return not (lhs == rhs);
 }
 
-} // namespace detail
+} // namespace detail_
 
-inline detail::all_devices devices()
+inline detail_::all_devices devices()
 {
-	return detail::all_devices();
+	return detail_::all_devices();
 }
 
 } // namespace cuda

src/cuda/api/error.hpp

@@ -134,7 +134,7 @@ constexpr inline bool is_failure(status_t status)  { return status != (status_t)
  */
 inline ::std::string describe(status_t status) { return cudaGetErrorString(status); }
 
-namespace detail {
+namespace detail_ {
 
 template <typename I, bool UpperCase = false>
 ::std::string as_hex(I x)
@@ -165,7 +165,7 @@ inline ::std::string ptr_as_hex(const I* ptr)
 	return as_hex((size_t) ptr);
 }
 
-} // namespace detail
+} // namespace detail_
 
 /**
  * A (base?) class for exceptions raised by CUDA code; these errors are thrown by