Fixes #92, regards #91: Dropped the grid_info namespace.

Also, preferring `lane::id()` over `lane::index()` when the latter was previously used.

README.md

@@ -92,22 +92,22 @@ num_full_warps_in_grid_block()
 ```
 the library groups these (and many other related) functions into relevant namespaces. We thus have:
 ```
-linear_grid::grid_info::thread::global_id()
-grid_info::lane::id()
-linear_grid::grid_info::thread::is_last_in_block()
-linear_grid::grid_info::block::num_full_warps()
+linear_grid::thread::global_id()
+lane::id()
+linear_grid::thread::is_last_in_block()
+linear_grid::block::num_full_warps()
 ```
-which is easier to browse through if you use auto-complete. The order comes at the expense of brevity... but we can alleviate this with an appropriate `namespace`. The above can then become simply:
+which is easier to browse through if you use auto-complete. The order comes at the expense of brevity... but we can alleviate this somewhat with a shorter namespace. The above can then become simply:
 ```
-namespace gi = kat::linear_grid::grid_info;
+namespace gi = kat::linear_grid;
 gi::thread::global_id()
 gi::lane::id()
 gi::thread::is_last_in_block()
 gi::block::num_full_warps()
 ```
 in your code. Now _this_ is how I want to write my kernels!
 
-You will note, that most similar phrases you could come up with about positions and sizes within the grid - already have implementations. For example: "I can get the number of full warps, but now I want the number of warps, period"; well, just replace `num_full_warps()` with `num_warps()` and it's there: `linear_grid::grid_info::block::num_warps()` is available.
+You will note, that most similar phrases you could come up with about positions and sizes within the grid - already have implementations. For example: "I can get the number of full warps, but now I want the number of warps, period"; well, just replace `num_full_warps()` with `num_warps()` and it's there: `linear_grid::block::num_warps()` is available.
 
 And as a final bonus - if you write a non-linear kernel, with blocks and grids having y and z dimensions other than 1 - you will only need to change your `namespace =` or `using` statements, to be able to write the same code and use 3-D implementations of these functions instead.
 

src/kat/on_device/collaboration/block.cuh

@@ -32,18 +32,6 @@ namespace collaborative {
 
 namespace block {
 
-///@cond
-// If we want to refer to other primitives, we'll make those references explicit;
-// but we do want to be able to say `warp::id()` without prefixing that with anything.
-
-namespace grid   = grid_info::grid;
-namespace block  = grid_info::block;
-namespace warp   = grid_info::warp;
-namespace thread = grid_info::thread;
-namespace lane   = grid_info::lane;
-
-///@endcond
-
 /*
  * TODO: Implement
  * KAT_FD  unsigned all_satisfy(unsigned int predicate, unsigned* scratch_area);
@@ -71,8 +59,8 @@ KAT_FD void share_per_warp_data(
 	T*  __restrict__  where_to_make_available,
 	unsigned          writing_lane_id)
 {
-	if (lane::index() == writing_lane_id) {
-		where_to_make_available[warp::id()] = datum;
+	if (lane::id() == writing_lane_id) {
+		where_to_make_available[kat::warp::id()] = datum;
 	}
 	if (Synchronize) __syncthreads();
 }
@@ -144,18 +132,6 @@ namespace linear_grid {
 namespace collaborative {
 namespace block {
 
-///@cond
-// If we want to refer to other collaboration primitives, we'll make those references explicit;
-// but we do want to be able to say `warp::id()` without prefixing that with anything.
-
-namespace grid   = grid_info::grid;
-namespace block  = grid_info::block;
-namespace warp   = grid_info::warp;
-namespace thread = grid_info::thread;
-namespace lane   = grid_info::lane;
-
-///@endcond
-
 /*
  * TODO: Implement
  * KAT_FD  unsigned all_satisfy(unsigned int predicate, unsigned* scratch_area);
@@ -206,8 +182,9 @@ KAT_FD void share_per_warp_data(
 	T*  __restrict__  where_to_make_available,
 	unsigned          writing_lane_id)
 {
-	if (lane::index() == writing_lane_id) {
-		where_to_make_available[warp::id()] = datum;
+	namespace gi = kat::linear_grid;
+	if (gi::lane::id() == writing_lane_id) {
+		where_to_make_available[gi::warp::id()] = datum;
 	}
 	if (Synchronize) __syncthreads();
 }

src/kat/on_device/collaboration/grid.cuh

@@ -32,11 +32,11 @@ namespace grid {
 // If we want to refer to other primitives, we'll make those references explicit;
 // but we do want to be able to say `warp::index()` without prefixing that with anything.
 
-namespace grid   = kat::linear_grid::grid_info::grid;
-namespace block  = kat::linear_grid::grid_info::block;
-namespace warp   = kat::linear_grid::grid_info::warp;
-namespace thread = kat::linear_grid::grid_info::thread;
-namespace lane   = kat::linear_grid::grid_info::lane;
+namespace grid   = kat::linear_grid::grid;
+namespace block  = kat::linear_grid::block;
+namespace warp   = kat::linear_grid::warp;
+namespace thread = kat::linear_grid::thread;
+namespace lane   = kat::linear_grid::lane;
 
 /**
  * Have all kernel threads perform some action over the linear range
@@ -92,7 +92,7 @@ namespace warp_per_input_element {
  *
  * @note This version of `at_grid_stride` is specific to linear grids,
  * even though the text of its code looks the same as that of
- * @ref kat::grid_info::collaborative::warp::at_grid_stride .
+ * @ref kat::collaborative::warp::at_grid_stride .
  *
  * @param length The length of the range of positions on which to act
  * @param f The callable for warps to use each position in the sequence

src/kat/on_device/collaboration/warp.cuh

@@ -78,8 +78,6 @@ KAT_FD int last_lane_in(lane_mask_t mask)
 namespace collaborative {
 namespace warp {
 
-namespace lane   = grid_info::lane;
-
 /**
  * @brief Guarantees all memory writes by (mask-specified) warp lanes are visible to the other (mask-specified) lanes.
  *
@@ -296,13 +294,13 @@ KAT_FD T get_from_lane(T value, int source_lane)
 template <typename T>
 KAT_FD T get_from_first_lane(T value)
 {
-	return get_from_lane(value, grid_info::warp::first_lane);
+	return get_from_lane(value, kat::warp::first_lane);
 }
 
 template <typename T>
 KAT_FD T get_from_last_lane(T value)
 {
-	return get_from_lane(value, grid_info::warp::last_lane);
+	return get_from_lane(value, kat::warp::last_lane);
 }
 
 /**
@@ -422,13 +420,13 @@ KAT_FD typename std::result_of<Function()>::type have_a_single_lane_compute(Func
 template <typename Function>
 KAT_FD typename std::result_of<Function()>::type have_first_lane_compute(Function f)
 {
-	return have_a_single_lane_compute<Function>(f, grid_info::warp::first_lane);
+	return have_a_single_lane_compute<Function>(f, kat::warp::first_lane);
 }
 
 template <typename Function>
 KAT_FD typename std::result_of<Function()>::type have_last_lane_compute(Function f)
 {
-	return have_a_single_lane_compute<Function>(f, grid_info::warp::last_lane);
+	return have_a_single_lane_compute<Function>(f, kat::warp::last_lane);
 }
 
 KAT_FD unsigned index_among_active_lanes()
@@ -554,11 +552,11 @@ KAT_FD search_result_t<T> multisearch(const T& lane_needle, const T& lane_hay_st
 		unsigned lower, upper; // lower  is inclusive, upper is exclusive
 	} bounds;
 	if (lane_needle <= lane_hay_straw) {
-		bounds.lower = grid_info::warp::first_lane;
-		bounds.upper = grid_info::lane::id();
+		bounds.lower = kat::warp::first_lane;
+		bounds.upper = kat::lane::id();
 	}
 	else {
-		bounds.lower = grid_info::lane::id() + 1;
+		bounds.lower = lane::id() + 1;
 		bounds.upper = warp_size;
 	}
 	enum : unsigned { cutoff_to_linear_search = 6 };
@@ -680,7 +678,7 @@ template <
 	auto full_warp_writes_output_length = (PossibilityOfSlack == detail::has_no_slack) ?
 		full_warp_reads_output_length :
 		round_down_to_full_warps(full_warp_reads_output_length);
-	const auto lane_index = grid_info::lane::id();
+	const auto lane_index = lane::id();
 
 
 	promoted_size_t<Size> input_pos = lane_index;

src/kat/on_device/grid_info.cuh

@@ -175,8 +175,6 @@ constexpr KAT_FHD bool dimensionality_is_canonical(dimensions_t dims)
 
 // TODO: Perhaps have functions for strided copy in and out
 
-namespace grid_info {
-
 namespace detail {
 
 template <unsigned Dimensionality = 3>
@@ -267,7 +265,7 @@ KAT_FD grid_block_dimension_t size()        { return size(dimensions()); }
 
 KAT_FD position_t   first_thread_position() { return position_t{0, 0, 0}; }
 template <unsigned Dimensionality = 3>
-KAT_FD position_t   last_thread_position()  { return grid_info::detail::last_position_for(blockDim); }
+KAT_FD position_t   last_thread_position()  { return detail::last_position_for(blockDim); }
 
 template <unsigned Dimensionality = 3>
 KAT_FD grid_block_dimension_t
@@ -306,7 +304,7 @@ namespace grid {
  */
 KAT_FD bool         is_linear()
 {
-	return gridDim.y == 1 and gridDim.z == 1 and grid_info::block::is_linear();
+	return gridDim.y == 1 and gridDim.z == 1 and block::is_linear();
 }
 
 // TODO: Consider templatizing this on the dimensions too
@@ -428,11 +426,11 @@ KAT_FD position_t  global_index()          { return position_in_grid<OuterDimens
 namespace warp {
 
 template <unsigned Dimensionality = 3>
-KAT_FD unsigned    id_in_block()           { return grid_info::thread::id_in_block<Dimensionality>() / warp_size; }
+KAT_FD unsigned    id_in_block()           { return thread::id_in_block<Dimensionality>() / warp_size; }
 template <unsigned Dimensionality = 3>
 KAT_FD unsigned    index_in_block()        { return id_in_block<Dimensionality>(); }
 template <unsigned OuterDimensionality = 3, unsigned InnerDimensionality = 3>
-KAT_FD unsigned    id_in_grid()            { return grid_info::thread::id_in_grid<OuterDimensionality, InnerDimensionality>() / warp_size; }
+KAT_FD unsigned    id_in_grid()            { return thread::id_in_grid<OuterDimensionality, InnerDimensionality>() / warp_size; }
 template <unsigned Dimensionality>
 KAT_FD unsigned    index()                 { return index_in_block<Dimensionality>(); }
 template <unsigned OuterDimensionality = 3, unsigned InnerDimensionality = 3>
@@ -531,13 +529,10 @@ KAT_FD bool     is_last_in_warp()   { return lane::id<Dimensionality>() == warp:
 
 } // namespace thread
 
-} // namespace grid_info
 
 // I couldn't use '1d, 2d, 3d since those aren't valid identifiers...
 namespace linear_grid {
 
-namespace grid_info {
-
 namespace grid {
 
 // TODO: Should we use the same return types as for the non-linear case?
@@ -556,7 +551,7 @@ KAT_FD grid_dimension_t  first_last_position()     { return index_of_last_block(
 
 namespace block {
 
-using kat::grid_info::block::dimensions;
+using kat::block::dimensions;
 KAT_FD unsigned                index_in_grid()           { return blockIdx.x; }
 KAT_FD grid_block_dimension_t  index()                   { return index_in_grid(); }
 KAT_FD unsigned                id_in_grid()              { return index_in_grid(); }
@@ -605,10 +600,10 @@ KAT_FD unsigned     num_warps_per_block()   { return block::num_warps(); }
 
 namespace warp {
 
-using kat::grid_info::warp::first_lane;
-using kat::grid_info::warp::last_lane;
-using kat::grid_info::warp::size;
-using kat::grid_info::warp::length;
+using kat::warp::first_lane;
+using kat::warp::last_lane;
+using kat::warp::size;
+using kat::warp::length;
 
 }
 
@@ -632,8 +627,8 @@ KAT_FD bool                    is_last_in_block()  { return index_in_block() ==
 KAT_FD bool                    is_first_in_grid()  { return block::is_first_in_grid() and thread::is_first_in_block(); }
 KAT_FD bool                    is_last_in_grid()   { return block::is_last_in_grid() and thread::is_last_in_block();   }
 
-using ::kat::grid_info::thread::is_first_in_warp;
-using ::kat::grid_info::thread::is_last_in_warp;
+using ::kat::thread::is_first_in_warp;
+using ::kat::thread::is_last_in_warp;
 
 
 /**
@@ -714,7 +709,7 @@ namespace lane {
 // directly - we have to separate the code for the linear-grid and
 // non-linear-grid cases.
 
-enum { half_warp_size = kat::grid_info::lane::half_warp_size };
+enum { half_warp_size = kat::lane::half_warp_size };
 
 KAT_FD unsigned id_of(unsigned thread_index)
 {
@@ -750,8 +745,6 @@ KAT_FD unsigned is_in_second_half_warp()  { return id_in_warp() >= half_warp_siz
 
 } // namespace lane
 
-} // namespace grid_info
-
 } // namespace linear_grid
 
 } // namespace kat

src/kat/on_device/ranges.cuh

@@ -28,7 +28,7 @@ template <typename Size>
 KAT_DEV kat::ranges::strided<Size> warp_stride(Size length)
 {
 	constexpr const auto stride = warp_size;
-	auto begin = grid_info::lane::id();
+	auto begin = lane::id();
 	return ::kat::ranges::strided<promoted_size_t<Size>>(begin, length, stride);
 }
 
@@ -58,7 +58,7 @@ template <typename Size>
 KAT_DEV kat::ranges::strided<Size> warp_stride(Size length)
 {
 	constexpr const auto stride = warp_size;
-	auto begin = grid_info::lane::id();
+	auto begin = lane::id();
 	return ::kat::ranges::strided<promoted_size_t<Size>>(begin, length, stride);
 }
 
@@ -76,8 +76,8 @@ KAT_DEV kat::ranges::strided<Size> warp_stride(Size length)
 template <typename Size>
 KAT_DEV kat::ranges::strided<Size> block_stride(Size length)
 {
-	const auto stride = linear_grid::grid_info::block::size();
-	const auto begin = grid_info::thread::id_in_block();
+	const auto stride = linear_grid::block::size();
+	const auto begin = thread::id_in_block();
 	return ::kat::ranges::strided<promoted_size_t<Size>>(begin, length, stride);
 }
 
@@ -94,8 +94,8 @@ KAT_DEV kat::ranges::strided<Size> block_stride(Size length)
 template <typename Size>
 KAT_DEV kat::ranges::strided<Size> grid_stride(Size length)
 {
-	const auto stride = grid_info::grid::total_size();
-	const auto begin = grid_info::thread::global_id();
+	const auto stride = grid::total_size();
+	const auto begin = thread::global_id();
 	return ::kat::ranges::strided<promoted_size_t<Size>>(begin, length, stride);
 }
 
@@ -104,8 +104,8 @@ namespace warp_per_input_element {
 template <typename Size>
 KAT_DEV kat::ranges::strided<Size> grid_stride(Size length)
 {
-	const auto stride = grid_info::grid::num_warps();
-	const auto begin = grid_info::warp::global_id();
+	const auto stride = grid::num_warps();
+	const auto begin = warp::global_id();
 	return ::kat::ranges::strided<promoted_size_t<Size>>(begin, length, stride);
 }