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
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74

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

const R1: array<u32, 80> = array<u32, 80>(
  0u,1u,2u,3u,4u,5u,6u,7u,8u,9u,10u,11u,12u,13u,14u,15u,
  7u,4u,13u,1u,10u,6u,15u,3u,12u,0u,9u,5u,2u,14u,11u,8u,
  3u,10u,14u,4u,9u,15u,8u,1u,2u,7u,0u,6u,13u,11u,5u,12u,
  1u,9u,11u,10u,0u,8u,12u,4u,13u,3u,7u,15u,14u,5u,6u,2u,
  4u,0u,5u,9u,7u,12u,2u,10u,14u,1u,3u,8u,11u,6u,15u,13u);
const R2: array<u32, 80> = array<u32, 80>(
  5u,14u,7u,0u,9u,2u,11u,4u,13u,6u,15u,8u,1u,10u,3u,12u,
  6u,11u,3u,7u,0u,13u,5u,10u,14u,15u,8u,12u,4u,9u,1u,2u,
  15u,5u,1u,3u,7u,14u,6u,9u,11u,8u,12u,2u,10u,0u,4u,13u,
  8u,6u,4u,1u,3u,11u,15u,0u,5u,12u,2u,13u,9u,7u,10u,14u,
  12u,15u,10u,4u,1u,5u,8u,7u,6u,2u,13u,14u,0u,3u,9u,11u);
const S1: array<u32, 80> = array<u32, 80>(
  11u,14u,15u,12u,5u,8u,7u,9u,11u,13u,14u,15u,6u,7u,9u,8u,
  7u,6u,8u,13u,11u,9u,7u,15u,7u,12u,15u,9u,11u,7u,13u,12u,
  11u,13u,6u,7u,14u,9u,13u,15u,14u,8u,13u,6u,5u,12u,7u,5u,
  11u,12u,14u,15u,14u,15u,9u,8u,9u,14u,5u,6u,8u,6u,5u,12u,
  9u,15u,5u,11u,6u,8u,13u,12u,5u,12u,13u,14u,11u,8u,5u,6u);
const S2: array<u32, 80> = array<u32, 80>(
  8u,9u,9u,11u,13u,15u,15u,5u,7u,7u,8u,11u,14u,14u,12u,6u,
  9u,13u,15u,7u,12u,8u,9u,11u,7u,7u,12u,7u,6u,15u,13u,11u,
  9u,7u,15u,11u,8u,6u,6u,14u,12u,13u,5u,14u,13u,13u,7u,5u,
  15u,5u,8u,11u,14u,14u,6u,14u,6u,9u,12u,9u,12u,5u,15u,8u,
  8u,5u,12u,9u,12u,5u,14u,6u,8u,13u,6u,5u,15u,13u,11u,11u);

fn ripemd_rot(x: u32, n: u32) -> u32 { return (x << n) | (x >> ((32u - n) & 31u)); }
fn ripemd_f(j: u32, x: u32, y: u32, z: u32) -> u32 {
  if (j < 16u) { return x ^ y ^ z; }
  if (j < 32u) { return (x & y) | ((~x) & z); }
  if (j < 48u) { return (x | (~y)) ^ z; }
  if (j < 64u) { return (x & z) | (y & (~z)); }
  return x ^ (y | (~z));
}
fn ripemd_k1(j: u32) -> u32 {
  if (j < 16u) { return 0u; }
  if (j < 32u) { return 0x5a827999u; }
  if (j < 48u) { return 0x6ed9eba1u; }
  if (j < 64u) { return 0x8f1bbcdcu; }
  return 0xa953fd4eu;
}
fn ripemd_k2(j: u32) -> u32 {
  if (j < 16u) { return 0x50a28be6u; }
  if (j < 32u) { return 0x5c4dd124u; }
  if (j < 48u) { return 0x6d703ef3u; }
  if (j < 64u) { return 0x7a6d76e9u; }
  return 0u;
}

@compute @workgroup_size(1, 1, 1)
fn hash_ripemd160(@builtin(global_invocation_id) id: vec3<u32>) {
  if (id.x != 0u) { return; }
  var x: array<u32, 16>;
  for (var i = 0u; i < 16u; i = i + 1u) { x[i] = 0u; }
  let words = min(arrayLength(&input_words), 14u);
  for (var i = 0u; i < words; i = i + 1u) { x[i] = input_words[i]; }
  let bytes = words * 4u;
  if (words < 16u) { x[words] = x[words] | 0x80u; }
  x[14u] = bytes * 8u;
  var h0 = 0x67452301u; var h1 = 0xefcdab89u; var h2 = 0x98badcfeu; var h3 = 0x10325476u; var h4 = 0xc3d2e1f0u;
  var a1 = h0; var b1 = h1; var c1 = h2; var d1 = h3; var e1 = h4;
  var a2 = h0; var b2 = h1; var c2 = h2; var d2 = h3; var e2 = h4;
  for (var j = 0u; j < 80u; j = j + 1u) {
    let t1 = ripemd_rot(a1 + ripemd_f(j, b1, c1, d1) + x[R1[j]] + ripemd_k1(j), S1[j]) + e1;
    a1 = e1; e1 = d1; d1 = ripemd_rot(c1, 10u); c1 = b1; b1 = t1;
    let t2 = ripemd_rot(a2 + ripemd_f(79u - j, b2, c2, d2) + x[R2[j]] + ripemd_k2(j), S2[j]) + e2;
    a2 = e2; e2 = d2; d2 = ripemd_rot(c2, 10u); c2 = b2; b2 = t2;
  }
  let t = h1 + c1 + d2;
  h1 = h2 + d1 + e2; h2 = h3 + e1 + a2; h3 = h4 + a1 + b2; h4 = h0 + b1 + c2; h0 = t;
  output_words[0u] = h0; output_words[1u] = h1; output_words[2u] = h2; output_words[3u] = h3; output_words[4u] = h4;
}