| … | … | @@ -14,6 +14,20 @@ __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) { |
| 20 | // Allocate memory |
| 21 | KeyValue* hashtable; |
| 22 | cudaMalloc(&hashtable, sizeof(KeyValue) * kHashTableCapacity); |
| 23 | |
| 24 | // Initialize hash table to empty |
| 25 | static_assert(kEmpty == 0xffffffff, "memset expected kEmpty=0xffffffff"); |
| 26 | cudaMemset(hashtable, 0xff, sizeof(KeyValue) * kHashTableCapacity); |
| 27 | |
| 28 | return hashtable; |
| 29 | } |
| 30 | |
| 17 | 31 | __global__ void gpu_hashtable_insert(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs) |
| 18 | 32 | { |
| 19 | 33 | unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x; |
| … | … | @@ -25,8 +39,8 @@ __global__ void gpu_hashtable_insert(KeyValue* hashtable, const KeyValue* kvs, u |
| 25 | 39 | |
| 26 | 40 | while (true) |
| 27 | 41 | { |
| 28 | | uint32_t prev = atomicCAS(&hashtable[slot].key, 0, key); |
| 29 | | if (prev == 0 || prev == key) |
| 42 | uint32_t prev = atomicCAS(&hashtable[slot].key, kEmpty, key); |
| 43 | if (prev == kEmpty || prev == key) |
| 30 | 44 | { |
| 31 | 45 | hashtable[slot].value = value; |
| 32 | 46 | break; |
| … | … | @@ -37,18 +51,6 @@ __global__ void gpu_hashtable_insert(KeyValue* hashtable, const KeyValue* kvs, u |
| 37 | 51 | } |
| 38 | 52 | } |
| 39 | 53 | |
| 40 | | KeyValue* 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 | | |
| 52 | 54 | void insert_hashtable(KeyValue* pHashTable, const KeyValue* kvs, uint32_t num_kvs) |
| 53 | 55 | { |
| 54 | 56 | // Copy the keyvalues to the GPU |
| … | … | @@ -68,7 +70,7 @@ void insert_hashtable(KeyValue* pHashTable, const KeyValue* kvs, uint32_t num_kv |
| 68 | 70 | |
| 69 | 71 | cudaEventRecord(start); |
| 70 | 72 | |
| 71 | | // Hash table insertion |
| 73 | // Insert all the keys into the hash table |
| 72 | 74 | int gridsize = ((uint32_t)num_kvs + threadblocksize - 1) / threadblocksize; |
| 73 | 75 | gpu_hashtable_insert<<<gridsize, threadblocksize>>>(pHashTable, device_kvs, (uint32_t)num_kvs); |
| 74 | 76 | |
| … | … | @@ -96,7 +98,7 @@ std::vector<KeyValue> iterate_hashtable(KeyValue* pHashTable) |
| 96 | 98 | |
| 97 | 99 | for (uint32_t i = 0; i < kHashTableCapacity; i++) |
| 98 | 100 | { |
| 99 | | if (pHostHashTable[i].key != 0) |
| 101 | if (pHostHashTable[i].key != kEmpty) |
| 100 | 102 | { |
| 101 | 103 | kvs.push_back(pHostHashTable[i]); |
| 102 | 104 | } |