Initial work. Added some tests for big nums in rust and OpenCL.
This commit is contained in:
Kirill Shakirov
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,29 @@
|
||||
// aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
|
||||
//
|
||||
// Copyright (C) 2025 Kirill Shakirov
|
||||
//
|
||||
// This program 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.
|
||||
//
|
||||
// This program 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 this program. If not, see <https://www.gnu.org/licenses/>.
|
||||
|
||||
|
||||
inline void aes128_InvertKey (uint *ks);
|
||||
inline void aes128_ExpandKey (uint *ks, const uint *ukey);
|
||||
inline void aes128_set_encrypt_key (uint *ks, const uint *ukey);
|
||||
inline void aes128_set_decrypt_key (uint *ks, const uint *ukey);
|
||||
inline void aes128_encrypt (const uint *ks, const uint *in, uint *out);
|
||||
inline void aes128_decrypt (const uint *ks, const uint *in, uint *out);
|
||||
|
||||
inline void xts_mul2 (uint *in, uint *out);
|
||||
inline void aes_xts256_gen_tweak (const uint *ks, const uint *sec_n, const uint block_n, uint *out);
|
||||
inline void aes_xts256_enc_block (const uint *ks, const uint *T, const uint *in, uint *out);
|
||||
inline void aes_xts256_dec_block (const uint *ks, const uint *T, const uint *in, uint *out);
|
||||
@@ -0,0 +1,174 @@
|
||||
// aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
|
||||
//
|
||||
// Copyright (C) 2025 Kirill Shakirov
|
||||
//
|
||||
// This program 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.
|
||||
//
|
||||
// This program 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 this program. If not, see <https://www.gnu.org/licenses/>.
|
||||
|
||||
|
||||
|
||||
typedef struct __attribute__((aligned(32))) {
|
||||
uint carry;
|
||||
uint carry4;
|
||||
} bigintRes;
|
||||
|
||||
typedef union {
|
||||
ulong l;
|
||||
uint i[2];
|
||||
} __attribute__((aligned(32))) ul_ui_union;
|
||||
|
||||
|
||||
|
||||
inline uint add_one_to_bigint4_(uint *_n)
|
||||
{
|
||||
ul_ui_union t;
|
||||
t.l = (ulong)_n[0] + 1ul;
|
||||
_n[0] = t.i[0];
|
||||
|
||||
t.l = (ulong)_n[1] + (ulong)t.i[1];
|
||||
_n[1] = t.i[0];
|
||||
t.l = (ulong)_n[2] + (ulong)t.i[1];
|
||||
_n[2] = t.i[0];
|
||||
t.l = (ulong)_n[3] + (ulong)t.i[1];
|
||||
_n[3] = t.i[0];
|
||||
|
||||
return t.i[1];
|
||||
}
|
||||
|
||||
inline uint add_uint_to_bigint4_ (uint* _n, const uint b)
|
||||
{
|
||||
ul_ui_union t;
|
||||
t.l = (ulong)_n[0] + (ulong)b;
|
||||
_n[0] = t.i[0];
|
||||
|
||||
t.l = (ulong)_n[1] + (ulong)t.i[1];
|
||||
_n[1] = t.i[0];
|
||||
t.l = (ulong)_n[2] + (ulong)t.i[1];
|
||||
_n[2] = t.i[0];
|
||||
t.l = (ulong)_n[3] + (ulong)t.i[1];
|
||||
_n[3] = t.i[0];
|
||||
|
||||
return t.i[1];
|
||||
}
|
||||
|
||||
inline bigintRes add_one_to_bigint8(const uint *n, uint *out)
|
||||
{
|
||||
bigintRes res;
|
||||
ul_ui_union t;
|
||||
t.l = (ulong)n[0] + 1ul;
|
||||
out[0] = t.i[0];
|
||||
|
||||
t.l = (ulong)n[1] + (ulong)t.i[1];
|
||||
out[1] = t.i[0];
|
||||
t.l = (ulong)n[2] + (ulong)t.i[1];
|
||||
out[2] = t.i[0];
|
||||
t.l = (ulong)n[3] + (ulong)t.i[1];
|
||||
out[3] = t.i[0];
|
||||
|
||||
res.carry4 = t.i[1];
|
||||
t.l = (ulong)n[4] + (ulong)t.i[1];
|
||||
out[4] = t.i[0];
|
||||
t.l = (ulong)n[5] + (ulong)t.i[1];
|
||||
out[5] = t.i[0];
|
||||
t.l = (ulong)n[6] + (ulong)t.i[1];
|
||||
out[6] = t.i[0];
|
||||
t.l = (ulong)n[7] + (ulong)t.i[1];
|
||||
out[7] = t.i[0];
|
||||
res.carry = t.i[1];
|
||||
return res;
|
||||
}
|
||||
|
||||
inline bigintRes add_one_to_bigint8_(uint *_n)
|
||||
{
|
||||
bigintRes res;
|
||||
ul_ui_union t;
|
||||
t.l = (ulong)_n[0] + 1ul;
|
||||
_n[0] = t.i[0];
|
||||
|
||||
t.l = (ulong)_n[1] + (ulong)t.i[1];
|
||||
_n[1] = t.i[0];
|
||||
t.l = (ulong)_n[2] + (ulong)t.i[1];
|
||||
_n[2] = t.i[0];
|
||||
t.l = (ulong)_n[3] + (ulong)t.i[1];
|
||||
_n[3] = t.i[0];
|
||||
|
||||
res.carry4 = t.i[1];
|
||||
t.l = (ulong)_n[4] + (ulong)t.i[1];
|
||||
_n[4] = t.i[0];
|
||||
t.l = (ulong)_n[5] + (ulong)t.i[1];
|
||||
_n[5] = t.i[0];
|
||||
t.l = (ulong)_n[6] + (ulong)t.i[1];
|
||||
_n[6] = t.i[0];
|
||||
t.l = (ulong)_n[7] + (ulong)t.i[1];
|
||||
_n[7] = t.i[0];
|
||||
|
||||
res.carry = t.i[1];
|
||||
return res;
|
||||
}
|
||||
|
||||
|
||||
inline bigintRes add_uint_to_bigint8 (const uint *n, const uint b, uint *out)
|
||||
{
|
||||
bigintRes res;
|
||||
ul_ui_union t;
|
||||
t.l = (ulong)n[0] + (ulong)b;
|
||||
out[0] = t.i[0];
|
||||
|
||||
t.l = (ulong)n[1] + (ulong)t.i[1];
|
||||
out[1] = t.i[0];
|
||||
t.l = (ulong)n[2] + (ulong)t.i[1];
|
||||
out[2] = t.i[0];
|
||||
t.l = (ulong)n[3] + (ulong)t.i[1];
|
||||
out[3] = t.i[0];
|
||||
|
||||
res.carry4 = t.i[1];
|
||||
t.l = (ulong)n[4] + (ulong)t.i[1];
|
||||
out[4] = t.i[0];
|
||||
t.l = (ulong)n[5] + (ulong)t.i[1];
|
||||
out[5] = t.i[0];
|
||||
t.l = (ulong)n[6] + (ulong)t.i[1];
|
||||
out[6] = t.i[0];
|
||||
t.l = (ulong)n[7] + (ulong)t.i[1];
|
||||
out[7] = t.i[0];
|
||||
|
||||
res.carry = t.i[1];
|
||||
return res;
|
||||
}
|
||||
|
||||
inline bigintRes add_uint_to_bigint8_ (uint* _n, const uint b)
|
||||
{
|
||||
bigintRes res;
|
||||
ul_ui_union t;
|
||||
t.l = (ulong)_n[0] + (ulong)b;
|
||||
_n[0] = t.i[0];
|
||||
|
||||
t.l = (ulong)_n[1] + (ulong)t.i[1];
|
||||
_n[1] = t.i[0];
|
||||
t.l = (ulong)_n[2] + (ulong)t.i[1];
|
||||
_n[2] = t.i[0];
|
||||
t.l = (ulong)_n[3] + (ulong)t.i[1];
|
||||
_n[3] = t.i[0];
|
||||
|
||||
res.carry4 = t.i[1];
|
||||
t.l = (ulong)_n[4] + (ulong)t.i[1];
|
||||
_n[4] = t.i[0];
|
||||
t.l = (ulong)_n[5] + (ulong)t.i[1];
|
||||
_n[5] = t.i[0];
|
||||
t.l = (ulong)_n[6] + (ulong)t.i[1];
|
||||
_n[6] = t.i[0];
|
||||
t.l = (ulong)_n[7] + (ulong)t.i[1];
|
||||
_n[7] = t.i[0];
|
||||
|
||||
res.carry = t.i[1];
|
||||
return res;
|
||||
}
|
||||
@@ -0,0 +1,101 @@
|
||||
// aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
|
||||
//
|
||||
// Copyright (C) 2025 Kirill Shakirov
|
||||
//
|
||||
// This program 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.
|
||||
//
|
||||
// This program 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 this program. If not, see <https://www.gnu.org/licenses/>.
|
||||
|
||||
|
||||
#include "aes256_xts_plain.cl"
|
||||
#include "num_utils.cl"
|
||||
|
||||
// batch_size uint - batch size
|
||||
// Ti - sector index ulong
|
||||
// Tj - encryption block number (16 bytes blocks)
|
||||
|
||||
// Tk - tweak key uint[4]
|
||||
// s_Dk - start of data key uint[4]
|
||||
|
||||
// t_e_d - target_enc_data uint[4]
|
||||
// u_d - unencrypted data to be encrypted uint[4]
|
||||
|
||||
// g_key_found uint[9] - 0 element - flag that sets to 1 if key found.
|
||||
// Other 8 elements is found key
|
||||
|
||||
__kernel void search_key_test(const uint batch_size, const ulong g_Ti, const uint g_Tj,
|
||||
__global const uint8* g_start_enc_key,
|
||||
__global const uint4* g_uenc_data,
|
||||
__global const uint4* g_target_data,
|
||||
__global uint* g_key_found)
|
||||
{
|
||||
const uint g_id = get_global_id(0);
|
||||
|
||||
uint enc_key[8];
|
||||
uint tweak[4];
|
||||
uint uenc_data[4];
|
||||
uint4 target_data = *g_target_data;
|
||||
uint4 enc_data = (uint4)(0);
|
||||
uint d_ks[44]; // data expanded key
|
||||
uint t_ks[44]; // tweak expanded key
|
||||
|
||||
uint sec_n[4] = {0};
|
||||
sec_n[0] = ((uint*)&g_Ti)[0];
|
||||
sec_n[1] = ((uint*)&g_Ti)[1];
|
||||
|
||||
uint Tj = g_Tj;
|
||||
|
||||
vstore4(*g_uenc_data, 0, uenc_data);
|
||||
vstore8(*g_start_enc_key, 0, enc_key);
|
||||
|
||||
|
||||
// Set initial start key for every work thread
|
||||
uint k_data_carry = add_uint_to_bigint4_ (enc_key, (g_id*batch_size));
|
||||
uint k_tweak_carry = add_one_to_bigint4_ (&enc_key[4]);
|
||||
if (k_tweak_carry != 0u) return; // if reached max key value exit thread
|
||||
|
||||
// Generate tweak
|
||||
aes128_set_encrypt_key (t_ks, &enc_key[4]);
|
||||
aes_xts256_gen_tweak (t_ks, sec_n, Tj, tweak);
|
||||
|
||||
for (uint batch_id = 0u; (batch_id < batch_size); batch_id++)
|
||||
{
|
||||
// Data encrypt key always changing because we increment from 0 index to 8
|
||||
aes128_set_encrypt_key (d_ks, enc_key);
|
||||
|
||||
// encrypt data
|
||||
aes_xts256_enc_block (d_ks, tweak, uenc_data, (uint*)&enc_data);
|
||||
|
||||
// check if we found the key!
|
||||
if (all(enc_data==target_data))
|
||||
{
|
||||
g_key_found[0] = 1;
|
||||
vstore8(vload8(0, enc_key), 0, &g_key_found[1]);
|
||||
return;
|
||||
}
|
||||
|
||||
// Increment data key part by 1.
|
||||
k_data_carry = add_one_to_bigint4_ (enc_key);
|
||||
|
||||
// Tweak changes only once in 2^128 times
|
||||
if (k_data_carry != 0u) {
|
||||
// Increment tweak part
|
||||
k_tweak_carry = add_one_to_bigint4_ (&enc_key[4]);
|
||||
if (k_tweak_carry != 0u) return; // if reached max key value exit thread
|
||||
// Gen new tweak
|
||||
aes128_set_encrypt_key (t_ks, &enc_key[4]);
|
||||
aes_xts256_gen_tweak (t_ks, sec_n, Tj, tweak);
|
||||
}
|
||||
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,98 @@
|
||||
// aes-xts-pur64 is OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
|
||||
//
|
||||
// Copyright (C) 2025 Kirill Shakirov
|
||||
//
|
||||
// This program 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.
|
||||
//
|
||||
// This program 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 this program. If not, see <https://www.gnu.org/licenses/>.
|
||||
|
||||
|
||||
#include "aes256_xts_plain.cl"
|
||||
|
||||
__kernel void encrypt_data(__global const ulong* g_Ti, __global const uint* g_Tj,
|
||||
__global const uint8* g_key,
|
||||
__global const uint4* g_u_data,
|
||||
__global uint* g_enc_data)
|
||||
{
|
||||
const size_t g_id = get_global_id(0);
|
||||
uint d_ks[44];
|
||||
uint t_ks[44];
|
||||
uint tweak[4];
|
||||
uint enc_key[8];
|
||||
uint u_data[4];
|
||||
uint enc_data[4] = { 0 };
|
||||
|
||||
uint sec_n[4];
|
||||
ulong Ti = g_Ti[g_id];
|
||||
sec_n[0] = ((uint*)&Ti)[0];
|
||||
sec_n[1] = ((uint*)&Ti)[1];
|
||||
sec_n[2] = 0;
|
||||
sec_n[3] = 0;
|
||||
|
||||
uint Tj = g_Tj[g_id];
|
||||
|
||||
vstore8(*g_key, 0, enc_key);
|
||||
vstore4(g_u_data[g_id], 0, u_data);
|
||||
|
||||
// printf("Ti: %lu\\n", Ti);
|
||||
// printf("Tj: %u\\n", Tj);
|
||||
// printf("enc_key: %v8u\\n", *(uint8*)enc_key);
|
||||
// printf("uenc_data: %v4u\\n", *(uint4*)uenc_data);
|
||||
|
||||
//calculate tweak value
|
||||
aes128_set_encrypt_key (t_ks, &enc_key[4]);
|
||||
aes_xts256_gen_tweak (t_ks, sec_n, Tj, tweak);
|
||||
|
||||
// encrypt data
|
||||
aes128_set_encrypt_key (d_ks, enc_key);
|
||||
aes_xts256_enc_block (d_ks, tweak, u_data, enc_data);
|
||||
// printf("enc_data: %v4u\\n", *(uint4*)enc_data);
|
||||
vstore4(*(uint4*)enc_data, g_id, g_enc_data);
|
||||
}
|
||||
|
||||
|
||||
__kernel void decrypt_data(__global const ulong* g_Ti, __global const uint* g_Tj,
|
||||
__global const uint8* g_key,
|
||||
__global const uint4* g_enc_data,
|
||||
__global uint* g_u_data)
|
||||
{
|
||||
const size_t g_id = get_global_id(0);
|
||||
uint d_ks[44];
|
||||
uint t_ks[44];
|
||||
uint tweak[4];
|
||||
uint enc_key[8];
|
||||
uint enc_data[4];
|
||||
uint u_data[4] = { 0 };
|
||||
|
||||
uint sec_n[4];
|
||||
ulong Ti = g_Ti[g_id];
|
||||
sec_n[0] = ((uint*)&Ti)[0];
|
||||
sec_n[1] = ((uint*)&Ti)[1];
|
||||
sec_n[2] = 0;
|
||||
sec_n[3] = 0;
|
||||
|
||||
uint Tj = g_Tj[g_id];
|
||||
|
||||
vstore8(*g_key, 0, enc_key);
|
||||
vstore4(g_enc_data[g_id], 0, enc_data);
|
||||
|
||||
|
||||
//calculate tweak value
|
||||
aes128_set_encrypt_key (t_ks, &enc_key[4]);
|
||||
aes_xts256_gen_tweak (t_ks, sec_n, Tj, tweak);
|
||||
|
||||
// decrypt data
|
||||
aes128_set_decrypt_key (d_ks, enc_key);
|
||||
aes_xts256_dec_block (d_ks, tweak, enc_data, u_data);
|
||||
// printf("enc_data: %v4u\\n", *(uint4*)enc_data);
|
||||
vstore4(*(uint4*)u_data, g_id, g_u_data);
|
||||
}
|
||||
@@ -0,0 +1,27 @@
|
||||
#include "num_utils.cl"
|
||||
|
||||
|
||||
__kernel void test_add(__global const uint* g_num_to_add,
|
||||
__global const uint8* g_t0,
|
||||
__global uint* g_t1,
|
||||
__global uint* g_t2)
|
||||
{
|
||||
const size_t g_id = get_global_id(0);
|
||||
uint t0[8] = {0};
|
||||
uint t1[8] = {0};
|
||||
uint t2[8] = {0};
|
||||
uint num_to_add = g_num_to_add[g_id];
|
||||
|
||||
vstore8(g_t0[g_id], 0, t0);
|
||||
add_one_to_bigint8 (t0, t1);
|
||||
add_uint_to_bigint8 (t0, num_to_add, t2);
|
||||
|
||||
// save results
|
||||
vstore8(*(uint8*)t1, 0, &g_t1[g_id*8]);
|
||||
vstore8(*(uint8*)t2, 0, &g_t2[g_id*8]);
|
||||
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
Reference in New Issue
Block a user