reference_deltas.cuh File Reference

#include <cub/device/device_radix_sort.cuh>
#include <cub/device/device_scan.cuh>
#include <cub/device/device_select.cuh>
#include <cuda/buffer>
#include <cuda/iterator>
#include <cuda/std/cstdint>
#include <cuda/std/functional>
#include <cuda/stream>
#include <optional>
#include <cuddl/detail/kernels.cuh>
#include <cuddl/error.hpp>

Go to the source code of this file.

Classes
struct	cuddl::detail::reference_delta_view
	Reference rows stored as a base reference plus the buckets that differ from it. More...
struct	cuddl::detail::reference_deltas

Namespaces
namespace	cuddl
	Batch operations over contiguous device rows.
namespace	cuddl::detail

Functions
template<size_t BucketCount>
__global__	cuddl::detail::__launch_bounds__ (block_size) void build_score_planes_kernel(uint16_t const *rows
	Transposes compact score rows into the bit-plane layout of score_plane_index.
	cuddl::detail::for (auto i=static_cast< size_t >(blockIdx.x) *block_size+threadIdx.x;i< total;i+=static_cast< size_t >(gridDim.x) *block_size)
__global__	cuddl::detail::__launch_bounds__ (block_size) static void build_sparse_directory_kernel(uint16_t const *sorted_keys
	Builds a sparse key directory: entry c of bucket b is the first position in b's sorted keys whose key is at least c << shift.

Variables
constexpr uint32_t	cuddl::detail::delta_change_limit = 160U
	Most buckets in which a child reference may differ from its base reference.
constexpr uint32_t	cuddl::detail::delta_signatures = 8U
	Exact bucket signatures that propose base candidates, and buckets per signature.
constexpr uint32_t	cuddl::detail::delta_signature_buckets = 8U
constexpr uint32_t	cuddl::detail::delta_signature_candidates = 4U
	Highest-ID rows taken from each matching signature.
constexpr uint32_t	cuddl::detail::delta_candidates = delta_signatures * delta_signature_candidates
constexpr uint32_t	cuddl::detail::delta_correction_warps = 24U
	Warps per apply_reference_deltas_kernel block, and the most and fewest children one block cell takes.
constexpr uint32_t	cuddl::detail::delta_child_tile = 2048U
constexpr uint32_t	cuddl::detail::delta_min_child_tile = 4U * delta_correction_warps
template<size_t BucketCount>
constexpr uint32_t	cuddl::detail::delta_query_group = static_cast<uint32_t>(32768U / BucketCount)
	Queries one correction cell stages bucket-major: 64 KiB of scores for every bucket count.
template<size_t BucketCount>
constexpr size_t	cuddl::detail::delta_correction_shared_bytes
__global__ uint32_t uint32_t	cuddl::detail::signature
__global__ uint32_t uint32_t uint64_t *	cuddl::detail::keys
__global__ uint32_t uint32_t uint64_t uint32_t *	cuddl::detail::ids