OverviewFilesCommitsIssuesDocsReleasesPeers

opencog-SimpleGPUHashTable commits / 333ce8c

Initial commit
333ce8c80caae2bf0888b5e797fa9628bee351f1
by David Farrell <nosferalatu@gmail.com> · 2020-03-01T00:33:41.000Z

Changed files (5) +323 -0

.gitignore +1 -0
……@@ -0,0 +1 @@
1build/
CMakeLists.txt +19 -0
……@@ -0,0 +1,19 @@
1cmake_minimum_required(VERSION 3.8 FATAL_ERROR)
2
3project(SimpleConcurrentGPUHashTable LANGUAGES CXX CUDA)
4
5# put predefined cmake projects in their own solution folder
6set_property(GLOBAL PROPERTY USE_FOLDERS ON)
7
8string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_35,code=sm_35")
9
10set(SOURCE_FILES src/test.cpp src/linearprobing.h src/linearprobing.cu)
11
12include_directories(${CMAKE_SOURCE_DIR}/src)
13
14add_executable(test ${SOURCE_FILES})
15
16# visual studio project should mimic directory structure
17# this isn't working for me; I think because
18# https://developercommunity.visualstudio.com/content/problem/777578/source-grouptree-no-longer-works.html
19source_group(TREE ${CMAKE_CURRENT_SOURCE_DIR} FILES ${SOURCE_FILES})
src/linearprobing.cu +112 -0
……@@ -0,0 +1,112 @@
1#include "stdio.h"
2#include "stdint.h"
3#include "vector"
4#include "linearprobing.h"
5
6// 32 bit Murmur3 hash
7__device__ uint32_t hash(uint32_t k)
8{
9 k ^= k >> 16;
10 k *= 0x85ebca6b;
11 k ^= k >> 13;
12 k *= 0xc2b2ae35;
13 k ^= k >> 16;
14 return k & (kHashTableCapacity-1);
15}
16
17__global__ void gpu_hashtable_insert(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
18{
19 unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
20 if (threadid < numkvs)
21 {
22 uint32_t key = kvs[threadid].key;
23 uint32_t value = kvs[threadid].value;
24 uint32_t slot = hash(key);
25
26 while (true)
27 {
28 uint32_t prev = atomicCAS(&hashtable[slot].key, 0, key);
29 if (prev == 0 || prev == key)
30 {
31 hashtable[slot].value = value;
32 break;
33 }
34
35 slot = (slot + 1) & (kHashTableCapacity-1);
36 }
37 }
38}
39
40KeyValue* create_hashtable(uint32_t capacity)
41{
42 // Allocate memory
43 KeyValue* hashtable;
44 cudaMalloc(&hashtable, sizeof(KeyValue)*kHashTableCapacity);
45
46 // Initialize hash table to empty
47 cudaMemset(hashtable, 0, sizeof(KeyValue)*kHashTableCapacity);
48
49 return hashtable;
50}
51
52void insert_hashtable(KeyValue* pHashTable, const KeyValue* kvs, uint32_t num_kvs)
53{
54 // Copy the keyvalues to the GPU
55 KeyValue* device_kvs;
56 cudaMalloc(&device_kvs, sizeof(KeyValue) * num_kvs);
57 cudaMemcpy(device_kvs, kvs, sizeof(KeyValue) * num_kvs, cudaMemcpyHostToDevice);
58
59 // Have CUDA calculate the thread block size
60 int mingridsize;
61 int threadblocksize;
62 cudaOccupancyMaxPotentialBlockSize(&mingridsize, &threadblocksize, gpu_hashtable_insert, 0, 0);
63
64 // Create events for GPU timing
65 cudaEvent_t start, stop;
66 cudaEventCreate(&start);
67 cudaEventCreate(&stop);
68
69 cudaEventRecord(start);
70
71 // Hash table insertion
72 int gridsize = ((uint32_t)num_kvs + threadblocksize - 1) / threadblocksize;
73 gpu_hashtable_insert<<<gridsize, threadblocksize>>>(pHashTable, device_kvs, (uint32_t)num_kvs);
74
75 cudaEventRecord(stop);
76
77 cudaEventSynchronize(stop);
78
79 float milliseconds = 0;
80 cudaEventElapsedTime(&milliseconds, start, stop);
81 float seconds = milliseconds / 1000.0f;
82 printf(" GPU inserted %d items in %f ms (%f million keys/second)\n", num_kvs, milliseconds, num_kvs / (double)seconds / 1000000.0f);
83
84 cudaFree(device_kvs);
85}
86
87std::vector<KeyValue> iterate_hashtable(KeyValue* pHashTable)
88{
89 KeyValue* pHostHashTable;
90 cudaHostAlloc(&pHostHashTable, sizeof(KeyValue) * kHashTableCapacity, 0);
91 cudaMemcpy(pHostHashTable, pHashTable, sizeof(KeyValue) * kHashTableCapacity, cudaMemcpyDeviceToHost);
92
93 std::vector<KeyValue> kvs;
94 kvs.reserve(kHashTableCapacity);
95
96 for (uint32_t i = 0; i < kHashTableCapacity; i++)
97 {
98 if (pHostHashTable[i].key != 0)
99 {
100 kvs.push_back(pHostHashTable[i]);
101 }
102 }
103
104 cudaFree(pHostHashTable);
105
106 return kvs;
107}
108
109void destroy_hashtable(KeyValue* pHashTable)
110{
111 cudaFree(pHashTable);
112}
src/linearprobing.h +21 -0
……@@ -0,0 +1,21 @@
1#pragma once
2
3#define ASSERT(X) if (!(X)) __debugbreak();
4
5struct KeyValue
6{
7 uint32_t key;
8 uint32_t value;
9};
10
11const uint32_t kHashTableCapacity = 128 * 1024 * 1024;
12
13const uint32_t kNumKeyValues = kHashTableCapacity / 2;
14
15KeyValue* create_hashtable(uint32_t initialCapacity);
16
17void insert_hashtable(KeyValue* hashtable, const KeyValue* kvs, uint32_t num_kvs);
18
19std::vector<KeyValue> iterate_hashtable(KeyValue* hashtable);
20
21void destroy_hashtable(KeyValue* hashtable);
src/test.cpp +170 -0
……@@ -0,0 +1,170 @@
1#include "windows.h"
2#include "stdio.h"
3#include "stdint.h"
4#include "unordered_set"
5#include "unordered_map"
6#include "vector"
7#include "algorithm"
8#include "random"
9#include "linearprobing.h"
10
11std::vector<KeyValue> generate_random_keyvalues()
12{
13 std::random_device rd;
14 std::mt19937 gen(rd()); // mersenne_twister_engine
15 std::uniform_int_distribution<uint32_t> dis(1, 0xffffffff); // 0 is reserved for empty keys
16
17 std::vector<KeyValue> kvs;
18 kvs.reserve(kNumKeyValues);
19
20 for (int i = 0; i < kNumKeyValues; i++)
21 {
22 uint32_t rand0 = dis(gen);
23 uint32_t rand1 = dis(gen);
24 kvs.push_back(KeyValue{ rand0, rand1 });
25 }
26
27 return kvs;
28}
29
30LARGE_INTEGER start_timer()
31{
32 LARGE_INTEGER timer;
33 QueryPerformanceCounter(&timer);
34 return timer;
35}
36
37double get_elapsed_time(LARGE_INTEGER start)
38{
39 LARGE_INTEGER end;
40 QueryPerformanceCounter(&end);
41
42 LARGE_INTEGER frequency;
43 QueryPerformanceFrequency(&frequency);
44
45 LARGE_INTEGER elapsed;
46 elapsed.QuadPart = end.QuadPart - start.QuadPart;
47 elapsed.QuadPart *= 1000000;
48 double milliseconds = elapsed.QuadPart / (double)frequency.QuadPart / 1000;
49
50 return milliseconds;
51}
52
53void test_unordered_map(std::vector<KeyValue> all_kvs)
54{
55 LARGE_INTEGER timer = start_timer();
56
57 std::unordered_map<uint32_t, uint32_t> kvs_map;
58 for (auto& kv : all_kvs)
59 {
60 kvs_map[kv.key] = kv.value;
61 }
62
63 double milliseconds = get_elapsed_time(timer);
64 double seconds = milliseconds / 1000.0f;
65 printf("std::unordered_map time : %f ms (%f million keys/second)\n", milliseconds, kNumKeyValues / seconds / 1000000.0f);
66}
67
68int main()
69{
70 printf("Testing insertion of %d elements\n", kNumKeyValues);
71
72 // Initialize keyvalue pairs with random numbers
73 std::vector<KeyValue> all_kvs = generate_random_keyvalues();
74
75 // Begin test
76 LARGE_INTEGER timer = start_timer();
77
78 KeyValue* pHashTable = create_hashtable(kHashTableCapacity);
79
80 // Insert all the elements; each batch is processed concurrently on the GPU
81 int num_batch_insertions = 10;
82 int num_kvs_in_batch = kNumKeyValues / num_batch_insertions;
83 for (int i = 0; i < num_batch_insertions; i++)
84 {
85 insert_hashtable(pHashTable, all_kvs.data() + i*num_kvs_in_batch, num_kvs_in_batch);
86 }
87
88 // Get all the key-values from the hash table
89 std::vector<KeyValue> kvs = iterate_hashtable(pHashTable);
90
91 destroy_hashtable(pHashTable);
92
93 // Summarize results
94 double milliseconds = get_elapsed_time(timer);
95 double seconds = milliseconds / 1000.0f;
96 printf("Total time (including memory copies, readback, etc.): %f ms (%f million keys/second)\n", milliseconds, kNumKeyValues / seconds / 1000000.0f);
97
98 test_unordered_map(all_kvs);
99
100 // Test the results
101 if (true)
102 {
103 printf("Testing that there are no duplicate keys...\n");
104 std::unordered_set<uint32_t> unique_keys;
105 for (uint32_t i = 0; i < kvs.size(); i++)
106 {
107 if (i % 10000000 == 0)
108 printf(" Verifying %d/%d\n", i, (uint32_t)kvs.size());
109
110 KeyValue* node = &kvs[i];
111 if (unique_keys.find(node->key) != unique_keys.end())
112 {
113 printf("Duplicate key found in GPU hash table at slot %d\n", i);
114 exit(-1);
115 }
116 unique_keys.insert(node->key);
117 }
118
119 printf("Building unordered_map from insertion list...\n");
120 std::unordered_map<uint32_t, std::vector<uint32_t>> all_kvs_map;
121 for (int i = 0; i < all_kvs.size(); i++)
122 {
123 if (i % 10000000 == 0)
124 printf(" Building %d/%d\n", i, (uint32_t)all_kvs.size());
125
126 auto iter = all_kvs_map.find(all_kvs[i].key);
127 if (iter == all_kvs_map.end())
128 {
129 all_kvs_map[all_kvs[i].key] = std::vector<uint32_t>({ all_kvs[i].value });
130 }
131 else
132 {
133 iter->second.push_back(all_kvs[i].value);
134 }
135 }
136
137 if (unique_keys.size() != all_kvs_map.size())
138 {
139 printf("# of unique keys in hashtable is incorrect\n");
140 exit(-1);
141 }
142
143 printf("Testing that each key/value in hashtable is in the insertion list...\n");
144 for (uint32_t i = 0; i < kvs.size(); i++)
145 {
146 if (i % 10000000 == 0)
147 printf(" Verifying %d/%d\n", i, (uint32_t)kvs.size());
148
149 auto iter = all_kvs_map.find(kvs[i].key);
150 if (iter == all_kvs_map.end())
151 {
152 printf("Hashtable key not found in insertion list\n");
153 exit(-1);
154 }
155
156 std::vector<uint32_t>& values = iter->second;
157 if (std::find(values.begin(), values.end(), kvs[i].value) == values.end())
158 {
159 printf("Hashtable value not found in insertion list\n");
160 exit(-1);
161 }
162 }
163
164 printf("Deleting std::unordered_map and std::unique_set...\n");
165 }
166
167 printf("Success\n");
168
169 return 0;
170}