hash_table.cuh File Reference

Implements kernel and device functions for a basic hash table. More...

#include "definitions.h"
#include "hash_table.h"
#include <driver_types.h>

Namespaces
	CudaHT
	Encapsulates the hash table library.

Functions
__device__ __host__ Entry	CudaHT::CuckooHashing::make_entry (unsigned key, unsigned value)
	Makes an 64-bit Entry out of a key-value pair for the hash table.
__device__ __host__ unsigned	CudaHT::CuckooHashing::get_key (Entry entry)
	Returns the key of an Entry.
__device__ __host__ unsigned	CudaHT::CuckooHashing::get_value (Entry entry)
	Returns the value of an Entry.
template<unsigned kNumHashFunctions>
__device__ unsigned	CudaHT::CuckooHashing::retrieve (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. More...
template<unsigned kNumHashFunctions>
__global__ void	CudaHT::CuckooHashing::hash_retrieve (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)
	Perform a retrieval from a basic hash table. Each thread manages a single query.
Internal
template<class T >
__global__ void	CudaHT::CuckooHashing::clear_table (const unsigned table_size, const T value, T *table)
	Fills the entire array with a specific value.
template<unsigned kNumHashFunctions>
__device__ void	CudaHT::CuckooHashing::KeyLocations (const Functions< kNumHashFunctions > constants, const unsigned table_size, const unsigned key, unsigned locations[kNumHashFunctions])
	Determine where in the hash table the key could be located.
template<unsigned kNumHashFunctions>
__device__ unsigned	CudaHT::CuckooHashing::determine_next_location (const Functions< kNumHashFunctions > constants, const unsigned table_size, const unsigned key, const unsigned previous_location)
	Determine where to insert the key next. The hash functions are used in round-robin order.
template<unsigned kNumHashFunctions>
__device__ bool	CudaHT::CuckooHashing::insert (const unsigned table_size, const Functions< kNumHashFunctions > constants, const uint2 stash_constants, const unsigned max_iteration_attempts, Entry *table, unsigned *stash_count, Entry entry, unsigned *iterations_used)
	Attempts to insert a single entry into the hash table. More...
template<unsigned kNumHashFunctions>
__global__ void	CudaHT::CuckooHashing::CuckooHash (const unsigned n_entries, const unsigned *keys, const unsigned *values, const unsigned table_size, const Functions< kNumHashFunctions > constants, const unsigned max_iteration_attempts, Entry *table, uint2 stash_constants, unsigned *stash_count, unsigned *failures, unsigned *iterations_taken)

Detailed Description

Implements kernel and device functions for a basic hash table.

Function Documentation

template<unsigned kNumHashFunctions>

__device__ unsigned CudaHT::CuckooHashing::retrieve	(	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.

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	The hash function used to build the stash
[in]	stash_count	The number of items in the stash
[out]	num_probes_required	Debug only: The number of probes required to resolve the query.

Returns

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

template<unsigned kNumHashFunctions>

__device__ bool CudaHT::CuckooHashing::insert	(	const unsigned	table_size,
		const Functions< kNumHashFunctions >	constants,
		const uint2	stash_constants,
		const unsigned	max_iteration_attempts,
		Entry *	table,
		unsigned *	stash_count,
		Entry	entry,
		unsigned *	iterations_used
	)

Attempts to insert a single entry into the hash table.

This process stops after a certain number of iterations. If the thread is still holding onto an item because of an eviction, it tries the stash. If it fails to enter the stash, it returns false. Otherwise, it succeeds and returns true.