hash_compacting.cu File Reference

Implements hash tables that assign each unique key an ID. More...

#include "debugging.h"
#include "hash_compacting.h"
#include "hash_functions.h"
#include "hash_table.cuh"
#include <cudpp.h>
#include "cuda_util.h"
#include <set>

Namespaces
	CudaHT
	Encapsulates the hash table library.

Functions
template<unsigned kNumHashFunctions>
__device__ unsigned	CudaHT::CuckooHashing::retrieve_compacting (const unsigned query_key, const unsigned table_size, const Entry *table, const Functions< kNumHashFunctions > constants, const uint2 stash_constants, const unsigned stash_count, unsigned *num_probes_required=NULL)
	Answers a single query from a compacting hash table. More...
template<unsigned kNumHashFunctions>
__global__ void	CudaHT::CuckooHashing::hash_retrieve_compacting (const unsigned n_queries, const unsigned *keys_in, const unsigned table_size, const Entry *table, const Functions< kNumHashFunctions > constants, const uint2 stash_constants, const unsigned stash_count, unsigned *values_out, unsigned *num_probes_required=NULL)
	Returns the unique identifier for every query key. Each thread manages a single query. More...
void	CudaHT::CuckooHashing::CUDAWrapper::CallHashBuildCompacting (const int n, const unsigned num_hash_functions, const unsigned *keys, const unsigned table_size, const Functions< 2 > constants_2, const Functions< 3 > constants_3, const Functions< 4 > constants_4, const Functions< 5 > constants_5, const uint2 stash_constants, const unsigned max_iteration_attempts, unsigned *table, unsigned *stash_count, unsigned *failures)
	Calls the compacting cuckoo hash construction kernel.
void	CudaHT::CuckooHashing::CUDAWrapper::CallHashRemoveDuplicates (const unsigned num_hash_functions, const unsigned table_size, const unsigned total_table_size, const Functions< 2 > constants_2, const Functions< 3 > constants_3, const Functions< 4 > constants_4, const Functions< 5 > constants_5, const uint2 stash_constants, unsigned *keys, unsigned *is_unique)
	Removes any duplicate keys from the table.
void	CudaHT::CuckooHashing::CUDAWrapper::CallHashCompactDown (const unsigned table_size, Entry *table_entry, unsigned *unique_keys, const unsigned *table, const unsigned *indices)
	Creates a compacted list of the unique keys and sets up the keys with their unique IDs.
void	CudaHT::CuckooHashing::CUDAWrapper::CallHashRetrieveCompacting (const unsigned n_queries, const unsigned num_hash_functions, const unsigned *keys_in, const unsigned table_size, const Entry *table, const Functions< 2 > constants_2, const Functions< 3 > constants_3, const Functions< 4 > constants_4, const Functions< 5 > constants_5, const uint2 stash_constants, const unsigned stash_count, unsigned *values_out)
	Calls the retrieval kernel.
void	CudaHT::CuckooHashing::CUDAWrapper::ClearTable (const unsigned slots_in_table, const unsigned fill_value, unsigned *d_contents)
	Fills an array with a particular value.
Internal
template<unsigned kNumHashFunctions>
__global__ void	CudaHT::CuckooHashing::hash_build_compacting (const int n, const unsigned *keys, const unsigned table_size, const Functions< kNumHashFunctions > constants, const uint2 stash_constants, const unsigned max_iteration_attempts, unsigned *table, unsigned *stash_count, unsigned *failures)
	Builds a compacting hash table.
template<unsigned kNumHashFunctions>
__global__ void	CudaHT::CuckooHashing::hash_remove_duplicates (const unsigned table_size, const unsigned total_table_size, const Functions< kNumHashFunctions > constants, const uint2 stash_constants, unsigned *keys, unsigned *is_unique)
	Removes all key duplicates from a compacting hash table. More...
__global__ void	CudaHT::CuckooHashing::hash_compact_down (const unsigned table_size, Entry *table_entry, unsigned *unique_keys, const unsigned *table, const unsigned *indices)
	Interleave the keys and their unique IDs in the cuckoo hash table, then compact down the keys.
Explicit template specializations
template<>
__global__ void	CudaHT::CuckooHashing::hash_remove_duplicates< 2 > (const unsigned table_size, const unsigned total_table_size, const Functions< 2 > constants, const uint2 stash_constants, unsigned *keys, unsigned *is_unique)
template<>
__global__ void	CudaHT::CuckooHashing::hash_remove_duplicates< 3 > (const unsigned table_size, const unsigned total_table_size, const Functions< 3 > constants, const uint2 stash_constants, unsigned *keys, unsigned *is_unique)
template<>
__global__ void	CudaHT::CuckooHashing::hash_remove_duplicates< 4 > (const unsigned table_size, const unsigned total_table_size, const Functions< 4 > constants, const uint2 stash_constants, unsigned *keys, unsigned *is_unique)
template<>
__global__ void	CudaHT::CuckooHashing::hash_remove_duplicates< 5 > (const unsigned table_size, const unsigned total_table_size, const Functions< 5 > constants, const uint2 stash_constants, unsigned *keys, unsigned *is_unique)

Detailed Description

Implements hash tables that assign each unique key an ID.

Function Documentation

template<unsigned kNumHashFunctions>

__device__ unsigned CudaHT::CuckooHashing::retrieve_compacting	(	const unsigned	query_key,
		const unsigned	table_size,
		const Entry *	table,
		const Functions< kNumHashFunctions >	constants,
		const uint2	stash_constants,
		const unsigned	stash_count,
		unsigned *	num_probes_required = NULL
	)

Answers a single query from a compacting hash table.

Parameters

[in]	key	Query key
[in]	table_size	Size of the hash table
[in]	table	The contents of the hash table
[in]	constants	The hash functions used to build the table
[in]	stash_constants	Constants used by the stash hash function
[in]	stash_count	Number of items contained in the stash
[out]	num_probes_required	Debug only: The number of probes required to resolve the query.

Returns

The ID of the query key is returned if the key exists in the table. Otherwise, kNotFound will be returned.

template<unsigned kNumHashFunctions>

__global__ void CudaHT::CuckooHashing::hash_retrieve_compacting	(	const unsigned	n_queries,
		const unsigned *	keys_in,
		const unsigned	table_size,
		const Entry *	table,
		const Functions< kNumHashFunctions >	constants,
		const uint2	stash_constants,
		const unsigned	stash_count,
		unsigned *	values_out,
		unsigned *	num_probes_required = NULL
	)

Returns the unique identifier for every query key. Each thread manages a single query.

Parameters

[in]	n_queries	Number of query keys
[in]	keys_in	Query keys
[in]	table_size	Size of the hash table
[in]	table	The contents of the hash table
[in]	constants	The hash functions used to build the table
[in]	stash_constants	Constants used by the stash hash function
[in]	stash_count	Number of items contained in the stash
[out]	values_out	The unique identifiers for each query key
[out]	num_probes_required	Debug only: The number of probes required to resolve the query.

The ID of the query key is written out if the key exists in the table. Otherwise, kNotFound will be.

template<unsigned kNumHashFunctions>

__global__ void CudaHT::CuckooHashing::hash_remove_duplicates	(	const unsigned	table_size,
		const unsigned	total_table_size,
		const Functions< kNumHashFunctions >	constants,
		const uint2	stash_constants,
		unsigned *	keys,
		unsigned *	is_unique
	)

Removes all key duplicates from a compacting hash table.

The unspecialized version is significantly slower than the explicitly specialized ones.