RainBow Table

1
#include <cuda_runtime.h>
2

3
__device__ unsigned int fnv1a_hash(unsigned int input) {
4
    const unsigned int FNV_PRIME = 16777619;
5
    const unsigned int OFFSET_BASIS = 2166136261;
6

7
    unsigned int hash = OFFSET_BASIS;
8

9
    for (int byte_pos = 0; byte_pos < 4; byte_pos++) {
10
        unsigned char byte = (input >> (byte_pos * 8)) & 0xFFu;
11
        hash = (hash ^ byte) * FNV_PRIME;
12
    }
13

14
    return hash;
15
}
16

17
__global__ void fnv1a_hash_kernel(const int* input, unsigned int* output, int N, int R) {
18
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
19
    if (idx >= N) return;
20

21
    unsigned int val = static_cast<unsigned int>(input[idx]);
22

23
    for (int r = 0; r < R; r++) {
24
        val = fnv1a_hash(val);
25
    }
26

27
    output[idx] = val;
28
}
29

30
// input, output are device pointers (i.e. pointers to memory on the GPU)
31
extern "C" void solve(const int* input, unsigned int* output, int N, int R) {
32
    int threadsPerBlock = 256;
33
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
34

35
    fnv1a_hash_kernel<<<blocksPerGrid, threadsPerBlock>>>(input, output, N, R);
36
    cudaDeviceSynchronize();
37
}