a2ae55c · opencog-SimpleGPUHashTable

Changed files (2) +61 -36

src/linearprobing.cu +35 -13

…	…	@@ -16,7 +16,8 @@ __device__ uint32_t hash(uint32_t k)
16	16
17	17	// Create a hash table. For linear probing, this is just an array
18	18	// of KeyValues. The hash table is
19		KeyValue* create_hashtable(uint32_t capacity) {
	19	KeyValue* create_hashtable(uint32_t capacity)
	20	{
20	21	// Allocate memory
21	22	KeyValue* hashtable;
22	23	cudaMalloc(&hashtable, sizeof(KeyValue) * kHashTableCapacity);
…	…	@@ -87,24 +88,45 @@ void insert_hashtable(KeyValue* pHashTable, const KeyValue* kvs, uint32_t num_kv
87	88	cudaFree(device_kvs);
88	89	}
89	90
90		std::vector<KeyValue> iterate_hashtable(KeyValue* pHashTable)
	91	__global__ void gpu_iterate_hashtable(KeyValue* pHashTable, KeyValue* kvs, uint32_t* kvs_size)
91	92	{
92		KeyValue* pHostHashTable;
93		cudaHostAlloc(&pHostHashTable, sizeof(KeyValue) * kHashTableCapacity, 0);
94		cudaMemcpy(pHostHashTable, pHashTable, sizeof(KeyValue) * kHashTableCapacity, cudaMemcpyDeviceToHost);
95
96		std::vector<KeyValue> kvs;
97		kvs.reserve(kHashTableCapacity);
98
99		for (uint32_t i = 0; i < kHashTableCapacity; i++)
	93	unsigned int threadid = blockIdx.x * blockDim.x + threadIdx.x;
	94	if (threadid < kHashTableCapacity)
100	95	{
101		if (pHostHashTable[i].key != kEmpty)
	96	if (pHashTable[threadid].key != kEmpty)
102	97	{
103		kvs.push_back(pHostHashTable[i]);
	98	uint32_t size = atomicAdd(kvs_size, 1);
	99	kvs[size] = pHashTable[threadid];
104	100	}
105	101	}
	102	}
	103
	104	std::vector<KeyValue> iterate_hashtable(KeyValue* pHashTable)
	105	{
	106	uint32_t* device_num_kvs;
	107	cudaMalloc(&device_num_kvs, sizeof(uint32_t));
	108	cudaMemset(device_num_kvs, 0, sizeof(uint32_t));
	109
	110	KeyValue* device_kvs;
	111	cudaMalloc(&device_kvs, sizeof(KeyValue) * kNumKeyValues);
	112
	113	int mingridsize;
	114	int threadblocksize;
	115	cudaOccupancyMaxPotentialBlockSize(&mingridsize, &threadblocksize, gpu_iterate_hashtable, 0, 0);
	116
	117	int gridsize = (kHashTableCapacity + threadblocksize - 1) / threadblocksize;
	118	gpu_iterate_hashtable<<<gridsize, threadblocksize>>>(pHashTable, device_kvs, device_num_kvs);
106	119
107		cudaFree(pHostHashTable);
	120	uint32_t num_kvs;
	121	cudaMemcpy(&num_kvs, device_num_kvs, sizeof(uint32_t), cudaMemcpyDeviceToHost);
	122
	123	std::vector<KeyValue> kvs;
	124	kvs.resize(num_kvs);
	125
	126	cudaMemcpy(kvs.data(), device_kvs, sizeof(KeyValue) * num_kvs, cudaMemcpyDeviceToHost);
	127
	128	cudaFree(device_kvs);
	129	cudaFree(device_num_kvs);
108	130
109	131	return kvs;
110	132	}

src/main.cpp +26 -23

…	…	@@ -1,42 +1,42 @@
1		#include "windows.h"
2		#include "stdio.h"
	1	#include "algorithm"
	2	#include "random"
3	3	#include "stdint.h"
4		#include "unordered_set"
	4	#include "stdio.h"
5	5	#include "unordered_map"
	6	#include "unordered_set"
6	7	#include "vector"
7		#include "algorithm"
8		#include "random"
	8	#include "windows.h"
9	9	#include "linearprobing.h"
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()
14	14	{
15	15	std::random_device rd;
16		std::mt19937 gen(rd()); // mersenne_twister_engine
17		std::uniform_int_distribution<uint32_t> dis(0, kEmpty-1);
	16	std::mt19937 gen(rd()); // mersenne_twister_engine
	17	std::uniform_int_distribution<uint32_t> dis(0, kEmpty - 1);
18	18
19	19	std::vector<KeyValue> kvs;
20	20	kvs.reserve(kNumKeyValues);
21	21
22		for (int i = 0; i < kNumKeyValues; i++)
	22	for (int i = 0; i < kNumKeyValues; i++)
23	23	{
24	24	uint32_t rand0 = dis(gen);
25	25	uint32_t rand1 = dis(gen);
26		kvs.push_back(KeyValue{ rand0, rand1 });
	26	kvs.push_back(KeyValue{rand0, rand1});
27	27	}
28	28
29	29	return kvs;
30	30	}
31	31
32		LARGE_INTEGER start_timer()
	32	LARGE_INTEGER start_timer()
33	33	{
34	34	LARGE_INTEGER timer;
35	35	QueryPerformanceCounter(&timer);
36	36	return timer;
37	37	}
38	38
39		double get_elapsed_time(LARGE_INTEGER start)
	39	double get_elapsed_time(LARGE_INTEGER start)
40	40	{
41	41	LARGE_INTEGER end;
42	42	QueryPerformanceCounter(&end);
…	…	@@ -52,27 +52,30 @@ double get_elapsed_time(LARGE_INTEGER start)
52	52	return milliseconds;
53	53	}
54	54
55		void test_unordered_map(std::vector<KeyValue> all_kvs)
	55	void test_unordered_map(std::vector<KeyValue> all_kvs)
56	56	{
57	57	LARGE_INTEGER timer = start_timer();
58	58
59		std::unordered_map<uint32_t, uint32_t> kvs_map;
60		for (auto& kv : all_kvs)
	59	printf("Testing insertion into std::unordered_map...\n");
	60
61	61	{
62		kvs_map[kv.key] = kv.value;
	62	std::unordered_map<uint32_t, uint32_t> kvs_map;
	63	for (auto& kv : all_kvs) {
	64	kvs_map[kv.key] = kv.value;
	65	}
63	66	}
64	67
65	68	double milliseconds = get_elapsed_time(timer);
66	69	double seconds = milliseconds / 1000.0f;
67		printf("std::unordered_map time : %f ms (%f million keys/second)\n",
	70	printf("Total time for std::unordered_map: %f ms (%f million keys/second)\n",
68	71	milliseconds, kNumKeyValues / seconds / 1000000.0f);
69	72	}
70	73
71	74	void test_correctness(std::vector<KeyValue>, std::vector<KeyValue>);
72	75
73		int main()
	76	int main()
74	77	{
75		printf("Testing insertion of %d elements\n", kNumKeyValues);
	78	printf("Testing insertion of %d elements into GPU hash table...\n", kNumKeyValues);
76	79
77	80	// Initialize keyvalue pairs with random numbers
78	81	std::vector<KeyValue> all_kvs = generate_random_keyvalues();
…	…	@@ -86,9 +89,9 @@ int main()
86	89	const uint32_t num_batch_insertions = 16;
87	90	static_assert(kNumKeyValues % num_batch_insertions == 0, "needs even divisor");
88	91	uint32_t num_kvs_in_batch = kNumKeyValues / num_batch_insertions;
89		for (int i = 0; i < num_batch_insertions; i++)
	92	for (int i = 0; i < num_batch_insertions; i++)
90	93	{
91		insert_hashtable(pHashTable, all_kvs.data() + i*num_kvs_in_batch, num_kvs_in_batch);
	94	insert_hashtable(pHashTable, all_kvs.data() + i * num_kvs_in_batch, num_kvs_in_batch);
92	95	}
93	96
94	97	// Get all the key-values from the hash table
…	…	@@ -99,8 +102,8 @@ int main()
99	102	// Summarize results
100	103	double milliseconds = get_elapsed_time(timer);
101	104	double seconds = milliseconds / 1000.0f;
102		printf("Total time (including memory copies, readback, etc): %f ms (%f million keys/second)\n",
103		milliseconds, kNumKeyValues / seconds / 1000000.0f);
	105	printf("Total time (including memory copies, readback, etc): %f ms (%f million keys/second)\n", milliseconds,
	106	kNumKeyValues / seconds / 1000000.0f);
104	107
105	108	test_unordered_map(all_kvs);
106	109

opencog-SimpleGPUHashTable commits / a2ae55c

Changed files (2) +61 -36