Initial commit, hopefully everything is fine :3
This commit is contained in:
Kira
@@ -1,2 +1,24 @@
|
||||
# aes-xts-pur64
|
||||
# aes-xts-pur64
|
||||
OpenCL code for aes-xts256-plain64 encryption compatible with LUKS
|
||||
|
||||
|
||||
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 code is based on work of Hashcat project and its contributors:
|
||||
https://github.com/hashcat/hashcat
|
||||
|
||||
It modified to be more simple and OpenCLish.
|
||||
It must be compatible with most OpenCL2.0+ implementations.
|
||||
|
||||
Code contains methods for encryption and decryption with aes-xts256-plain64 algorithm.
|
||||
And also small python test suite.
|
||||
Code tested to be compatible with LUKS encryption.
|
||||
|
||||
It supposed to be executed inside individual work items.
|
||||
See test_aes_xts256_plain.cl for example usage.
|
||||
|
||||
|
||||
Copyright (C) 2025 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,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);
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,81 @@
|
||||
# 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/>.
|
||||
|
||||
import json
|
||||
import subprocess
|
||||
import nyanger.simple.static as nya_stat
|
||||
import random
|
||||
|
||||
# Init logger
|
||||
log = nya_stat.get_logger("nyan")
|
||||
|
||||
LUKS_VOL_FILE_NAME = "./data/test_vol.img"
|
||||
UNENCRYPT_DATA_FILE_NAME= "./data/unencrypt.img"
|
||||
OUT_FILE_NAME = "./data/test_vectors.json"
|
||||
KEY_FILE_NAME = "./data/master.key"
|
||||
VECTORS_NUM=256
|
||||
|
||||
def read_metadata(file_name: str) -> dict:
|
||||
#cryptsetup luksDump --dump-json-metadata /dev/loop0
|
||||
luks_cmd: list[str] = ["cryptsetup", "luksDump", "--dump-json-metadata", file_name]
|
||||
|
||||
result = subprocess.run(luks_cmd, capture_output=True, encoding="UTF-8")
|
||||
if result.returncode == 0 and result.stdout is not None:
|
||||
metadata = json.loads(result.stdout)
|
||||
return metadata
|
||||
else:
|
||||
raise Exception(f"Error executing 'cryptsetup' binary! {result.stderr}")
|
||||
|
||||
|
||||
def main():
|
||||
log.info("Start :3")
|
||||
metad = read_metadata(LUKS_VOL_FILE_NAME)
|
||||
log.info(f"metadata:\n{metad}")
|
||||
|
||||
segments_offset_bytes = int(metad["segments"]["0"]["offset"])
|
||||
|
||||
with open(LUKS_VOL_FILE_NAME, 'rb') as luks_file:
|
||||
luks_file.seek(segments_offset_bytes)
|
||||
enc_data = luks_file.read(512*VECTORS_NUM)
|
||||
|
||||
with open(UNENCRYPT_DATA_FILE_NAME, 'rb') as data_file:
|
||||
uenc_data = data_file.read(512*VECTORS_NUM)
|
||||
|
||||
with open(KEY_FILE_NAME, "rb") as key_file:
|
||||
key_data = key_file.read(32)
|
||||
|
||||
data_list = []
|
||||
for sec_n in range(VECTORS_NUM):
|
||||
block_num = random.randint(0, 31)
|
||||
_offset = ((sec_n*512)+(block_num*16))
|
||||
entry = {
|
||||
"sector_number": sec_n,
|
||||
"block_number": block_num,
|
||||
"unencrypted_data": uenc_data[_offset:_offset+16].hex(),
|
||||
"encrypted_data": enc_data[_offset:_offset+16].hex()
|
||||
}
|
||||
data_list.append(entry)
|
||||
|
||||
with open(OUT_FILE_NAME, 'wt') as out_file:
|
||||
json.dump({"encryption_key": key_data.hex(sep=' ', bytes_per_sep=16), "vectors": data_list},
|
||||
out_file, indent=" ")
|
||||
|
||||
log.info(" --- Finish --- ")
|
||||
|
||||
|
||||
if __name__ == '__main__':
|
||||
main()
|
||||
@@ -0,0 +1,9 @@
|
||||
[project]
|
||||
name = "aes-xts-pur64"
|
||||
version = "0.1.0"
|
||||
description = "Tests fro aes-xts-pur64"
|
||||
requires-python = ">=3.13"
|
||||
dependencies = [
|
||||
"nyanger==0.9.2",
|
||||
"pyopencl==2025.2.7"
|
||||
]
|
||||
@@ -0,0 +1,227 @@
|
||||
# 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/>.
|
||||
|
||||
|
||||
import json
|
||||
import time
|
||||
import argparse
|
||||
from nyanger.simple.log_writers.file_writer import FileWriter
|
||||
from nyanger.simple.log_writers.console_writer import ConsoleWriter
|
||||
import nyanger.simple.static as nya_stat
|
||||
import pyopencl as cl
|
||||
|
||||
# Init logger
|
||||
log: nya_stat.Nyanger
|
||||
|
||||
def load_program(include_dirs: list[str], file_name: str, ctx: cl.Context) -> cl.Program:
|
||||
with open(file_name, 'rt') as ft:
|
||||
prog_src = ft.read()
|
||||
opt_list = [f"-I {d}" for d in include_dirs]
|
||||
prg = cl.Program(ctx, prog_src).build(options=opt_list)
|
||||
return prg
|
||||
|
||||
|
||||
def test(cl_src_file: str, cl_incl_dir: str, vectors_file: str):
|
||||
log.info("Starting tests...")
|
||||
total_errors = 0
|
||||
start_time = time.time()
|
||||
|
||||
log.info(f"Loading test vectors from {vectors_file}")
|
||||
|
||||
with open(vectors_file, "rt") as vec_file:
|
||||
test_json = json.load(vec_file)
|
||||
|
||||
b_key = bytes.fromhex(test_json["encryption_key"])
|
||||
test_vectors = test_json["vectors"]
|
||||
test_vectors_num = len(test_vectors)
|
||||
|
||||
log.info(f"Data key: {b_key[:16].hex()}")
|
||||
log.info(f"Tweak key: {b_key[16:].hex()}")
|
||||
|
||||
b_u_data = bytearray(test_vectors_num * 16)
|
||||
b_enc_data = bytearray(test_vectors_num * 16)
|
||||
b_sec_nums = bytearray(test_vectors_num * 8)
|
||||
b_blk_nums = bytearray(test_vectors_num * 4)
|
||||
|
||||
for i, _vect in enumerate(test_vectors):
|
||||
of16 = i * 16
|
||||
of8 = i * 8
|
||||
of4 = i * 4
|
||||
b_u_data[of16:of16 + 16] = bytes.fromhex(_vect["unencrypted_data"])
|
||||
b_enc_data[of16:of16 + 16] = bytes.fromhex(_vect["encrypted_data"])
|
||||
b_sec_nums[of8:of8 + 8] = _vect["sector_number"].to_bytes(length=8, byteorder="little")
|
||||
b_blk_nums[of4:of4 + 4] = _vect["block_number"].to_bytes(length=4, byteorder="little")
|
||||
|
||||
log.info(f"{test_vectors_num} test vectors loaded.")
|
||||
|
||||
log.info("Getting OCL devices...")
|
||||
# get OCL devices
|
||||
cl_platforms = cl.get_platforms()
|
||||
log.info(f"{len(cl_platforms)} platforms found.")
|
||||
cl_devices = []
|
||||
for platform in cl_platforms:
|
||||
_dev = platform.get_devices(cl.device_type.ALL)
|
||||
log.info(f"Platform \"{platform.name}\" have {len(_dev)} devices: {[_d.name for _d in _dev]}")
|
||||
cl_devices.extend(_dev)
|
||||
|
||||
print()
|
||||
|
||||
for cl_device in cl_devices:
|
||||
log.info(f"Testing on {cl.device_type.to_string(cl_device.type)} device: {cl_device.name}")
|
||||
cl_ctx = cl.Context(devices=[cl_device])
|
||||
cl_queue = cl.CommandQueue(cl_ctx)
|
||||
cl_prg = load_program([cl_incl_dir], cl_src_file, cl_ctx)
|
||||
|
||||
encrypt_data_kernel = cl_prg.encrypt_data
|
||||
decrypt_data_kernel = cl_prg.decrypt_data
|
||||
|
||||
mf = cl.mem_flags
|
||||
|
||||
clb_key = cl.Buffer(cl_ctx, mf.READ_ONLY, size=32)
|
||||
clb_sec_nums = cl.Buffer(cl_ctx, mf.READ_ONLY, size=8 * test_vectors_num)
|
||||
clb_blk_nums = cl.Buffer(cl_ctx, mf.READ_ONLY, size=4 * test_vectors_num)
|
||||
clb_u_data = cl.Buffer(cl_ctx, mf.READ_ONLY, size=16 * test_vectors_num)
|
||||
clb_enc_data = cl.Buffer(cl_ctx, mf.WRITE_ONLY, size=16 * test_vectors_num)
|
||||
|
||||
log.info("Executing encrypt kernel...")
|
||||
|
||||
cl.enqueue_copy(cl_queue, clb_sec_nums, b_sec_nums)
|
||||
cl.enqueue_copy(cl_queue, clb_blk_nums, b_blk_nums)
|
||||
cl.enqueue_copy(cl_queue, clb_key, b_key)
|
||||
cl.enqueue_copy(cl_queue, clb_u_data, b_u_data)
|
||||
|
||||
knl_e = encrypt_data_kernel(cl_queue, (test_vectors_num,), None,
|
||||
clb_sec_nums, clb_blk_nums, clb_key, clb_u_data, clb_enc_data)
|
||||
|
||||
b_test_enc_data = bytearray(test_vectors_num * 16)
|
||||
cl.enqueue_copy(cl_queue, b_test_enc_data, clb_enc_data, wait_for=[knl_e])
|
||||
|
||||
log.info("Executing decrypt kernel...")
|
||||
|
||||
cl.enqueue_copy(cl_queue, clb_sec_nums, b_sec_nums)
|
||||
cl.enqueue_copy(cl_queue, clb_blk_nums, b_blk_nums)
|
||||
cl.enqueue_copy(cl_queue, clb_key, b_key)
|
||||
cl.enqueue_copy(cl_queue, clb_u_data, b_enc_data)
|
||||
|
||||
knl_e = decrypt_data_kernel(cl_queue, (test_vectors_num,), None,
|
||||
clb_sec_nums, clb_blk_nums, clb_key, clb_u_data, clb_enc_data)
|
||||
|
||||
b_test_u_data = bytearray(test_vectors_num * 16)
|
||||
cl.enqueue_copy(cl_queue, b_test_u_data, clb_enc_data, wait_for=[knl_e])
|
||||
|
||||
|
||||
cl_queue.finish()
|
||||
|
||||
log.info("Comparing encryption results...")
|
||||
right = 0
|
||||
wrong = 0
|
||||
for i in range(test_vectors_num):
|
||||
of16 = i * 16
|
||||
if b_enc_data[of16:of16 + 16] == b_test_enc_data[of16:of16 + 16]:
|
||||
right += 1
|
||||
else:
|
||||
wrong += 1
|
||||
log.error(f"Missmatch of test vector {i} !")
|
||||
log.error(f" expected:{b_enc_data[of16:of16 + 16].hex()}")
|
||||
log.error(f" actual:{b_test_enc_data[of16:of16 + 16].hex()}")
|
||||
|
||||
if wrong > 0:
|
||||
log.error(f"Test failed, total of {wrong} mismatch results!")
|
||||
else:
|
||||
log.info("All match! Success!")
|
||||
total_errors += wrong
|
||||
|
||||
|
||||
log.info("Comparing decryption results...")
|
||||
right = 0
|
||||
wrong = 0
|
||||
for i in range(test_vectors_num):
|
||||
of16 = i * 16
|
||||
if b_u_data[of16:of16 + 16] == b_test_u_data[of16:of16 + 16]:
|
||||
right += 1
|
||||
else:
|
||||
wrong += 1
|
||||
log.error(f"Missmatch of test vector {i} !")
|
||||
log.error(f" expected:{b_u_data[of16:of16 + 16].hex()}")
|
||||
log.error(f" actual:{b_test_u_data[of16:of16 + 16].hex()}")
|
||||
|
||||
if wrong > 0:
|
||||
log.error(f"Test failed, total of {wrong} mismatch results!\n")
|
||||
else:
|
||||
log.info("All match! Success!\n")
|
||||
total_errors += wrong
|
||||
|
||||
run_time = time.time() - start_time
|
||||
if total_errors > 0:
|
||||
log.error(f"--- Finish in {run_time} seconds with {total_errors} failed tests ---\n\n")
|
||||
else:
|
||||
log.info(f"--- Finish in {run_time} seconds without errors ---\n\n")
|
||||
|
||||
|
||||
|
||||
def cmd_parse():
|
||||
# Initialize arguments parser
|
||||
parser = argparse.ArgumentParser(
|
||||
prog="test_vectors.py",
|
||||
description="This test suite for OpenCL AES-XTS implementation",
|
||||
epilog="Have a nice day!")
|
||||
|
||||
parser.add_argument("-s", "--cl-src-file",
|
||||
action="store",
|
||||
default="../src/test_aes_xts256_plain.cl",
|
||||
help="OpenCL src file with defined test kernels.",
|
||||
required=False)
|
||||
parser.add_argument("-d", "--cl-incl-dir",
|
||||
action="store",
|
||||
default="../src",
|
||||
help="OpenCL include dir with aes256_xts_plain.cl src file.",
|
||||
required=False)
|
||||
parser.add_argument("-v", "--test-vectors",
|
||||
action="store",
|
||||
default="./data/test_vectors.json",
|
||||
help="Path to json file with test vectors.",
|
||||
required=False)
|
||||
|
||||
parser.add_argument("-l", "--log-file",
|
||||
action="store",
|
||||
default=None,
|
||||
help="Path to optional log file.",
|
||||
required=False)
|
||||
|
||||
return parser.parse_args()
|
||||
|
||||
|
||||
def main():
|
||||
argumets = cmd_parse()
|
||||
# Init logger
|
||||
global log
|
||||
if argumets.log_file is not None:
|
||||
log = nya_stat.get_logger("nyan", log_writers=[ConsoleWriter(), FileWriter(argumets.log_file)])
|
||||
else:
|
||||
log = nya_stat.get_logger("nyan")
|
||||
|
||||
log.start()
|
||||
|
||||
test(cl_src_file=argumets.cl_src_file,
|
||||
cl_incl_dir=argumets.cl_incl_dir,
|
||||
vectors_file=argumets.test_vectors)
|
||||
|
||||
log.stop()
|
||||
|
||||
|
||||
if __name__ == '__main__':
|
||||
main()
|
||||
Reference in New Issue
Block a user