vyre 0.4.0

GPU compute intermediate representation with a standard operation library
Documentation 
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
44
45
46
47
48
49
50

@group(0) @binding(0) var<storage, read> input_words: array<u64>;
@group(0) @binding(1) var<storage, read_write> output_words: array<u64>;

const B2B_IV: array<u64, 8> = array<u64, 8>(
  0x6a09e667f3bcc908u, 0xbb67ae8584caa73bu, 0x3c6ef372fe94f82bu, 0xa54ff53a5f1d36f1u,
  0x510e527fade682d1u, 0x9b05688c2b3e6c1fu, 0x1f83d9abfb41bd6bu, 0x5be0cd19137e2179u);
const B2B_SIGMA: array<u32, 160> = array<u32, 160>(
  0u,1u,2u,3u,4u,5u,6u,7u,8u,9u,10u,11u,12u,13u,14u,15u, 14u,10u,4u,8u,9u,15u,13u,6u,1u,12u,0u,2u,11u,7u,5u,3u,
  11u,8u,12u,0u,5u,2u,15u,13u,10u,14u,3u,6u,7u,1u,9u,4u, 7u,9u,3u,1u,13u,12u,11u,14u,2u,6u,5u,10u,4u,0u,15u,8u,
  9u,0u,5u,7u,2u,4u,10u,15u,14u,1u,11u,12u,6u,8u,3u,13u, 2u,12u,6u,10u,0u,11u,8u,3u,4u,13u,7u,5u,15u,14u,1u,9u,
  12u,5u,1u,15u,14u,13u,4u,10u,0u,7u,6u,3u,9u,2u,8u,11u, 13u,11u,7u,14u,12u,1u,3u,9u,5u,0u,15u,4u,8u,6u,2u,10u,
  6u,15u,14u,9u,11u,3u,0u,8u,12u,2u,13u,7u,1u,4u,10u,5u, 10u,2u,8u,4u,7u,6u,1u,5u,15u,11u,9u,14u,3u,12u,13u,0u);

fn b2b_rotr(x: u64, n: u32) -> u64 { return (x >> n) | (x << ((64u - n) & 63u)); }
fn b2b_g(v_in: array<u64, 16>, a: u32, b: u32, c: u32, d: u32, x: u64, y: u64) -> array<u64, 16> {
  var v = v_in;
  v[a] = v[a] + v[b] + x; v[d] = b2b_rotr(v[d] ^ v[a], 32u);
  v[c] = v[c] + v[d];     v[b] = b2b_rotr(v[b] ^ v[c], 24u);
  v[a] = v[a] + v[b] + y; v[d] = b2b_rotr(v[d] ^ v[a], 16u);
  v[c] = v[c] + v[d];     v[b] = b2b_rotr(v[b] ^ v[c], 63u);
  return v;
}

@compute @workgroup_size(1, 1, 1)
fn hash_blake2b(@builtin(global_invocation_id) id: vec3<u32>) {
  if (id.x != 0u) { return; }
  var m: array<u64, 16>;
  for (var i = 0u; i < 16u; i = i + 1u) { m[i] = 0u; }
  let words = min(arrayLength(&input_words), 16u);
  for (var i = 0u; i < words; i = i + 1u) { m[i] = input_words[i]; }
  var h: array<u64, 8>;
  for (var i = 0u; i < 8u; i = i + 1u) { h[i] = B2B_IV[i]; }
  h[0u] = h[0u] ^ 0x01010040u;
  var v: array<u64, 16>;
  for (var i = 0u; i < 8u; i = i + 1u) { v[i] = h[i]; v[i + 8u] = B2B_IV[i]; }
  v[12u] = v[12u] ^ (u64(words) * 8u);
  v[14u] = ~v[14u];
  for (var r = 0u; r < 12u; r = r + 1u) {
    let o = (r % 10u) * 16u;
    v = b2b_g(v, 0u,4u,8u,12u, m[B2B_SIGMA[o]], m[B2B_SIGMA[o + 1u]]);
    v = b2b_g(v, 1u,5u,9u,13u, m[B2B_SIGMA[o + 2u]], m[B2B_SIGMA[o + 3u]]);
    v = b2b_g(v, 2u,6u,10u,14u, m[B2B_SIGMA[o + 4u]], m[B2B_SIGMA[o + 5u]]);
    v = b2b_g(v, 3u,7u,11u,15u, m[B2B_SIGMA[o + 6u]], m[B2B_SIGMA[o + 7u]]);
    v = b2b_g(v, 0u,5u,10u,15u, m[B2B_SIGMA[o + 8u]], m[B2B_SIGMA[o + 9u]]);
    v = b2b_g(v, 1u,6u,11u,12u, m[B2B_SIGMA[o + 10u]], m[B2B_SIGMA[o + 11u]]);
    v = b2b_g(v, 2u,7u,8u,13u, m[B2B_SIGMA[o + 12u]], m[B2B_SIGMA[o + 13u]]);
    v = b2b_g(v, 3u,4u,9u,14u, m[B2B_SIGMA[o + 14u]], m[B2B_SIGMA[o + 15u]]);
  }
  for (var i = 0u; i < 8u; i = i + 1u) { output_words[i] = h[i] ^ v[i] ^ v[i + 8u]; }
}