cuDDL
Loading...
Searching...
No Matches
Public Types | Public Member Functions | Static Public Member Functions | List of all members
cuddl::detail::sketch_view< K, BucketCount, Layout > Class Template Reference

Non-owning implementation view of a DDL sketch. More...

Public Types

using register_type = uint32_t
 
using layout_type = Layout
 

Public Member Functions

__host__ __device__ constexpr sketch_view (device_span< register_type > registers, cuda::std::optional< blacklist_view > blacklist=cuda::std::nullopt) noexcept
 Constructs a reference over registers.
 
__host__ __device__ constexpr device_span< register_type > data () const noexcept
 Device registers.
 
Result< void > clear_async (cuda::stream_ref stream) const noexcept
 Resets every register to zero.
 
Result< void > add_async (device_span< uint64_t const > input, cuda::stream_ref stream) const noexcept
 Accumulates packed device k-mers into the existing sketch without clearing it.
 
Result< void > add_sequence_async (device_span< char const > sequence, cuda::stream_ref stream) const noexcept
 Accumulates device-resident raw ASCII bases without clearing.
 
template<bool IncludeCardinality = false>
Result< void > summary_async (sketch_view other, pairwise_summary &output, cuda::stream_ref stream) const noexcept
 Computes the raw pairwise summary into caller-owned device storage.
 
Result< void > cardinality_async (uint64_t *empty_out, double *estimate_out, cuda::stream_ref stream) const noexcept
 Computes this sketch's cardinality reduction on the GPU.
 
Result< void > hybrid_cardinality_async (hybrid_cardinality_estimates *output, cuda::stream_ref stream) const noexcept
 Computes BBTools and paper-style HybridDDL estimates in one GPU register scan.
 

Static Public Member Functions

static constexpr size_t bucket_count () noexcept
 Number of registers in the sketch.
 
static constexpr uint32_t kmer_length () noexcept
 Compile-time k-mer length.
 

Detailed Description

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
class cuddl::detail::sketch_view< K, BucketCount, Layout >

Non-owning implementation view of a DDL sketch.

Provides allocation-free, stream-ordered operations on an external allocation of BucketCount uint32_t registers, each holding its bucket's winning score. Pass by value into device or host code.

Definition at line 25 of file sketch_view.cuh.

Member Typedef Documentation

◆ layout_type

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
using cuddl::detail::sketch_view< K, BucketCount, Layout >::layout_type = Layout

Definition at line 32 of file sketch_view.cuh.

◆ register_type

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
using cuddl::detail::sketch_view< K, BucketCount, Layout >::register_type = uint32_t

Definition at line 31 of file sketch_view.cuh.

Constructor & Destructor Documentation

◆ sketch_view()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
__host__ __device__ constexpr cuddl::detail::sketch_view< K, BucketCount, Layout >::sketch_view ( device_span< register_type >  registers,
cuda::std::optional< blacklist_view >  blacklist = cuda::std::nullopt 
)
inlineexplicitconstexprnoexcept

Constructs a reference over registers.

Definition at line 35 of file sketch_view.cuh.

39 : registers_(registers), blacklist_(blacklist) {}
__global__ size_t uint32_t bool blacklist_view blacklist
Definition kernels.cuh:96
__global__ size_t uint32_t * registers
Definition kernels.cuh:94

Member Function Documentation

◆ add_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
Result< void > cuddl::detail::sketch_view< K, BucketCount, Layout >::add_async ( device_span< uint64_t const >  input,
cuda::stream_ref  stream 
) const
inlinenoexcept

Accumulates packed device k-mers into the existing sketch without clearing it.

The input must remain valid until stream completes.

Definition at line 67 of file sketch_view.cuh.

67 {
68 return detail::launch_construction<BucketCount, Layout>(
69 input, registers_, stream, blacklist_
70 );
71 }

◆ add_sequence_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
Result< void > cuddl::detail::sketch_view< K, BucketCount, Layout >::add_sequence_async ( device_span< char const >  sequence,
cuda::stream_ref  stream 
) const
inlinenoexcept

Accumulates device-resident raw ASCII bases without clearing.

