warpcore 0.0.1
Hashing at the Speed of Light on modern CUDA-accelerators
base.cuh
Go to the documentation of this file.
1 #ifndef WARPCORE_BASE_CUH
2 #define WARPCORE_BASE_CUH
3 
4 // modified by Markus Vieth
5 
6 #include <cstdint>
7 #include <algorithm>
8 #include <assert.h>
9 #include <limits>
10 #include <cuda_runtime.h>
11 #include <cooperative_groups.h>
12 #include <cuda_helpers.cuh>
13 #include <packed_types.cuh>
14 #include <cub/cub.cuh>
15 
16 #include "primes.hpp"
17 
18 namespace warpcore
19 {
20 
21 namespace cg = cooperative_groups;
22 
23 using index_t = std::uint64_t;
24 using status_base_t = std::uint32_t;
25 
26 namespace detail
27 {
28 
29 HOSTQUALIFIER INLINEQUALIFIER
30 index_t get_valid_capacity(index_t min_capacity, index_t cg_size) noexcept
31 {
32  const auto x = SDIV(min_capacity, cg_size);
33  const auto y =
34  std::lower_bound(primes.begin(), primes.end(), x);
35  return (y == primes.end()) ? 0 : (*y) * cg_size;
36 }
37 
38 } // namespace detail
39 
40 } // namespace warpcore
41 
42 // TODO move to defaults and expose as constexpr
43 #ifdef __CUDACC_DEBUG__
44  #ifndef WARPCORE_BLOCKSIZE
45  #define WARPCORE_BLOCKSIZE 128
46  #endif
47 #else
48  #ifndef WARPCORE_BLOCKSIZE
49  #define WARPCORE_BLOCKSIZE MAXBLOCKSIZE // MAXBLOCKSIZE defined in cuda_helpers
50  #endif
51 #endif
52 
53 #include "tags.cuh"
54 #include "checks.cuh"
55 #include "status.cuh"
56 #include "hashers.cuh"
57 #include "probing_schemes.cuh"
58 #include "storage.cuh"
59 #include "defaults.cuh"
60 #include "gpu_engine.cuh"
61 
62 #endif /* WARPCORE_BASE_CUH */