sketch_view.cuh

Go to the documentation of this file.

#pragma once
 
#include <cuda/algorithm>
#include <cuda/std/cstdint>
#include <cuda/stream>
 
#include <cuddl/detail/construction.cuh>
#include <cuddl/detail/kernels.cuh>
#include <cuddl/detail/sequence_encode.cuh>
#include <cuddl/device_span.cuh>
#include <cuddl/error.hpp>
#include <cuddl/hybrid_cardinality.cuh>
#include <cuddl/pairwise_counts.cuh>
 
namespace cuddl::detail {
 
template <uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
class sketch_view {
    static_assert(K >= 1 && K <= 31);
    static_assert(BucketCount >= (size_t{1} << 11) && BucketCount <= (size_t{1} << 13));
    static_assert((BucketCount & (BucketCount - 1)) == 0);
 
   public:
    using register_type = uint32_t;
    using layout_type = Layout;
 
    __host__ __device__ constexpr explicit sketch_view(
        device_span<register_type> registers,
        cuda::std::optional<blacklist_view> blacklist = cuda::std::nullopt
    ) noexcept
        : registers_(registers), blacklist_(blacklist) {}
 
    [[nodiscard]] __host__ __device__ constexpr device_span<register_type> data() const noexcept {
        return registers_;
    }
 
    [[nodiscard]] static constexpr size_t bucket_count() noexcept {
        return BucketCount;
    }
 
    [[nodiscard]] static constexpr uint32_t kmer_length() noexcept {
        return K;
    }
 
    [[nodiscard]] Result<void> clear_async(cuda::stream_ref stream) const noexcept {
        return cuda_try([&] {
            cuda::fill_bytes(stream, cuda::std::span{registers_.data(), registers_.size()}, 0);
        });
    }
 
    [[nodiscard]] Result<void>
    add_async(device_span<uint64_t const> input, cuda::stream_ref stream) const noexcept {
        return detail::launch_construction<BucketCount, Layout>(
            input, registers_, stream, blacklist_
        );
    }
 
    [[nodiscard]] Result<void>
    add_sequence_async(device_span<char const> sequence, cuda::stream_ref stream) const noexcept {
        return detail::launch_sequence_add<BucketCount, Layout>(
            sequence, K, registers_, stream, blacklist_
        );
    }
 
    template <bool IncludeCardinality = false>
    [[nodiscard]] Result<void> summary_async(
        sketch_view other,
        pairwise_summary& output,
        cuda::stream_ref stream
    ) const noexcept {
        detail::summary_kernel<BucketCount, IncludeCardinality, Layout>
            <<<1, detail::block_size, 0, stream.get()>>>(
                registers_.data(), other.registers_.data(), output
            );
        return cuda_try(cudaGetLastError());
    }
 
    [[nodiscard]] Result<void> cardinality_async(
        uint64_t* empty_out,
        double* estimate_out,
        cuda::stream_ref stream
    ) const noexcept {
        detail::cardinality_kernel<BucketCount, Layout><<<1, detail::block_size, 0, stream.get()>>>(
            registers_.data(), empty_out, estimate_out
        );
        return cuda_try(cudaGetLastError());
    }
 
    [[nodiscard]] Result<void> hybrid_cardinality_async(
        hybrid_cardinality_estimates* output,
        cuda::stream_ref stream
    ) const noexcept {
        detail::hybrid_cardinality_kernel<BucketCount, Layout>
            <<<1, detail::block_size, 0, stream.get()>>>(registers_.data(), output);
        return cuda_try(cudaGetLastError());
    }
 
   private:
    device_span<register_type> registers_;
    cuda::std::optional<blacklist_view> blacklist_;
};
 
}  // namespace cuddl::detail