975c0f2 · opencog-SimpleGPUHashTable

Changed files (4) +226 -46

src/linearprobing.cu +136 -4

…	…	@@ -29,6 +29,7 @@ KeyValue* create_hashtable(uint32_t capacity)
29	29	return hashtable;
30	30	}
31	31
	32	// Insert the key/values in kvs into the hashtable
32	33	__global__ void gpu_hashtable_insert(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
33	34	{
34	35	unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
…	…	@@ -44,7 +45,7 @@ __global__ void gpu_hashtable_insert(KeyValue* hashtable, const KeyValue* kvs, u
44	45	if (prev == kEmpty \|\| prev == key)
45	46	{
46	47	hashtable[slot].value = value;
47		break;
	48	return;
48	49	}
49	50
50	51	slot = (slot + 1) & (kHashTableCapacity-1);
…	…	@@ -88,15 +89,145 @@ void insert_hashtable(KeyValue* pHashTable, const KeyValue* kvs, uint32_t num_kv
88	89	cudaFree(device_kvs);
89	90	}
90	91
91		__global__ void gpu_iterate_hashtable(KeyValue* pHashTable, KeyValue* kvs, uint32_t* kvs_size)
	92	// Lookup keys in the hashtable, and return the values
	93	__global__ void gpu_hashtable_lookup(KeyValue* hashtable, KeyValue* kvs, unsigned int numkvs)
	94	{
	95	unsigned int threadid = blockIdx.x * blockDim.x + threadIdx.x;
	96	if (threadid < kHashTableCapacity)
	97	{
	98	uint32_t key = kvs[threadid].key;
	99	uint32_t slot = hash(key);
	100
	101	while (true)
	102	{
	103	if (hashtable[slot].key == key)
	104	{
	105	kvs[threadid].value = hashtable[slot].value;
	106	return;
	107	}
	108	if (hashtable[slot].key == kEmpty)
	109	{
	110	kvs[threadid].value = kEmpty;
	111	return;
	112	}
	113	slot = (slot + 1) & (kHashTableCapacity - 1);
	114	}
	115	}
	116	}
	117
	118	void lookup_hashtable(KeyValue* pHashTable, KeyValue* kvs, uint32_t num_kvs)
	119	{
	120	// Copy the keyvalues to the GPU
	121	KeyValue* device_kvs;
	122	cudaMalloc(&device_kvs, sizeof(KeyValue) * num_kvs);
	123	cudaMemcpy(device_kvs, kvs, sizeof(KeyValue) * num_kvs, cudaMemcpyHostToDevice);
	124
	125	// Have CUDA calculate the thread block size
	126	int mingridsize;
	127	int threadblocksize;
	128	cudaOccupancyMaxPotentialBlockSize(&mingridsize, &threadblocksize, gpu_hashtable_insert, 0, 0);
	129
	130	// Create events for GPU timing
	131	cudaEvent_t start, stop;
	132	cudaEventCreate(&start);
	133	cudaEventCreate(&stop);
	134
	135	cudaEventRecord(start);
	136
	137	// Insert all the keys into the hash table
	138	int gridsize = ((uint32_t)num_kvs + threadblocksize - 1) / threadblocksize;
	139	gpu_hashtable_insert << <gridsize, threadblocksize >> > (pHashTable, device_kvs, (uint32_t)num_kvs);
	140
	141	cudaEventRecord(stop);
	142
	143	cudaEventSynchronize(stop);
	144
	145	float milliseconds = 0;
	146	cudaEventElapsedTime(&milliseconds, start, stop);
	147	float seconds = milliseconds / 1000.0f;
	148	printf(" GPU lookup %d items in %f ms (%f million keys/second)\n",
	149	num_kvs, milliseconds, num_kvs / (double)seconds / 1000000.0f);
	150
	151	cudaFree(device_kvs);
	152	}
	153
	154	// Delete each key in kvs from the hash table, if the key exists
	155	// A deleted key is left in the hash table, but its value is set to kEmpty
	156	// Deleted keys are not reused; once a key is assigned a slot, it never moves
	157	__global__ void gpu_hashtable_delete(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
	158	{
	159	unsigned int threadid = blockIdx.x * blockDim.x + threadIdx.x;
	160	if (threadid < kHashTableCapacity)
	161	{
	162	uint32_t key = kvs[threadid].key;
	163	uint32_t slot = hash(key);
	164
	165	while (true)
	166	{
	167	if (hashtable[slot].key == key)
	168	{
	169	hashtable[slot].value = kEmpty;
	170	return;
	171	}
	172	if (hashtable[slot].key == kEmpty)
	173	{
	174	return;
	175	}
	176	slot = (slot + 1) & (kHashTableCapacity - 1);
	177	}
	178	}
	179	}
	180
	181	void delete_hashtable(KeyValue* pHashTable, const KeyValue* kvs, uint32_t num_kvs)
	182	{
	183	// Copy the keyvalues to the GPU
	184	KeyValue* device_kvs;
	185	cudaMalloc(&device_kvs, sizeof(KeyValue) * num_kvs);
	186	cudaMemcpy(device_kvs, kvs, sizeof(KeyValue) * num_kvs, cudaMemcpyHostToDevice);
	187
	188	// Have CUDA calculate the thread block size
	189	int mingridsize;
	190	int threadblocksize;
	191	cudaOccupancyMaxPotentialBlockSize(&mingridsize, &threadblocksize, gpu_hashtable_insert, 0, 0);
	192
	193	// Create events for GPU timing
	194	cudaEvent_t start, stop;
	195	cudaEventCreate(&start);
	196	cudaEventCreate(&stop);
	197
	198	cudaEventRecord(start);
	199
	200	// Insert all the keys into the hash table
	201	int gridsize = ((uint32_t)num_kvs + threadblocksize - 1) / threadblocksize;
	202	gpu_hashtable_delete<< <gridsize, threadblocksize >> > (pHashTable, device_kvs, (uint32_t)num_kvs);
	203
	204	cudaEventRecord(stop);
	205
	206	cudaEventSynchronize(stop);
	207
	208	float milliseconds = 0;
	209	cudaEventElapsedTime(&milliseconds, start, stop);
	210	float seconds = milliseconds / 1000.0f;
	211	printf(" GPU delete %d items in %f ms (%f million keys/second)\n",
	212	num_kvs, milliseconds, num_kvs / (double)seconds / 1000000.0f);
	213
	214	cudaFree(device_kvs);
	215	}
	216
	217	// Iterate over every item in the hashtable; return non-empty key/values
	218	__global__ void gpu_iterate_hashtable(KeyValue* pHashTable, KeyValue* kvs, uint32_t* kvs_size)
92	219	{
93	220	unsigned int threadid = blockIdx.x * blockDim.x + threadIdx.x;
94	221	if (threadid < kHashTableCapacity)
95	222	{
96	223	if (pHashTable[threadid].key != kEmpty)
97	224	{
98		uint32_t size = atomicAdd(kvs_size, 1);
99		kvs[size] = pHashTable[threadid];
	225	uint32_t value = pHashTable[threadid].value;
	226	if (value != kEmpty)
	227	{
	228	uint32_t size = atomicAdd(kvs_size, 1);
	229	kvs[size] = pHashTable[threadid];
	230	}
100	231	}
101	232	}
102	233	}
…	…	@@ -131,6 +262,7 @@ std::vector<KeyValue> iterate_hashtable(KeyValue* pHashTable)
131	262	return kvs;
132	263	}
133	264
	265	// Free the memory of the hashtable
134	266	void destroy_hashtable(KeyValue* pHashTable)
135	267	{
136	268	cudaFree(pHashTable);

src/linearprobing.h +4 -0

…	…	@@ -16,6 +16,10 @@ KeyValue* create_hashtable(uint32_t initialCapacity);
16	16
17	17	void insert_hashtable(KeyValue* hashtable, const KeyValue* kvs, uint32_t num_kvs);
18	18
	19	void lookup_hashtable(KeyValue* hashtable, KeyValue* kvs, uint32_t num_kvs);
	20
	21	void delete_hashtable(KeyValue* hashtable, const KeyValue* kvs, uint32_t num_kvs);
	22
19	23	std::vector<KeyValue> iterate_hashtable(KeyValue* hashtable);
20	24
21	25	void destroy_hashtable(KeyValue* hashtable);

src/main.cpp +64 -32

…	…	@@ -10,16 +10,16 @@
10	10
11	11	// Create random keys/values in the range [0, kEmpty)
12	12	// kEmpty is used to indicate an empty slot
13		std::vector<KeyValue> generate_random_keyvalues()
	13	std::vector<KeyValue> generate_random_keyvalues(uint32_t numkvs)
14	14	{
15	15	std::random_device rd;
16	16	std::mt19937 gen(rd()); // mersenne_twister_engine
17	17	std::uniform_int_distribution<uint32_t> dis(0, kEmpty - 1);
18	18
19	19	std::vector<KeyValue> kvs;
20		kvs.reserve(kNumKeyValues);
	20	kvs.reserve(numkvs);
21	21
22		for (int i = 0; i < kNumKeyValues; i++)
	22	for (uint32_t i = 0; i < numkvs; i++)
23	23	{
24	24	uint32_t rand0 = dis(gen);
25	25	uint32_t rand1 = dis(gen);
…	…	@@ -29,6 +29,19 @@ std::vector<KeyValue> generate_random_keyvalues()
29	29	return kvs;
30	30	}
31	31
	32	// return numshuffledkvs random items from kvs
	33	std::vector<KeyValue> shuffle_keyvalues(std::vector<KeyValue> kvs, uint32_t numshuffledkvs)
	34	{
	35	std::random_shuffle(kvs.begin(), kvs.end());
	36
	37	std::vector<KeyValue> shuffled_kvs;
	38	shuffled_kvs.resize(numshuffledkvs);
	39
	40	std::copy(kvs.begin(), kvs.begin() + numshuffledkvs, shuffled_kvs.begin());
	41
	42	return shuffled_kvs;
	43	}
	44
32	45	LARGE_INTEGER start_timer()
33	46	{
34	47	LARGE_INTEGER timer;
…	…	@@ -52,17 +65,24 @@ double get_elapsed_time(LARGE_INTEGER start)
52	65	return milliseconds;
53	66	}
54	67
55		void test_unordered_map(std::vector<KeyValue> all_kvs)
	68	void test_unordered_map(std::vector<KeyValue> insert_kvs, std::vector<KeyValue> delete_kvs)
56	69	{
57	70	LARGE_INTEGER timer = start_timer();
58	71
59		printf("Testing insertion into std::unordered_map...\n");
	72	printf("Timing std::unordered_map...\n");
60	73
61	74	{
62	75	std::unordered_map<uint32_t, uint32_t> kvs_map;
63		for (auto& kv : all_kvs) {
	76	for (auto& kv : insert_kvs)
	77	{
64	78	kvs_map[kv.key] = kv.value;
65	79	}
	80	for (auto& kv : delete_kvs)
	81	{
	82	auto i = kvs_map.find(kv.key);
	83	if (i != kvs_map.end())
	84	kvs_map.erase(i);
	85	}
66	86	}
67	87
68	88	double milliseconds = get_elapsed_time(timer);
…	…	@@ -71,45 +91,57 @@ void test_unordered_map(std::vector<KeyValue> all_kvs)
71	91	milliseconds, kNumKeyValues / seconds / 1000000.0f);
72	92	}
73	93
74		void test_correctness(std::vector<KeyValue>, std::vector<KeyValue>);
	94	void test_correctness(std::vector<KeyValue>, std::vector<KeyValue>, std::vector<KeyValue>);
75	95
76	96	int main()
77	97	{
78		printf("Testing insertion of %d elements into GPU hash table...\n", kNumKeyValues);
	98	while (true)
	99	{
	100	// Initialize keyvalue pairs with random numbers
	101	std::vector<KeyValue> insert_kvs = generate_random_keyvalues(kNumKeyValues);
	102	std::vector<KeyValue> delete_kvs = shuffle_keyvalues(insert_kvs, kNumKeyValues / 2);
79	103
80		// Initialize keyvalue pairs with random numbers
81		std::vector<KeyValue> all_kvs = generate_random_keyvalues();
	104	printf("Testing insertion/deletion of %d/%d elements into GPU hash table...\n",
	105	(uint32_t)insert_kvs.size(), (uint32_t)delete_kvs.size());
82	106
83		// Begin test
84		LARGE_INTEGER timer = start_timer();
	107	// Begin test
	108	LARGE_INTEGER timer = start_timer();
85	109
86		KeyValue* pHashTable = create_hashtable(kHashTableCapacity);
	110	KeyValue* pHashTable = create_hashtable(kHashTableCapacity);
87	111
88		// Insert all the elements; each batch is processed concurrently on the GPU
89		const uint32_t num_batch_insertions = 16;
90		static_assert(kNumKeyValues % num_batch_insertions == 0, "needs even divisor");
91		uint32_t num_kvs_in_batch = kNumKeyValues / num_batch_insertions;
92		for (int i = 0; i < num_batch_insertions; i++)
93		{
94		insert_hashtable(pHashTable, all_kvs.data() + i * num_kvs_in_batch, num_kvs_in_batch);
95		}
	112	// Insert items into the hash table
	113	const uint32_t num_insert_batches = 16;
	114	uint32_t num_inserts_per_batch = (uint32_t)insert_kvs.size() / num_insert_batches;
	115	for (uint32_t i = 0; i < num_insert_batches; i++)
	116	{
	117	insert_hashtable(pHashTable, insert_kvs.data() + i * num_inserts_per_batch, num_inserts_per_batch);
	118	}
	119
	120	// Delete items from the hash table
	121	const uint32_t num_delete_batches = 8;
	122	uint32_t num_deletes_per_batch = (uint32_t)delete_kvs.size() / num_delete_batches;
	123	for (uint32_t i = 0; i < num_delete_batches; i++)
	124	{
	125	delete_hashtable(pHashTable, delete_kvs.data() + i * num_deletes_per_batch, num_deletes_per_batch);
	126	}
96	127
97		// Get all the key-values from the hash table
98		std::vector<KeyValue> kvs = iterate_hashtable(pHashTable);
	128	// Get all the key-values from the hash table
	129	std::vector<KeyValue> kvs = iterate_hashtable(pHashTable);
99	130
100		destroy_hashtable(pHashTable);
	131	destroy_hashtable(pHashTable);
101	132
102		// Summarize results
103		double milliseconds = get_elapsed_time(timer);
104		double seconds = milliseconds / 1000.0f;
105		printf("Total time (including memory copies, readback, etc): %f ms (%f million keys/second)\n", milliseconds,
106		kNumKeyValues / seconds / 1000000.0f);
	133	// Summarize results
	134	double milliseconds = get_elapsed_time(timer);
	135	double seconds = milliseconds / 1000.0f;
	136	printf("Total time (including memory copies, readback, etc): %f ms (%f million keys/second)\n", milliseconds,
	137	kNumKeyValues / seconds / 1000000.0f);
107	138
108		test_unordered_map(all_kvs);
	139	test_unordered_map(insert_kvs, delete_kvs);
109	140
110		test_correctness(all_kvs, kvs);
	141	test_correctness(insert_kvs, delete_kvs, kvs);
111	142
112		printf("Success\n");
	143	printf("Success\n");
	144	}
113	145
114	146	return 0;
115	147	}

src/test.cpp +22 -10

…	…	@@ -8,7 +8,7 @@
8	8	#include "random"
9	9	#include "linearprobing.h"
10	10
11		void test_correctness(std::vector<KeyValue> all_kvs, std::vector<KeyValue> kvs)
	11	void test_correctness(std::vector<KeyValue> insert_kvs, std::vector<KeyValue> delete_kvs, std::vector<KeyValue> kvs)
12	12	{
13	13	printf("Testing that there are no duplicate keys...\n");
14	14	std::unordered_set<uint32_t> unique_keys;
…	…	@@ -26,21 +26,33 @@ void test_correctness(std::vector<KeyValue> all_kvs, std::vector<KeyValue> kvs)
26	26	unique_keys.insert(node->key);
27	27	}
28	28
29		printf("Building unordered_map from insertion list...\n");
	29	printf("Building unordered_map from original list...\n");
30	30	std::unordered_map<uint32_t, std::vector<uint32_t>> all_kvs_map;
31		for (int i = 0; i < all_kvs.size(); i++)
	31	for (int i = 0; i < insert_kvs.size(); i++)
32	32	{
33	33	if (i % 10000000 == 0)
34		printf(" Building %d/%d\n", i, (uint32_t)all_kvs.size());
	34	printf(" Inserting %d/%d\n", i, (uint32_t)insert_kvs.size());
35	35
36		auto iter = all_kvs_map.find(all_kvs[i].key);
	36	auto iter = all_kvs_map.find(insert_kvs[i].key);
37	37	if (iter == all_kvs_map.end())
38	38	{
39		all_kvs_map[all_kvs[i].key] = std::vector<uint32_t>({ all_kvs[i].value });
	39	all_kvs_map[insert_kvs[i].key] = std::vector<uint32_t>({ insert_kvs[i].value });
40	40	}
41	41	else
42	42	{
43		iter->second.push_back(all_kvs[i].value);
	43	iter->second.push_back(insert_kvs[i].value);
	44	}
	45	}
	46
	47	for (int i = 0; i < delete_kvs.size(); i++)
	48	{
	49	if (i % 10000000 == 0)
	50	printf(" Deleting %d/%d\n", i, (uint32_t)delete_kvs.size());
	51
	52	auto iter = all_kvs_map.find(delete_kvs[i].key);
	53	if (iter != all_kvs_map.end())
	54	{
	55	all_kvs_map.erase(iter);
44	56	}
45	57	}
46	58
…	…	@@ -50,7 +62,7 @@ void test_correctness(std::vector<KeyValue> all_kvs, std::vector<KeyValue> kvs)
50	62	exit(-1);
51	63	}
52	64
53		printf("Testing that each key/value in hashtable is in the insertion list...\n");
	65	printf("Testing that each key/value in hashtable is in the original list...\n");
54	66	for (uint32_t i = 0; i < kvs.size(); i++)
55	67	{
56	68	if (i % 10000000 == 0)
…	…	@@ -59,14 +71,14 @@ void test_correctness(std::vector<KeyValue> all_kvs, std::vector<KeyValue> kvs)
59	71	auto iter = all_kvs_map.find(kvs[i].key);
60	72	if (iter == all_kvs_map.end())
61	73	{
62		printf("Hashtable key not found in insertion list\n");
	74	printf("Hashtable key not found in original list\n");
63	75	exit(-1);
64	76	}
65	77
66	78	std::vector<uint32_t>& values = iter->second;
67	79	if (std::find(values.begin(), values.end(), kvs[i].value) == values.end())
68	80	{
69		printf("Hashtable value not found in insertion list\n");
	81	printf("Hashtable value not found in original list\n");
70	82	exit(-1);
71	83	}
72	84	}

opencog-SimpleGPUHashTable commits / 975c0f2

Changed files (4) +226 -46