450895f · opencog-SimpleGPUHashTable

Changed files (5) +26 -6

README.md +22 -1

…	…	@@ -5,7 +5,8 @@
5	5	This project shows how to implement a simple GPU hash table. Thanks to the high bandwidth and massive parallelism of
6	6	GPU's, the result is a high performance hash table capable of hundreds of millions of operations per second.
7	7
8		The code achieves an average insertion rate of 309 million key/second on my development laptop with an NVIDIA GTX 1060.
	8	The code achieves an average insertion rate of 326 million key/second on my development laptop with an NVIDIA GTX 1060,
	9	measured by inserting 64 million elements.
9	10
10	11	The code implements a lock free hash table using linear probing. Concurrent inserts, deletes, and lookups are supported by
11	12	this hash table. The hash table works on 32 bit keys and 32 bit values (although 0xffffffff is reserved for both keys
…	…	@@ -26,6 +27,26 @@ For more information on the lock free hash table design, read:
26	27	Table](https://preshing.com/20130605/the-worlds-simplest-lock-free-hash-table/)
27	28	* Cliff Click's slides on [A Lock-Free Wait-Free Hash Table](https://web.stanford.edu/class/ee380/Abstracts/070221_LockFreeHash.pdf)
28	29
	30	# How To Use
	31
	32	If you build and run the executable, it enters an infinite loop of inserting and deleting random numbers into the
	33	GPU hash table and verifying that the results are correct. The seed used to generate random numbers changes every time
	34	you run the executable, but you can set the seed to a specific value in code if you'd like to reproduce results across
	35	runs.
	36
	37	This is how you insert a vector of `KeyValue` pairs into the hash table and then retrieve all the `KeyValue` pairs back:
	38
	39	```cpp
	40	std::vector<KeyValue> things_to_insert = { {0,1}, {1,2}, {2,3}, {3,4} };
	41
	42	KeyValue* pHashTable2 = create_hashtable(kHashTableCapacity);
	43	insert_hashtable(pHashTable2, things_to_insert.data(), (uint32_t)things_to_insert.size());
	44	std::vector<KeyValue> result = iterate_hashtable(pHashTable2);
	45	destroy_hashtable(pHashTable2);
	46	```
	47
	48	After that runs, the vectors `things_to_insert` and `result` should be the same, but possibly in a different order.
	49
29	50	# Prerequisites
30	51
31	52	* CMake

screenshot.png +0 -0

(no content changes — binary or rename)

src/linearprobing.cu +2 -3

…	…	@@ -14,9 +14,8 @@ __device__ uint32_t hash(uint32_t k)
14	14	return k & (kHashTableCapacity-1);
15	15	}
16	16
17		// Create a hash table. For linear probing, this is just an array
18		// of KeyValues. The hash table is
19		KeyValue* create_hashtable(uint32_t capacity)
	17	// Create a hash table. For linear probing, this is just an array of KeyValues
	18	KeyValue* create_hashtable()
20	19	{
21	20	// Allocate memory
22	21	KeyValue* hashtable;

src/linearprobing.h +1 -1

…	…	@@ -12,7 +12,7 @@ const uint32_t kNumKeyValues = kHashTableCapacity / 2;
12	12
13	13	const uint32_t kEmpty = 0xffffffff;
14	14
15		KeyValue* create_hashtable(uint32_t initialCapacity);
	15	KeyValue* create_hashtable();
16	16
17	17	void insert_hashtable(KeyValue* hashtable, const KeyValue* kvs, uint32_t num_kvs);
18	18

src/main.cpp +1 -1

…	…	@@ -107,7 +107,7 @@ int main()
107	107
108	108	Time timer = start_timer();
109	109
110		KeyValue* pHashTable = create_hashtable(kHashTableCapacity);
	110	KeyValue* pHashTable = create_hashtable();
111	111
112	112	// Insert items into the hash table
113	113	const uint32_t num_insert_batches = 16;

opencog-SimpleGPUHashTable commits / 450895f

Changed files (5) +26 -6