summaryrefslogtreecommitdiff
path: root/keccak.cuh
blob: 68dbde1557fe026815f08906c631854d27fefc5c (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
#pragma once

/* Compute a keccak-like hash in 1024 threads via values supplied
   locally in each thread (never leaving register memory). The input
   stream given by each warp has to be synchronised: the same amount of
   data has to be hashed at the same time).

   Approach:
   1. Compute a keccak-hash of an input stream of the local warp until
      end of data is reached. 25 (of the 32 threads) store one 32-bit
      word each.
   2. Compute the Keccak-Hash of all those hashes by streaming them to
      one warp via shared memory.
*/

#include <inttypes.h>

namespace keccak {

const uint32_t digest_words = 256 / 32;
struct Result {
  uint32_t digest[digest_words];
};

union State {
  uint32_t
    a[5][5], 
    s[25];
};

struct SharedState {
  Result digests[1024 / 32]; // digest from 32 warps
};

union RegisterState {
  uint32_t a;
};

__device__ void init(RegisterState &s);
__device__ void add(RegisterState &s, uint32_t data);
__device__ void finish(RegisterState &s, SharedState &ss, Result &dst);

} // namespace keccak
contact: Jan Huwald // Impressum