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

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

const HMD5_S: array<u32, 64> = array<u32, 64>(
  7u,12u,17u,22u,7u,12u,17u,22u,7u,12u,17u,22u,7u,12u,17u,22u,5u,9u,14u,20u,5u,9u,14u,20u,5u,9u,14u,20u,5u,9u,14u,20u,
  4u,11u,16u,23u,4u,11u,16u,23u,4u,11u,16u,23u,4u,11u,16u,23u,6u,10u,15u,21u,6u,10u,15u,21u,6u,10u,15u,21u,6u,10u,15u,21u);
const HMD5_K: array<u32, 64> = array<u32, 64>(
  0xd76aa478u,0xe8c7b756u,0x242070dbu,0xc1bdceeeu,0xf57c0fafu,0x4787c62au,0xa8304613u,0xfd469501u,
  0x698098d8u,0x8b44f7afu,0xffff5bb1u,0x895cd7beu,0x6b901122u,0xfd987193u,0xa679438eu,0x49b40821u,
  0xf61e2562u,0xc040b340u,0x265e5a51u,0xe9b6c7aau,0xd62f105du,0x02441453u,0xd8a1e681u,0xe7d3fbc8u,
  0x21e1cde6u,0xc33707d6u,0xf4d50d87u,0x455a14edu,0xa9e3e905u,0xfcefa3f8u,0x676f02d9u,0x8d2a4c8au,
  0xfffa3942u,0x8771f681u,0x6d9d6122u,0xfde5380cu,0xa4beea44u,0x4bdecfa9u,0xf6bb4b60u,0xbebfbc70u,
  0x289b7ec6u,0xeaa127fau,0xd4ef3085u,0x04881d05u,0xd9d4d039u,0xe6db99e5u,0x1fa27cf8u,0xc4ac5665u,
  0xf4292244u,0x432aff97u,0xab9423a7u,0xfc93a039u,0x655b59c3u,0x8f0ccc92u,0xffeff47du,0x85845dd1u,
  0x6fa87e4fu,0xfe2ce6e0u,0xa3014314u,0x4e0811a1u,0xf7537e82u,0xbd3af235u,0x2ad7d2bbu,0xeb86d391u);

fn hmd5_rotl(x: u32, n: u32) -> u32 { return (x << n) | (x >> ((32u - n) & 31u)); }
fn hmd5_f(i: u32, b: u32, c: u32, d: u32) -> u32 {
  if (i < 16u) { return (b & c) | ((~b) & d); }
  if (i < 32u) { return (d & b) | ((~d) & c); }
  if (i < 48u) { return b ^ c ^ d; }
  return c ^ (b | (~d));
}
fn hmd5_g(i: u32) -> u32 {
  if (i < 16u) { return i; }
  if (i < 32u) { return (5u * i + 1u) & 15u; }
  if (i < 48u) { return (3u * i + 5u) & 15u; }
  return (7u * i) & 15u;
}
fn hmd5_compress(block: array<u32, 16>, seed: array<u32, 4>) -> array<u32, 4> {
  var a = seed[0u]; var b = seed[1u]; var c = seed[2u]; var d = seed[3u];
  for (var i = 0u; i < 64u; i = i + 1u) {
    let next = b + hmd5_rotl(a + hmd5_f(i, b, c, d) + HMD5_K[i] + block[hmd5_g(i)], HMD5_S[i]);
    a = d; d = c; c = b; b = next;
  }
  return array<u32, 4>(seed[0u] + a, seed[1u] + b, seed[2u] + c, seed[3u] + d);
}
fn hmd5_be(x: u32) -> u32 { return ((x & 0xffu) << 24u) | ((x & 0xff00u) << 8u) | ((x >> 8u) & 0xff00u) | ((x >> 24u) & 0xffu); }

@compute @workgroup_size(1, 1, 1)
fn hash_hmac_md5(@builtin(global_invocation_id) id: vec3<u32>) {
  if (id.x != 0u) { return; }
  var ipad: array<u32, 16>;
  var opad: array<u32, 16>;
  for (var i = 0u; i < 16u; i = i + 1u) { ipad[i] = 0x36363636u; opad[i] = 0x5c5c5c5cu; }
  let kwords = min(arrayLength(&key_words), 16u);
  for (var i = 0u; i < kwords; i = i + 1u) { ipad[i] = ipad[i] ^ key_words[i]; opad[i] = opad[i] ^ key_words[i]; }
  var seed = array<u32, 4>(0x67452301u, 0xefcdab89u, 0x98badcfeu, 0x10325476u);
  var inner = hmd5_compress(ipad, seed);
  var msg_block: array<u32, 16>;
  for (var i = 0u; i < 16u; i = i + 1u) { msg_block[i] = 0u; }
  let mwords = min(arrayLength(&msg_words), 14u);
  for (var i = 0u; i < mwords; i = i + 1u) { msg_block[i] = msg_words[i]; }
  msg_block[mwords] = msg_block[mwords] | 0x80u;
  msg_block[14u] = (64u + mwords * 4u) * 8u;
  inner = hmd5_compress(msg_block, inner);
  var outer_block = opad;
  var digest_block: array<u32, 16>;
  for (var i = 0u; i < 16u; i = i + 1u) { digest_block[i] = 0u; }
  digest_block[0u] = inner[0u]; digest_block[1u] = inner[1u]; digest_block[2u] = inner[2u]; digest_block[3u] = inner[3u];
  digest_block[4u] = 0x80u; digest_block[14u] = 80u * 8u;
  let outer_seed = hmd5_compress(outer_block, seed);
  let out = hmd5_compress(digest_block, outer_seed);
  output_words[0u] = hmd5_be(out[0u]); output_words[1u] = hmd5_be(out[1u]); output_words[2u] = hmd5_be(out[2u]); output_words[3u] = hmd5_be(out[3u]);
}