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

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

fn sip_rotl(x: u64, n: u32) -> u64 { return (x << n) | (x >> ((64u - n) & 63u)); }
fn sip_key0() -> u64 {
  if (arrayLength(&key_words) > 0u) {
    return key_words[0u];
  }
  return 0u;
}
fn sip_key1(k0: u64) -> u64 {
  return k0 ^ 0x73616e74682d7679u;
}
fn sip_round(v_in: array<u64, 4>) -> array<u64, 4> {
  var v = v_in;
  v[0u] = v[0u] + v[1u]; v[1u] = sip_rotl(v[1u], 13u); v[1u] = v[1u] ^ v[0u]; v[0u] = sip_rotl(v[0u], 32u);
  v[2u] = v[2u] + v[3u]; v[3u] = sip_rotl(v[3u], 16u); v[3u] = v[3u] ^ v[2u];
  v[0u] = v[0u] + v[3u]; v[3u] = sip_rotl(v[3u], 21u); v[3u] = v[3u] ^ v[0u];
  v[2u] = v[2u] + v[1u]; v[1u] = sip_rotl(v[1u], 17u); v[1u] = v[1u] ^ v[2u]; v[2u] = sip_rotl(v[2u], 32u);
  return v;
}

@compute @workgroup_size(1, 1, 1)
fn hash_siphash13(@builtin(global_invocation_id) id: vec3<u32>) {
  if (id.x != 0u) { return; }
  let k0 = sip_key0();
  let k1 = sip_key1(k0);
  var v: array<u64, 4>;
  v[0u] = 0x736f6d6570736575u ^ k0;
  v[1u] = 0x646f72616e646f6du ^ k1;
  v[2u] = 0x6c7967656e657261u ^ k0;
  v[3u] = 0x7465646279746573u ^ k1;
  let len = arrayLength(&input_words);
  for (var i = 0u; i < len; i = i + 1u) {
    v[3u] = v[3u] ^ input_words[i];
    v = sip_round(v);
    v[0u] = v[0u] ^ input_words[i];
  }
  let last = (u64(len) * 8u) << 56u;
  v[3u] = v[3u] ^ last;
  v = sip_round(v);
  v[0u] = v[0u] ^ last;
  v[2u] = v[2u] ^ 0xffu;
  v = sip_round(v);
  v = sip_round(v);
  v = sip_round(v);
  output_words[0u] = v[0u] ^ v[1u] ^ v[2u] ^ v[3u];
}