Expects raw contiguous ASCII (no FASTA/newline stripping). Canonicalises case-insensitively with the packed path; ambiguity breaks windows and emits no k-mer across it. Accumulates without reset; short input below K is a no-op. Windows per call are capped at UINT32_MAX, so size must not exceed UINT32_MAX + K - 1. No hidden carry: callers supply the K-1 overlap within one record and no overlap between distinct records. No host copies. The input must remain valid until stream completes.

Definition at line 82 of file sketch_view.cuh.

82 {
83 return detail::launch_sequence_add<BucketCount, Layout>(
84 sequence, K, registers_, stream, blacklist_
85 );
86 }

◆ bucket_count()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
static constexpr size_t cuddl::detail::sketch_view< K, BucketCount, Layout >::bucket_count ( )
inlinestaticconstexprnoexcept

Number of registers in the sketch.

Definition at line 47 of file sketch_view.cuh.

47 {
48 return BucketCount;
49 }

◆ cardinality_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
Result< void > cuddl::detail::sketch_view< K, BucketCount, Layout >::cardinality_async ( uint64_t *  empty_out,
double *  estimate_out,
cuda::stream_ref  stream 
) const
inlinenoexcept

Computes this sketch's cardinality reduction on the GPU.

empty_out receives the empty-register count and estimate_out the cardinality estimate.

Definition at line 107 of file sketch_view.cuh.

111 {
112 detail::cardinality_kernel<BucketCount, Layout><<<1, detail::block_size, 0, stream.get()>>>(
113 registers_.data(), empty_out, estimate_out
114 );
115 return cuda_try(cudaGetLastError());
116 }
constexpr uint32_t block_size
Threads per CTA for the single-pair and cardinality reduction kernels.
Definition kernels.cuh:28
Result< void > cuda_try(cudaError_t error, std::source_location location=std::source_location::current())
Checks a CUDA runtime call and returns an error on failure.
Definition error.hpp:187
Here is the call graph for this function:
Here is the caller graph for this function:

◆ clear_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
Result< void > cuddl::detail::sketch_view< K, BucketCount, Layout >::clear_async ( cuda::stream_ref  stream) const
inlinenoexcept

Resets every register to zero.

Definition at line 57 of file sketch_view.cuh.

57 {
58 return cuda_try([&] {
59 cuda::fill_bytes(stream, cuda::std::span{registers_.data(), registers_.size()}, 0);
60 });
61 }
Here is the call graph for this function:

◆ data()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
__host__ __device__ constexpr device_span< register_type > cuddl::detail::sketch_view< K, BucketCount, Layout >::data ( ) const
inlineconstexprnoexcept

Device registers.

Definition at line 42 of file sketch_view.cuh.

42 {
43 return registers_;
44 }

◆ hybrid_cardinality_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
Result< void > cuddl::detail::sketch_view< K, BucketCount, Layout >::hybrid_cardinality_async ( hybrid_cardinality_estimates *  output,
cuda::stream_ref  stream 
) const
inlinenoexcept

Computes BBTools and paper-style HybridDDL estimates in one GPU register scan.

Definition at line 119 of file sketch_view.cuh.

122 {
123 detail::hybrid_cardinality_kernel<BucketCount, Layout>
124 <<<1, detail::block_size, 0, stream.get()>>>(registers_.data(), output);
125 return cuda_try(cudaGetLastError());
126 }
Here is the call graph for this function:

◆ kmer_length()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
static constexpr uint32_t cuddl::detail::sketch_view< K, BucketCount, Layout >::kmer_length ( )
inlinestaticconstexprnoexcept

Compile-time k-mer length.

Definition at line 52 of file sketch_view.cuh.

52 {
53 return K;
54 }

◆ summary_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
template<bool IncludeCardinality = false>
Result< void > cuddl::detail::sketch_view< K, BucketCount, Layout >::summary_async ( sketch_view< K, BucketCount, Layout >  other,
pairwise_summary &  output,
cuda::stream_ref  stream 
) const
inlinenoexcept

Computes the raw pairwise summary into caller-owned device storage.

output must point at device memory valid until stream completes.

Definition at line 92 of file sketch_view.cuh.

96 {
97 detail::summary_kernel<BucketCount, IncludeCardinality, Layout>
98 <<<1, detail::block_size, 0, stream.get()>>>(
99 registers_.data(), other.registers_.data(), output
100 );
101 return cuda_try(cudaGetLastError());
102 }
Here is the call graph for this function:

The documentation for this class was generated from the following file: