REDACTED-rig/src/backend/opencl/cl/rx/blake2b.cl
2024-08-07 00:18:51 +02:00

348 lines
12 KiB
Common Lisp

/*
Copyright (c) 2019 SChernykh
This file is part of RandomX OpenCL.
RandomX OpenCL is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or
(at your option) any later version.
RandomX OpenCL is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with RandomX OpenCL. If not, see <http://www.gnu.org/licenses/>.
*/
__constant static const uchar blake2b_sigma[12 * 16] = {
0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15,
14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3,
11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4,
7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8,
9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13,
2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9,
12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11,
13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10,
6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5,
10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13, 0,
0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15,
14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3,
};
enum Blake2b_IV
{
iv0 = 0x6a09e667f3bcc908ul,
iv1 = 0xbb67ae8584caa73bul,
iv2 = 0x3c6ef372fe94f82bul,
iv3 = 0xa54ff53a5f1d36f1ul,
iv4 = 0x510e527fade682d1ul,
iv5 = 0x9b05688c2b3e6c1ful,
iv6 = 0x1f83d9abfb41bd6bul,
iv7 = 0x5be0cd19137e2179ul,
};
ulong rotr64(ulong a, ulong shift) { return rotate(a, 64 - shift); }
#define G(r, i, a, b, c, d) \
do { \
a = a + b + m[blake2b_sigma[r * 16 + 2 * i + 0]]; \
d = rotr64(d ^ a, 32); \
c = c + d; \
b = rotr64(b ^ c, 24); \
a = a + b + m[blake2b_sigma[r * 16 + 2 * i + 1]]; \
d = rotr64(d ^ a, 16); \
c = c + d; \
b = rotr64(b ^ c, 63); \
} while (0)
#define ROUND(r) \
do { \
G(r, 0, v[0], v[4], v[8], v[12]); \
G(r, 1, v[1], v[5], v[9], v[13]); \
G(r, 2, v[2], v[6], v[10], v[14]); \
G(r, 3, v[3], v[7], v[11], v[15]); \
G(r, 4, v[0], v[5], v[10], v[15]); \
G(r, 5, v[1], v[6], v[11], v[12]); \
G(r, 6, v[2], v[7], v[8], v[13]); \
G(r, 7, v[3], v[4], v[9], v[14]); \
} while (0)
#define BLAKE2B_ROUNDS() ROUND(0);ROUND(1);ROUND(2);ROUND(3);ROUND(4);ROUND(5);ROUND(6);ROUND(7);ROUND(8);ROUND(9);ROUND(10);ROUND(11);
void blake2b_512_process_single_block(ulong *h, const ulong* m, uint blockTemplateSize)
{
ulong v[16] =
{
iv0 ^ 0x01010040, iv1, iv2, iv3, iv4 , iv5, iv6, iv7,
iv0 , iv1, iv2, iv3, iv4 ^ blockTemplateSize, iv5, ~iv6, iv7,
};
BLAKE2B_ROUNDS();
h[0] = v[0] ^ v[ 8] ^ iv0 ^ 0x01010040;
h[1] = v[1] ^ v[ 9] ^ iv1;
h[2] = v[2] ^ v[10] ^ iv2;
h[3] = v[3] ^ v[11] ^ iv3;
h[4] = v[4] ^ v[12] ^ iv4;
h[5] = v[5] ^ v[13] ^ iv5;
h[6] = v[6] ^ v[14] ^ iv6;
h[7] = v[7] ^ v[15] ^ iv7;
}
__attribute__((reqd_work_group_size(64, 1, 1)))
__kernel void blake2b_initial_hash(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
{
const uint global_index = get_global_id(0);
__global const ulong* p = (__global const ulong*) blockTemplate;
ulong m[16] = {
(blockTemplateSize > 0) ? p[ 0] : 0,
(blockTemplateSize > 8) ? p[ 1] : 0,
(blockTemplateSize > 16) ? p[ 2] : 0,
(blockTemplateSize > 24) ? p[ 3] : 0,
(blockTemplateSize > 32) ? p[ 4] : 0,
(blockTemplateSize > 40) ? p[ 5] : 0,
(blockTemplateSize > 48) ? p[ 6] : 0,
(blockTemplateSize > 56) ? p[ 7] : 0,
(blockTemplateSize > 64) ? p[ 8] : 0,
(blockTemplateSize > 72) ? p[ 9] : 0,
(blockTemplateSize > 80) ? p[10] : 0,
(blockTemplateSize > 88) ? p[11] : 0,
(blockTemplateSize > 96) ? p[12] : 0,
(blockTemplateSize > 104) ? p[13] : 0,
(blockTemplateSize > 112) ? p[14] : 0,
(blockTemplateSize > 120) ? p[15] : 0,
};
if (blockTemplateSize % sizeof(ulong))
m[blockTemplateSize / sizeof(ulong)] &= (ulong)(-1) >> (64 - (blockTemplateSize % sizeof(ulong)) * 8);
const ulong nonce = start_nonce + global_index;
m[4] = (m[4] & ((ulong)(-1) >> 8)) | (nonce << 56);
m[5] = (m[5] & ((ulong)(-1) << 24)) | (nonce >> 8);
ulong hash[8];
blake2b_512_process_single_block(hash, m, blockTemplateSize);
__global ulong* t = ((__global ulong*) out) + global_index * 8;
t[0] = hash[0];
t[1] = hash[1];
t[2] = hash[2];
t[3] = hash[3];
t[4] = hash[4];
t[5] = hash[5];
t[6] = hash[6];
t[7] = hash[7];
}
void blake2b_512_process_double_block_variable(ulong *out, ulong* m, __global const ulong* in, uint in_len, uint out_len)
{
ulong v[16] =
{
iv0 ^ (0x01010000u | out_len), iv1, iv2, iv3, iv4 , iv5, iv6, iv7,
iv0 , iv1, iv2, iv3, iv4 ^ 128, iv5, iv6, iv7,
};
BLAKE2B_ROUNDS();
ulong h[8];
v[0] = h[0] = v[0] ^ v[8] ^ iv0 ^ (0x01010000u | out_len);
v[1] = h[1] = v[1] ^ v[9] ^ iv1;
v[2] = h[2] = v[2] ^ v[10] ^ iv2;
v[3] = h[3] = v[3] ^ v[11] ^ iv3;
v[4] = h[4] = v[4] ^ v[12] ^ iv4;
v[5] = h[5] = v[5] ^ v[13] ^ iv5;
v[6] = h[6] = v[6] ^ v[14] ^ iv6;
v[7] = h[7] = v[7] ^ v[15] ^ iv7;
v[8] = iv0;
v[9] = iv1;
v[10] = iv2;
v[11] = iv3;
v[12] = iv4 ^ in_len;
v[13] = iv5;
v[14] = ~iv6;
v[15] = iv7;
m[ 0] = (in_len > 128) ? in[16] : 0;
m[ 1] = (in_len > 136) ? in[17] : 0;
m[ 2] = (in_len > 144) ? in[18] : 0;
m[ 3] = (in_len > 152) ? in[19] : 0;
m[ 4] = (in_len > 160) ? in[20] : 0;
m[ 5] = (in_len > 168) ? in[21] : 0;
m[ 6] = (in_len > 176) ? in[22] : 0;
m[ 7] = (in_len > 184) ? in[23] : 0;
m[ 8] = (in_len > 192) ? in[24] : 0;
m[ 9] = (in_len > 200) ? in[25] : 0;
m[10] = (in_len > 208) ? in[26] : 0;
m[11] = (in_len > 216) ? in[27] : 0;
m[12] = (in_len > 224) ? in[28] : 0;
m[13] = (in_len > 232) ? in[29] : 0;
m[14] = (in_len > 240) ? in[30] : 0;
m[15] = (in_len > 248) ? in[31] : 0;
if (in_len % sizeof(ulong))
m[(in_len - 128) / sizeof(ulong)] &= (ulong)(-1) >> (64 - (in_len % sizeof(ulong)) * 8);
BLAKE2B_ROUNDS();
if (out_len > 0) out[0] = h[0] ^ v[0] ^ v[8];
if (out_len > 8) out[1] = h[1] ^ v[1] ^ v[9];
if (out_len > 16) out[2] = h[2] ^ v[2] ^ v[10];
if (out_len > 24) out[3] = h[3] ^ v[3] ^ v[11];
if (out_len > 32) out[4] = h[4] ^ v[4] ^ v[12];
if (out_len > 40) out[5] = h[5] ^ v[5] ^ v[13];
if (out_len > 48) out[6] = h[6] ^ v[6] ^ v[14];
if (out_len > 56) out[7] = h[7] ^ v[7] ^ v[15];
}
__attribute__((reqd_work_group_size(64, 1, 1)))
__kernel void blake2b_initial_hash_double(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce)
{
const uint global_index = get_global_id(0);
__global const ulong* p = (__global const ulong*) blockTemplate;
ulong m[16] = { p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], p[8], p[9], p[10], p[11], p[12], p[13], p[14], p[15] };
const ulong nonce = start_nonce + global_index;
m[4] = (m[4] & ((ulong)(-1) >> 8)) | (nonce << 56);
m[5] = (m[5] & ((ulong)(-1) << 24)) | (nonce >> 8);
ulong hash[8];
blake2b_512_process_double_block_variable(hash, m, p, blockTemplateSize, 64);
__global ulong* t = ((__global ulong*) out) + global_index * 8;
t[0] = hash[0];
t[1] = hash[1];
t[2] = hash[2];
t[3] = hash[3];
t[4] = hash[4];
t[5] = hash[5];
t[6] = hash[6];
t[7] = hash[7];
}
void blake2b_512_process_big_block(ulong *out, __global const ulong* in, uint in_len, uint out_len, uint nonce, uint nonce_offset)
{
ulong h[8] = { iv0 ^ (0x01010000u | out_len), iv1, iv2, iv3, iv4, iv5, iv6, iv7 };
for (uint t = 128; t < in_len; t += 128, in += 16) {
ulong m[16] = { in[0], in[1], in[2], in[3], in[4], in[5], in[6], in[7], in[8], in[9], in[10], in[11], in[12], in[13], in[14], in[15] };
const uint k0 = (nonce_offset + 0) - (t - 128);
const uint k1 = (nonce_offset + 1) - (t - 128);
const uint k2 = (nonce_offset + 2) - (t - 128);
const uint k3 = (nonce_offset + 3) - (t - 128);
if (k0 < 128) m[k0 / 8] |= (ulong)((nonce >> 0) & 255) << ((k0 % 8) * 8);
if (k1 < 128) m[k1 / 8] |= (ulong)((nonce >> 8) & 255) << ((k1 % 8) * 8);
if (k2 < 128) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
if (k3 < 128) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
ulong v[16] = { h[0], h[1], h[2], h[3], h[4], h[5], h[6], h[7], iv0, iv1, iv2, iv3, iv4 ^ t, iv5, iv6, iv7 };
BLAKE2B_ROUNDS();
h[0] ^= v[0] ^ v[ 8];
h[1] ^= v[1] ^ v[ 9];
h[2] ^= v[2] ^ v[10];
h[3] ^= v[3] ^ v[11];
h[4] ^= v[4] ^ v[12];
h[5] ^= v[5] ^ v[13];
h[6] ^= v[6] ^ v[14];
h[7] ^= v[7] ^ v[15];
}
uint k = in_len & 127;
if (k == 0) k = 128;
ulong m[16] = {
(k > 0) ? in[ 0] : 0,
(k > 8) ? in[ 1] : 0,
(k > 16) ? in[ 2] : 0,
(k > 24) ? in[ 3] : 0,
(k > 32) ? in[ 4] : 0,
(k > 40) ? in[ 5] : 0,
(k > 48) ? in[ 6] : 0,
(k > 56) ? in[ 7] : 0,
(k > 64) ? in[ 8] : 0,
(k > 72) ? in[ 9] : 0,
(k > 80) ? in[10] : 0,
(k > 88) ? in[11] : 0,
(k > 96) ? in[12] : 0,
(k > 104) ? in[13] : 0,
(k > 112) ? in[14] : 0,
(k > 120) ? in[15] : 0
};
const uint t = in_len - k;
const uint k0 = nonce_offset + 0 - t;
const uint k1 = nonce_offset + 1 - t;
const uint k2 = nonce_offset + 2 - t;
const uint k3 = nonce_offset + 3 - t;
if (k0 < k) m[k0 / 8] |= (ulong)((nonce >> 0) & 255) << ((k0 % 8) * 8);
if (k1 < k) m[k1 / 8] |= (ulong)((nonce >> 8) & 255) << ((k1 % 8) * 8);
if (k2 < k) m[k2 / 8] |= (ulong)((nonce >> 16) & 255) << ((k2 % 8) * 8);
if (k3 < k) m[k3 / 8] |= (ulong)((nonce >> 24) & 255) << ((k3 % 8) * 8);
if (k % 8) {
m[k / 8] &= (ulong)(-1) >> (64 - (k % 8) * 8);
}
ulong v[16] = { h[0], h[1], h[2], h[3], h[4], h[5], h[6], h[7], iv0, iv1, iv2, iv3, iv4 ^ in_len, iv5, ~iv6, iv7 };
BLAKE2B_ROUNDS();
if (out_len > 0) out[0] = h[0] ^ v[0] ^ v[8];
if (out_len > 8) out[1] = h[1] ^ v[1] ^ v[9];
if (out_len > 16) out[2] = h[2] ^ v[2] ^ v[10];
if (out_len > 24) out[3] = h[3] ^ v[3] ^ v[11];
if (out_len > 32) out[4] = h[4] ^ v[4] ^ v[12];
if (out_len > 40) out[5] = h[5] ^ v[5] ^ v[13];
if (out_len > 48) out[6] = h[6] ^ v[6] ^ v[14];
if (out_len > 56) out[7] = h[7] ^ v[7] ^ v[15];
}
__attribute__((reqd_work_group_size(64, 1, 1)))
__kernel void blake2b_initial_hash_big(__global void *out, __global const void* blockTemplate, uint blockTemplateSize, uint start_nonce, uint nonce_offset)
{
const uint global_index = get_global_id(0);
__global const ulong* p = (__global const ulong*) blockTemplate;
ulong hash[8];
blake2b_512_process_big_block(hash, p, blockTemplateSize, 64, start_nonce + global_index, nonce_offset);
__global ulong* t = ((__global ulong*) out) + global_index * 8;
t[0] = hash[0];
t[1] = hash[1];
t[2] = hash[2];
t[3] = hash[3];
t[4] = hash[4];
t[5] = hash[5];
t[6] = hash[6];
t[7] = hash[7];
}
#define in_len 256
#define out_len 32
#define blake2b_512_process_double_block_name blake2b_512_process_double_block_32
#define blake2b_hash_registers_name blake2b_hash_registers_32
#include "blake2b_double_block.cl"
#undef blake2b_hash_registers_name
#undef blake2b_512_process_double_block_name
#undef out_len
#define out_len 64
#define blake2b_512_process_double_block_name blake2b_512_process_double_block_64
#define blake2b_hash_registers_name blake2b_hash_registers_64
#include "blake2b_double_block.cl"
#undef blake2b_hash_registers_name
#undef blake2b_512_process_double_block_name
#undef out_len