   }
 
   /**
-   * @brief Constructor specialised for static temporary storage, initializing using the runs' offsets.
+   * \brief Constructor specialised for static temporary storage, initializing using the runs' offsets.
    */
   template <typename UserRunOffsetT>
   __device__ __forceinline__ BlockRunLengthDecode(ItemT (&run_values)[RUNS_PER_THREAD],

 * retrieving a "window" from the run-length decoded array. The window's offset can be specified and BLOCK_THREADS *
 * DECODED_ITEMS_PER_THREAD (i.e., referred to as window_size) decoded items from the specified window will be returned.
 *
 * @note: Trailing runs of length 0 are supported (i.e., they may only appear at the end of the run_lengths array).
 * A run of length zero may not be followed by a run length that is not zero.
 *
 * @tparam ItemT The data type of the items being run-length decoded
 * @tparam BLOCK_DIM_X The thread block length in threads along the X dimension
 * @tparam RUNS_PER_THREAD The number of consecutive runs that each thread contributes
 * @tparam DECODED_ITEMS_PER_THREAD The maximum number of decoded items that each thread holds
 * @tparam DecodedOffsetT Type used to index into the block's decoded items (large enough to hold the sum over all the
 * runs' lengths)
 * @tparam BLOCK_DIM_Y The thread block length in threads along the Y dimension
 * @tparam BLOCK_DIM_Z The thread block length in threads along the Z dimension
 */
template <typename ItemT,
          int BLOCK_DIM_X,

  }

  /**
   * \brief Run-length decodes the runs previously passed via a call to Init(...) and returns the run-length decoded
   * items in a blocked arrangement to \p decoded_items. If the number of run-length decoded items exceeds the
   * run-length decode buffer (i.e., <b>DECODED_ITEMS_PER_THREAD * BLOCK_THREADS</b>), only the items that fit within
   * the buffer are returned. Subsequent calls to <b>RunLengthDecode</b> adjusting \p from_decoded_offset can be
   * used to retrieve the remaining run-length decoded items. Calling __syncthreads() between any two calls to
   * <b>RunLengthDecode</b> is not required.
   *
   * \param[out] decoded_items The run-length decoded items to be returned in a blocked arrangement
   * \param[in] from_decoded_offset If invoked with from_decoded_offset that is larger than total_decoded_size results
   * in undefined behavior.
   */
  __device__ __forceinline__ void RunLengthDecode(ItemT (&decoded_items)[DECODED_ITEMS_PER_THREAD],
