Hashtable.cu

//%%cuda -c "-I /does/not/exist -arch=sm_75"
#include <cuda_runtime.h>
#include <cstdint>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

#define TABLE_SIZE 1048576
typedef uint64_t Fingerprint;

__global__ void CudaHashInsert(Fingerprint* d_fingerprints, Fingerprint* d_hashTable, bool *d_isDuplicate, int numberOfItems)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	if(idx >= numberOfItems) return;

	Fingerprint fp = d_fingerprints[idx];
	uint32_t slot = fp % TABLE_SIZE;

	//Linear probing
	for(int i = 0; i < 8; i++)
	{
		uint32_t probeSlot = (slot + i) % TABLE_SIZE;
		if(d_hashTable[probeSlot] == 0)
		{
			d_hashTable[probeSlot] = fp;
			d_isDuplicate[idx] = false;
			return;
		}
		else if(d_hashTable[probeSlot] == fp)
		{
			d_isDuplicate[idx] = true;
			return;
		}
	}

	d_isDuplicate[idx] = true;
}

void CudaDeduplicate(Fingerprint* h_fingerprints, bool* h_results, int n)
{
	Fingerprint *d_fingerprints, *d_table;
	bool *d_results;

	cudaMalloc(&d_fingerprints, n * sizeof(Fingerprint));
	cudaMalloc(&d_results, n * sizeof(bool));
	cudaMalloc(&d_table, TABLE_SIZE * sizeof(Fingerprint));

	cudaMemcpy(d_fingerprints, h_fingerprints, n * sizeof(Fingerprint), cudaMemcpyHostToDevice);
	cudaMemset(d_table, 0, TABLE_SIZE * sizeof(Fingerprint));

	int threads = 256;
	int blocks = (n + threads - 1) / threads;
	CudaHashInsert<<<blocks, threads>>>(d_fingerprints, d_table, d_results, n);

	cudaMemcpy(h_results, d_results, n * sizeof(bool), cudaMemcpyDeviceToHost);

	cudaFree(d_fingerprints);
	cudaFree(d_results);
	cudaFree(d_table);
}

int main()
{
	int n = 1000000;
	Fingerprint* fps = (Fingerprint*)malloc(n * sizeof(Fingerprint));
	bool* results = (bool*)malloc(n * sizeof(bool));

	srand(time(NULL));

	//Create test data 
	for(int i = 0; i < n; i++) 
	{
		if (i % 5 == 0 && i > 0)
		{
			//Duplicate
			fps[i] = fps[i-1];
		}
		else
		{
			// New unique value
			fps[i] = ((uint64_t)rand() << 32) | rand();
		}
	}

	clock_t start = clock();
	CudaDeduplicate(fps, results, n);
	clock_t end = clock();
    
	// Count results
	int unique = 0, duplicates = 0;
	for(int i = 0; i < n; i++)
	{
		if(results[i]) duplicates++;
		else unique++;
	}
    
	printf("Unique: %d, Duplicates: %d\n", unique, duplicates);
	printf("Time: %.3f seconds\n", (double)(end - start) / CLOCKS_PER_SEC);

	printf("\nFirst 10 results:\n");
	for(int i = 0; i < 10; i++)
	{
		printf("Item %d: FP=%lu, Duplicate=%s\n", i, fps[i], results[i] ? "YES" : "NO");
	}

	free(fps);
	free(results);
	return 0;
}