diff --git a/include/bloom_filter.cuh b/include/bloom_filter.cuh index 39a4c24..454ff2e 100644 --- a/include/bloom_filter.cuh +++ b/include/bloom_filter.cuh @@ -260,14 +260,14 @@ public: * \param[in] group cooperative group this operation is executed in * \param[out] flag whether the key was already inside the filter before insertion */ - template< + template< index_type CGSize_ = cg_size(), class = std::enable_if_t> - DEVICEQUALIFIER INLINEQUALIFIER - bool insert_and_query( + DEVICEQUALIFIER INLINEQUALIFIER + bool insert_and_query( const key_type key, const cg::thread_block_tile& group) noexcept - { + { const index_type slot_index = ((Hasher::hash(key+seed_) % num_blocks_) * cg_size() + group.thread_rank()) % num_slots_; @@ -294,7 +294,7 @@ public: { return true; } - } + } /*! \brief get number of bits (m) * \return number of bits (m) @@ -317,11 +317,11 @@ public: /*! \brief get number of blocks * \return number of blocks */ - HOSTDEVICEQUALIFIER INLINEQUALIFIER - index_type num_blocks() const noexcept - { - return num_blocks_; - } + HOSTDEVICEQUALIFIER INLINEQUALIFIER + index_type num_blocks() const noexcept + { + return num_blocks_; + } /*! \brief get number of hash functions (k) * \return number of hash functions (k) diff --git a/include/bucket_list_hash_table.cuh b/include/bucket_list_hash_table.cuh index b451bdf..50c314d 100644 --- a/include/bucket_list_hash_table.cuh +++ b/include/bucket_list_hash_table.cuh @@ -50,20 +50,20 @@ public: /*! \brief get empty key * \return empty key */ - HOSTDEVICEQUALIFIER INLINEQUALIFIER - static constexpr key_type empty_key() noexcept - { - return EmptyKey; - } + HOSTDEVICEQUALIFIER INLINEQUALIFIER + static constexpr key_type empty_key() noexcept + { + return EmptyKey; + } /*! \brief get tombstone key * \return tombstone key */ - HOSTDEVICEQUALIFIER INLINEQUALIFIER - static constexpr key_type tombstone_key() noexcept - { - return TombstoneKey; - } + HOSTDEVICEQUALIFIER INLINEQUALIFIER + static constexpr key_type tombstone_key() noexcept + { + return TombstoneKey; + } /*! \brief checks if \c key is equal to \c (EmptyKey||TombstoneKey) * \return \c bool @@ -612,11 +612,11 @@ public: * \param stream CUDA stream in which this operation is executed in * \return load factor */ - HOSTQUALIFIER INLINEQUALIFIER - float value_load_factor(const cudaStream_t stream = 0) const noexcept - { - return value_store_.load_factor(stream); - } + HOSTQUALIFIER INLINEQUALIFIER + float value_load_factor(const cudaStream_t stream = 0) const noexcept + { + return value_store_.load_factor(stream); + } /*! \brief get the the total number of bytes occupied by this data structure * \return bytes diff --git a/include/hash_set.cuh b/include/hash_set.cuh index 22775ca..5b459f2 100644 --- a/include/hash_set.cuh +++ b/include/hash_set.cuh @@ -467,31 +467,31 @@ public: * \param[in] stream CUDA stream in which this operation is executed in * \param[in] size of shared memory to reserve for this execution */ - template - HOSTQUALIFIER INLINEQUALIFIER - void for_each( - Func f, - cudaStream_t stream = 0, - index_type smem_bytes = 0) const noexcept - { - if(!is_initialized_) return; - - helpers::lambda_kernel - <<>> - ([=, *this] DEVICEQUALIFIER // TODO mutable? - { - const index_type tid = helpers::global_thread_id(); - - if(tid < capacity()) - { - const key_type key = keys_[tid]; - if(is_valid_key(key)) - { - f(key); - } - } - }); - } + template + HOSTQUALIFIER INLINEQUALIFIER + void for_each( + Func f, + cudaStream_t stream = 0, + index_type smem_bytes = 0) const noexcept + { + if(!is_initialized_) return; + + helpers::lambda_kernel + <<>> + ([=, *this] DEVICEQUALIFIER // TODO mutable? + { + const index_type tid = helpers::global_thread_id(); + + if(tid < capacity()) + { + const key_type key = keys_[tid]; + if(is_valid_key(key)) + { + f(key); + } + } + }); + } /*! \brief number of key/value pairs stored inside the hash set * \return the number of key/value pairs inside the hash table diff --git a/include/multi_bucket_hash_table.cuh b/include/multi_bucket_hash_table.cuh index 3f27ac3..a2e123a 100644 --- a/include/multi_bucket_hash_table.cuh +++ b/include/multi_bucket_hash_table.cuh @@ -897,7 +897,7 @@ public: * \return \c warpcore::HashSet */ HOSTQUALIFIER INLINEQUALIFIER - const key_set_type get_key_set( + key_set_type get_key_set( const cudaStream_t stream = 0, const float size_fraction = 0.9) const noexcept { diff --git a/include/multi_value_hash_table.cuh b/include/multi_value_hash_table.cuh index 593c845..943388d 100755 --- a/include/multi_value_hash_table.cuh +++ b/include/multi_value_hash_table.cuh @@ -201,13 +201,13 @@ public: * \param[in] probing_length maximum number of probing attempts * \return status (per thread) */ - DEVICEQUALIFIER INLINEQUALIFIER - status_type insert( - const key_type key_in, - const value_type& value_in, - const cg::thread_block_tile& group, - const index_type probing_length = defaults::probing_length()) noexcept - { + DEVICEQUALIFIER INLINEQUALIFIER + status_type insert( + const key_type key_in, + const value_type& value_in, + const cg::thread_block_tile& group, + const index_type probing_length = defaults::probing_length()) noexcept + { if(!is_initialized_) { return status_type::not_initialized(); @@ -290,7 +290,7 @@ public: status_type::probing_length_exceeded(); device_join_status(status); return status; - } + } /*! \brief insert a set of keys into the hash table * \tparam StatusHandler handles returned status per key (see \c status_handlers) @@ -669,7 +669,7 @@ public: * \return \c warpcore::HashSet */ HOSTQUALIFIER INLINEQUALIFIER - const key_set_type get_key_set( + key_set_type get_key_set( const cudaStream_t stream = 0, const float size_fraction = 0.9) const noexcept { diff --git a/include/single_value_hash_table.cuh b/include/single_value_hash_table.cuh index 7068b39..3bd1de9 100644 --- a/include/single_value_hash_table.cuh +++ b/include/single_value_hash_table.cuh @@ -207,13 +207,13 @@ public: * \param[in] probing_length maximum number of probing attempts * \return status (per thread) */ - DEVICEQUALIFIER INLINEQUALIFIER - status_type insert( - const key_type key_in, - const value_type& value_in, - const cg::thread_block_tile& group, - const index_type probing_length = defaults::probing_length()) noexcept - { + DEVICEQUALIFIER INLINEQUALIFIER + status_type insert( + const key_type key_in, + const value_type& value_in, + const cg::thread_block_tile& group, + const index_type probing_length = defaults::probing_length()) noexcept + { status_type status = status_type::unknown_error(); value_type * value_ptr = @@ -225,7 +225,7 @@ public: } return status; - } + } /*! \brief insert a set of keys into the hash table * \tparam StatusHandler handles returned status per key (see \c status_handlers) diff --git a/include/storage.cuh b/include/storage.cuh index ba17894..8752e2c 100644 --- a/include/storage.cuh +++ b/include/storage.cuh @@ -1375,26 +1375,26 @@ public: * \param[in] stream CUDA stream in which this operation is executed in * \return load factor */ - HOSTDEVICEQUALIFIER INLINEQUALIFIER - float load_factor(const cudaStream_t stream = 0) const noexcept - { - index_type load = 0; + HOSTDEVICEQUALIFIER INLINEQUALIFIER + float load_factor(const cudaStream_t stream = 0) const noexcept + { + index_type load = 0; - cudaMemcpyAsync( - &load, next_free_bucket_, sizeof(index_type), D2H, stream); + cudaMemcpyAsync( + &load, next_free_bucket_, sizeof(index_type), D2H, stream); - cudaStreamSynchronize(stream); + cudaStreamSynchronize(stream); - return float(load) / float(capacity()); - } + return float(load) / float(capacity()); + } /*! \brief get the number of occupied bytes * \param[in] stream CUDA stream in which this operation is executed in * \return bytes */ - HOSTDEVICEQUALIFIER INLINEQUALIFIER - index_type bytes_occupied(const cudaStream_t stream = 0) const noexcept - { + HOSTDEVICEQUALIFIER INLINEQUALIFIER + index_type bytes_occupied(const cudaStream_t stream = 0) const noexcept + { index_type occupied = 0; cudaMemcpyAsync( @@ -1403,7 +1403,7 @@ public: cudaStreamSynchronize(stream); return occupied * sizeof(bucket_type); - } + } /*! \brief get bucket growth factor * \return factor