From 9c832092df2bd1582b7e3cf695706ac11142bfdc Mon Sep 17 00:00:00 2001 From: jsteube Date: Tue, 10 Oct 2017 01:27:32 +0200 Subject: [PATCH] Fix some function declarations inside LUKS kernel --- OpenCL/inc_luks_aes.cl | 116 ++++++++++++++++++------------------- OpenCL/inc_luks_serpent.cl | 116 ++++++++++++++++++------------------- OpenCL/inc_luks_twofish.cl | 116 ++++++++++++++++++------------------- OpenCL/m14612.cl | 6 +- OpenCL/m14613.cl | 6 +- OpenCL/m14621.cl | 6 +- OpenCL/m14622.cl | 6 +- OpenCL/m14623.cl | 6 +- OpenCL/m14631.cl | 6 +- OpenCL/m14632.cl | 6 +- OpenCL/m14633.cl | 6 +- OpenCL/m14641.cl | 6 +- OpenCL/m14642.cl | 6 +- OpenCL/m14643.cl | 6 +- 14 files changed, 207 insertions(+), 207 deletions(-) diff --git a/OpenCL/inc_luks_aes.cl b/OpenCL/inc_luks_aes.cl index 3b8499820..53ecf6812 100644 --- a/OpenCL/inc_luks_aes.cl +++ b/OpenCL/inc_luks_aes.cl @@ -31,7 +31,7 @@ static void aes256_decrypt_cbc (const u32 *ks1, const u32 *in, u32 *out, u32 *es essiv[3] = in[3]; } -static void luks_decrypt_sector_aes_cbc_essiv128 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -65,7 +65,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128 (__global u32 *in, u32 *out, co } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -102,7 +102,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1 (__global u32 *in, u32 } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -162,7 +162,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1_final (__global u32 *in } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -199,7 +199,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -259,7 +259,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -296,7 +296,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -356,7 +356,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -393,7 +393,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160 (__global u32 *in, } } -static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -453,7 +453,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160_final (__global u3 } } -static void luks_decrypt_sector_aes_cbc_essiv256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -487,7 +487,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256 (__global u32 *in, u32 *out, co } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -524,7 +524,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1 (__global u32 *in, u32 } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -584,7 +584,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1_final (__global u32 *in } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -621,7 +621,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -681,7 +681,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -718,7 +718,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -778,7 +778,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -815,7 +815,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160 (__global u32 *in, } } -static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -877,7 +877,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160_final (__global u3 // cbc-plain -static void luks_decrypt_sector_aes_cbc_plain128 (__global u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -907,7 +907,7 @@ static void luks_decrypt_sector_aes_cbc_plain128 (__global u32 *in, u32 *out, co } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -940,7 +940,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1 (__global u32 *in, u32 } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -996,7 +996,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1_final (__global u32 *in } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1029,7 +1029,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1085,7 +1085,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1118,7 +1118,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1174,7 +1174,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1207,7 +1207,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160 (__global u32 *in, } } -static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1263,7 +1263,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160_final (__global u3 } } -static void luks_decrypt_sector_aes_cbc_plain256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1293,7 +1293,7 @@ static void luks_decrypt_sector_aes_cbc_plain256 (__global u32 *in, u32 *out, co } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1326,7 +1326,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1 (__global u32 *in, u32 } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1382,7 +1382,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1_final (__global u32 *in } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1415,7 +1415,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1471,7 +1471,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1504,7 +1504,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1560,7 +1560,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512_final (__global u32 * } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1593,7 +1593,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160 (__global u32 *in, } } -static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1695,7 +1695,7 @@ static void aes256_decrypt_xts (const u32 *ks1, const u32 *in, u32 *out, u32 *T, xts_mul2 (T, T); } -static void luks_decrypt_sector_aes_xts_plain256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1729,7 +1729,7 @@ static void luks_decrypt_sector_aes_xts_plain256 (__global u32 *in, u32 *out, co } } -static void luks_decrypt_sector_aes_xts_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1766,7 +1766,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha1 (__global u32 *in, u32 } } -static void luks_decrypt_sector_aes_xts_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1826,7 +1826,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha1_final (__global u32 *in } } -static void luks_decrypt_sector_aes_xts_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1863,7 +1863,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha256 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_xts_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1923,7 +1923,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha256_final (__global u32 * } } -static void luks_decrypt_sector_aes_xts_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1960,7 +1960,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha512 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_xts_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2020,7 +2020,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha512_final (__global u32 * } } -static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2057,7 +2057,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160 (__global u32 *in, } } -static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2117,7 +2117,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160_final (__global u3 } } -static void luks_decrypt_sector_aes_xts_plain512 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2151,7 +2151,7 @@ static void luks_decrypt_sector_aes_xts_plain512 (__global u32 *in, u32 *out, co } } -static void luks_decrypt_sector_aes_xts_plain512_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2188,7 +2188,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha1 (__global u32 *in, u32 } } -static void luks_decrypt_sector_aes_xts_plain512_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2248,7 +2248,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha1_final (__global u32 *in } } -static void luks_decrypt_sector_aes_xts_plain512_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2285,7 +2285,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha256 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_xts_plain512_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2345,7 +2345,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha256_final (__global u32 * } } -static void luks_decrypt_sector_aes_xts_plain512_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2382,7 +2382,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha512 (__global u32 *in, u3 } } -static void luks_decrypt_sector_aes_xts_plain512_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2442,7 +2442,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha512_final (__global u32 * } } -static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2479,7 +2479,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160 (__global u32 *in, } } -static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2541,7 +2541,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160_final (__global u3 // luks helper -static void luks_af_sha1_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_af_sha1_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -2932,7 +2932,7 @@ static void luks_af_sha1_then_aes_decrypt (__global luks_t *luks_bufs, __global } } -static void luks_af_sha256_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_af_sha256_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -3291,7 +3291,7 @@ static void luks_af_sha256_then_aes_decrypt (__global luks_t *luks_bufs, __globa } } -static void luks_af_sha512_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_af_sha512_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -3669,7 +3669,7 @@ static void luks_af_sha512_then_aes_decrypt (__global luks_t *luks_bufs, __globa } } -static void luks_af_ripemd160_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) +static void luks_af_ripemd160_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; diff --git a/OpenCL/inc_luks_serpent.cl b/OpenCL/inc_luks_serpent.cl index 00068776b..9f0dcf2cf 100644 --- a/OpenCL/inc_luks_serpent.cl +++ b/OpenCL/inc_luks_serpent.cl @@ -31,7 +31,7 @@ static void serpent256_decrypt_cbc (const u32 *ks1, const u32 *in, u32 *out, u32 essiv[3] = in[3]; } -static void luks_decrypt_sector_serpent_cbc_essiv128 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -65,7 +65,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -102,7 +102,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -162,7 +162,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -199,7 +199,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -259,7 +259,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -296,7 +296,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -356,7 +356,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -393,7 +393,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -453,7 +453,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160_final (__globa } } -static void luks_decrypt_sector_serpent_cbc_essiv256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -487,7 +487,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -524,7 +524,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -584,7 +584,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -621,7 +621,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -681,7 +681,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -718,7 +718,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -778,7 +778,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -815,7 +815,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -877,7 +877,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160_final (__globa // cbc-plain -static void luks_decrypt_sector_serpent_cbc_plain128 (__global u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128 (__global const u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -907,7 +907,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -940,7 +940,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -996,7 +996,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1029,7 +1029,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1085,7 +1085,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1118,7 +1118,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1174,7 +1174,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1207,7 +1207,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1263,7 +1263,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160_final (__globa } } -static void luks_decrypt_sector_serpent_cbc_plain256 (__global u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1293,7 +1293,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1326,7 +1326,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1382,7 +1382,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1415,7 +1415,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1471,7 +1471,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1504,7 +1504,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1560,7 +1560,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512_final (__global u } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1593,7 +1593,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) +static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1695,7 +1695,7 @@ static void serpent256_decrypt_xts (const u32 *ks1, const u32 *in, u32 *out, u32 xts_mul2 (T, T); } -static void luks_decrypt_sector_serpent_xts_plain256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1729,7 +1729,7 @@ static void luks_decrypt_sector_serpent_xts_plain256 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1766,7 +1766,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1826,7 +1826,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1863,7 +1863,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1923,7 +1923,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256_final (__global u } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1960,7 +1960,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2020,7 +2020,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512_final (__global u } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2057,7 +2057,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2117,7 +2117,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160_final (__globa } } -static void luks_decrypt_sector_serpent_xts_plain512 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2151,7 +2151,7 @@ static void luks_decrypt_sector_serpent_xts_plain512 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2188,7 +2188,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2248,7 +2248,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2285,7 +2285,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2345,7 +2345,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256_final (__global u } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2382,7 +2382,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2442,7 +2442,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512_final (__global u } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2479,7 +2479,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) +static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2541,7 +2541,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160_final (__globa // luks helper -static void luks_af_sha1_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_sha1_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -2900,7 +2900,7 @@ static void luks_af_sha1_then_serpent_decrypt (__global luks_t *luks_bufs, __glo } } -static void luks_af_sha256_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_sha256_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -3259,7 +3259,7 @@ static void luks_af_sha256_then_serpent_decrypt (__global luks_t *luks_bufs, __g } } -static void luks_af_sha512_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_sha512_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -3637,7 +3637,7 @@ static void luks_af_sha512_then_serpent_decrypt (__global luks_t *luks_bufs, __g } } -static void luks_af_ripemd160_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_ripemd160_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; diff --git a/OpenCL/inc_luks_twofish.cl b/OpenCL/inc_luks_twofish.cl index 38800a872..e17ac1c1b 100644 --- a/OpenCL/inc_luks_twofish.cl +++ b/OpenCL/inc_luks_twofish.cl @@ -31,7 +31,7 @@ static void twofish256_decrypt_cbc (const u32 *sk1, const u32 *lk1, const u32 *i essiv[3] = in[3]; } -static void luks_decrypt_sector_twofish_cbc_essiv128 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -65,7 +65,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -102,7 +102,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -162,7 +162,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -199,7 +199,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -259,7 +259,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -296,7 +296,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -356,7 +356,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -393,7 +393,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -453,7 +453,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160_final (__globa } } -static void luks_decrypt_sector_twofish_cbc_essiv256 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -487,7 +487,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -524,7 +524,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -584,7 +584,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -621,7 +621,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -681,7 +681,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -718,7 +718,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -778,7 +778,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -815,7 +815,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -877,7 +877,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160_final (__globa // cbc-plain -static void luks_decrypt_sector_twofish_cbc_plain128 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -907,7 +907,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -940,7 +940,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -996,7 +996,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1029,7 +1029,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1085,7 +1085,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1118,7 +1118,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1174,7 +1174,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1207,7 +1207,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1263,7 +1263,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160_final (__globa } } -static void luks_decrypt_sector_twofish_cbc_plain256 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1293,7 +1293,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1326,7 +1326,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1382,7 +1382,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1415,7 +1415,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1471,7 +1471,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1504,7 +1504,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1560,7 +1560,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512_final (__global u } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1593,7 +1593,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) +static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1695,7 +1695,7 @@ static void twofish256_decrypt_xts (const u32 *sk1, const u32 *lk1, const u32 *i xts_mul2 (T, T); } -static void luks_decrypt_sector_twofish_xts_plain256 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1729,7 +1729,7 @@ static void luks_decrypt_sector_twofish_xts_plain256 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1766,7 +1766,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1826,7 +1826,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1863,7 +1863,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1923,7 +1923,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256_final (__global u } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -1960,7 +1960,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2020,7 +2020,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512_final (__global u } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2057,7 +2057,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2117,7 +2117,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160_final (__globa } } -static void luks_decrypt_sector_twofish_xts_plain512 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2151,7 +2151,7 @@ static void luks_decrypt_sector_twofish_xts_plain512 (__global u32 *in, u32 *out } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2188,7 +2188,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1 (__global u32 *in, } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2248,7 +2248,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1_final (__global u32 } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2285,7 +2285,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256 (__global u32 *in } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2345,7 +2345,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256_final (__global u } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2382,7 +2382,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512 (__global u32 *in } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2442,7 +2442,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512_final (__global u } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2479,7 +2479,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160 (__global u32 } } -static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) +static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector) { u32 S[4] = { sector, 0, 0, 0 }; @@ -2541,7 +2541,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160_final (__globa // luks helper -static void luks_af_sha1_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_sha1_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -2900,7 +2900,7 @@ static void luks_af_sha1_then_twofish_decrypt (__global luks_t *luks_bufs, __glo } } -static void luks_af_sha256_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_sha256_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -3259,7 +3259,7 @@ static void luks_af_sha256_then_twofish_decrypt (__global luks_t *luks_bufs, __g } } -static void luks_af_sha512_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_sha512_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; @@ -3637,7 +3637,7 @@ static void luks_af_sha512_then_twofish_decrypt (__global luks_t *luks_bufs, __g } } -static void luks_af_ripemd160_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) +static void luks_af_ripemd160_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf) { const u32 key_size = luks_bufs->key_size; const u32 cipher_mode = luks_bufs->cipher_mode; diff --git a/OpenCL/m14612.cl b/OpenCL/m14612.cl index 501a5daf5..992429ca2 100644 --- a/OpenCL/m14612.cl +++ b/OpenCL/m14612.cl @@ -61,7 +61,7 @@ void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[ sha1_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -135,7 +135,7 @@ __kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14612_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14612_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -222,7 +222,7 @@ __kernel void m14612_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14612_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14612_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); diff --git a/OpenCL/m14613.cl b/OpenCL/m14613.cl index e99737a09..d328e3cb6 100644 --- a/OpenCL/m14613.cl +++ b/OpenCL/m14613.cl @@ -61,7 +61,7 @@ void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[ sha1_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -135,7 +135,7 @@ __kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14613_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14613_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -222,7 +222,7 @@ __kernel void m14613_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14613_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14613_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); diff --git a/OpenCL/m14621.cl b/OpenCL/m14621.cl index 3c83d7cb6..622945efd 100644 --- a/OpenCL/m14621.cl +++ b/OpenCL/m14621.cl @@ -67,7 +67,7 @@ void hmac_sha256_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipa sha256_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -153,7 +153,7 @@ __kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14621_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14621_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -261,7 +261,7 @@ __kernel void m14621_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14621_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14621_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m14622.cl b/OpenCL/m14622.cl index 03ff787c6..a66ec0e93 100644 --- a/OpenCL/m14622.cl +++ b/OpenCL/m14622.cl @@ -67,7 +67,7 @@ void hmac_sha256_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipa sha256_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -153,7 +153,7 @@ __kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14622_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14622_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -261,7 +261,7 @@ __kernel void m14622_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14622_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14622_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); diff --git a/OpenCL/m14623.cl b/OpenCL/m14623.cl index 7859fa949..79937b817 100644 --- a/OpenCL/m14623.cl +++ b/OpenCL/m14623.cl @@ -67,7 +67,7 @@ void hmac_sha256_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipa sha256_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -153,7 +153,7 @@ __kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14623_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14623_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -261,7 +261,7 @@ __kernel void m14623_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14623_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14623_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); diff --git a/OpenCL/m14631.cl b/OpenCL/m14631.cl index bc1c01bab..0d4a7f84f 100644 --- a/OpenCL/m14631.cl +++ b/OpenCL/m14631.cl @@ -83,7 +83,7 @@ void hmac_sha512_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[ sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest); } -__kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -189,7 +189,7 @@ __kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14631_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14631_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -317,7 +317,7 @@ __kernel void m14631_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14631_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14631_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m14632.cl b/OpenCL/m14632.cl index 38d2ef114..86f273ee9 100644 --- a/OpenCL/m14632.cl +++ b/OpenCL/m14632.cl @@ -83,7 +83,7 @@ void hmac_sha512_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[ sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest); } -__kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -189,7 +189,7 @@ __kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14632_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14632_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -317,7 +317,7 @@ __kernel void m14632_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14632_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14632_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); diff --git a/OpenCL/m14633.cl b/OpenCL/m14633.cl index 5aabd1007..ae2e98a85 100644 --- a/OpenCL/m14633.cl +++ b/OpenCL/m14633.cl @@ -83,7 +83,7 @@ void hmac_sha512_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[ sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest); } -__kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -189,7 +189,7 @@ __kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14633_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14633_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -317,7 +317,7 @@ __kernel void m14633_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14633_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14633_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); diff --git a/OpenCL/m14641.cl b/OpenCL/m14641.cl index 12d3f7215..240c1fb43 100644 --- a/OpenCL/m14641.cl +++ b/OpenCL/m14641.cl @@ -61,7 +61,7 @@ void hmac_ripemd160_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ripemd160_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -135,7 +135,7 @@ __kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14641_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14641_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -222,7 +222,7 @@ __kernel void m14641_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14641_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14641_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); diff --git a/OpenCL/m14642.cl b/OpenCL/m14642.cl index 7bb6caf41..5bad9397d 100644 --- a/OpenCL/m14642.cl +++ b/OpenCL/m14642.cl @@ -61,7 +61,7 @@ void hmac_ripemd160_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ripemd160_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -135,7 +135,7 @@ __kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14642_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14642_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -222,7 +222,7 @@ __kernel void m14642_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14642_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14642_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); diff --git a/OpenCL/m14643.cl b/OpenCL/m14643.cl index 0ef4addbf..0bb93ccee 100644 --- a/OpenCL/m14643.cl +++ b/OpenCL/m14643.cl @@ -61,7 +61,7 @@ void hmac_ripemd160_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ripemd160_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { /** * base @@ -135,7 +135,7 @@ __kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14643_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14643_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0); @@ -222,7 +222,7 @@ __kernel void m14643_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m14643_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) +__kernel void m14643_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { const u64 gid = get_global_id (0);