▶ 使用CPU和GPU分别实现散列表
● CPU方法
#include <stdio.h>
#include <time.h>
#include "cuda_runtime.h"
#include "D:\Code\CUDA\book\common\book.h" #define SIZE (100*1024*1024)
#define ELEMENTS (SIZE / sizeof(unsigned int))
#define HASH_ENTRIES (1024) struct Entry
{
unsigned int key;
void *value;
Entry *next;
}; struct Table
{
size_t count;
Entry **entries;
Entry *pool;
Entry *firstFree;
}; size_t hash(unsigned int key, size_t count)
{
return key % count;
} void initialize_table(Table &table, int entries, int elements)
{
table.count = entries;
table.entries = (Entry**)calloc(entries, sizeof(Entry*));
table.pool = (Entry*)malloc(elements * sizeof(Entry));
table.firstFree = table.pool;
} void free_table(Table &table)
{
free(table.entries);
free(table.pool);
} void add_to_table(Table &table, unsigned int key, void *value)
{
size_t hashValue = hash(key, table.count);
Entry *location = table.firstFree++;
location->key = key;
location->value = value;
location->next = table.entries[hashValue];// 插到该分支的头部而不是尾部
table.entries[hashValue] = location;
} void verify_table(const Table &table)
{
int count = ;
for (size_t i = ; i<table.count; i++)
{
Entry *current = table.entries[i];
while (current != NULL)
{
++count;
if (hash(current->key, table.count) != i)
printf("\n\t%d hashed to %ld, but was located at %ld\n", current->key, hash(current->key, table.count), i);
current = current->next;
}
}
if (count != ELEMENTS)
printf("\n\t%d elements found in hash table. Should be %ld\n",
count, ELEMENTS);
else
printf("\n\tAll %d elements found in hash table.\n", count);
} int main(void)
{
unsigned int *buffer =(unsigned int*)big_random_block(SIZE);
Table table;
clock_t start, stop; initialize_table(table, HASH_ENTRIES, ELEMENTS); start = clock();
for (int i = ; i<ELEMENTS; i++)
add_to_table(table, buffer[i], (void*)NULL); stop = clock();
printf("\n\tBuilding the table: %3.1f ms\n", (float)(stop - start) / (float)CLOCKS_PER_SEC * 1000.0f); verify_table(table);
free_table(table);
free(buffer);
getchar();
return ;
}
● GPU方法（用到了前面的原子锁）
#include <stdio.h>
#include <time.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda.h"
#include "D:\Code\CUDA\book\common\book.h" #define SIZE (100*1024*1024)
#define ELEMENTS (SIZE / sizeof(unsigned int))
#define HASH_ENTRIES (1024) struct Lock
{
int *mutex;
Lock(void)
{
int state = ;
cudaMalloc((void **)&mutex, sizeof(int));
cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
}
~Lock(void)
{
cudaFree(mutex);
}
__device__ void lock(void)
{
while (atomicCAS(mutex, , ) != );
}
__device__ void unlock(void)
{
atomicExch(mutex, );
}
}; struct Entry
{
unsigned int key;
void *value;
Entry *next;
}; struct Table
{
size_t count;
Entry **entries;
Entry *pool;
Entry *firstFree;
}; __device__ __host__ size_t hash(unsigned int key, size_t count)
{
return key % count;
} void initialize_table(Table &table, int entries, int elements)
{
table.count = entries;
cudaMalloc((void**)&table.entries, entries * sizeof(Entry*));
cudaMemset(table.entries, , entries * sizeof(Entry*));
cudaMalloc((void**)&table.pool, elements * sizeof(Entry));
} void free_table(Table &table)
{
cudaFree(table.entries);
cudaFree(table.pool);
} __global__ void add_to_table(unsigned int *keys, void **values, Table table, Lock *lock)
// 锁数组用于锁定散列表中的每一个桶
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
while (tid < ELEMENTS)
{
unsigned int key = keys[tid];
size_t hashValue = hash(key, table.count);
for (int i = ; i<; i++)// 利用循环来分散线程束，使同一线程束中的32个线程在循环的不同次数时进行写入
{
if ((tid % ) == i)
{
Entry *location = &(table.pool[tid]);
location->key = key;
location->value = values[tid];
lock[hashValue].lock();
location->next = table.entries[hashValue];
table.entries[hashValue] = location;
lock[hashValue].unlock();
}
}
tid += stride;
}
} void copy_table_to_host(const Table &table, Table &hostTable)
{
hostTable.count = table.count;
hostTable.entries = (Entry**)calloc(table.count, sizeof(Entry*));
hostTable.pool = (Entry*)malloc(ELEMENTS * sizeof(Entry)); cudaMemcpy(hostTable.entries, table.entries, table.count * sizeof(Entry*), cudaMemcpyDeviceToHost);
cudaMemcpy(hostTable.pool, table.pool, ELEMENTS * sizeof(Entry), cudaMemcpyDeviceToHost); for (int i = ; i < table.count; i++)
{
if (hostTable.entries[i] != NULL)
hostTable.entries[i] = (Entry*)((size_t)hostTable.entries[i] - (size_t)table.pool + (size_t)hostTable.pool);
// 从从显存到内存的地址线性偏移 x - adressGPU + addressCPU
}
for (int i = ; i < ELEMENTS; i++)
{
if (hostTable.pool[i].next != NULL)
hostTable.pool[i].next = (Entry*)((size_t)hostTable.pool[i].next - (size_t)table.pool + (size_t)hostTable.pool);
// 同样是做偏移，但是要找到下一个元素的地址
}
} void verify_table(const Table &dev_table)
{
Table table;
copy_table_to_host(dev_table, table); int count = ;
for (size_t i = ; i < table.count; i++)
{
Entry *current = table.entries[i];
while (current != NULL)
{
++count;
if (hash(current->key, table.count) != i)
printf("%d hashed to %ld, but was located at %ld\n", current->key, hash(current->key, table.count), i);
current = current->next;
}
}
if (count != ELEMENTS)
printf("%d elements found in hash table. Should be %ld\n", count, ELEMENTS);
else
printf("All %d elements found in hash table.\n", count);
} int main(void)
{
unsigned int *buffer = (unsigned int*)big_random_block(SIZE); unsigned int *dev_keys;
void **dev_values;
cudaMalloc((void**)&dev_keys, SIZE);
cudaMalloc((void**)&dev_values, SIZE);
cudaMemcpy(dev_keys, buffer, SIZE, cudaMemcpyHostToDevice); Table table;
initialize_table(table, HASH_ENTRIES, ELEMENTS); Lock lock[HASH_ENTRIES];// 准备锁列表
Lock *dev_lock;
cudaMalloc((void**)&dev_lock, HASH_ENTRIES * sizeof(Lock));
cudaMemcpy(dev_lock, lock, HASH_ENTRIES * sizeof(Lock), cudaMemcpyHostToDevice); cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, ); add_to_table << <, >> >(dev_keys, dev_values, table, dev_lock); cudaEventRecord(stop, );
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Time to hash: %3.1f ms\n", elapsedTime); verify_table(table);
free_table(table); cudaEventDestroy(start);
cudaEventDestroy(stop);
free_table(table);
cudaFree(dev_lock);
cudaFree(dev_keys);
cudaFree(dev_values);
free(buffer);
getchar();
return ;
}