From c4333c4025463308bcdc5a55b809f7017879e64f Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 21 May 2016 21:23:22 +0200 Subject: [PATCH] Get SHA256 KDF for VeraCrypt to work --- OpenCL/m13751.cl | 79 ++++++++++++++------------ OpenCL/m13752.cl | 110 +++++++++++++++++++----------------- OpenCL/m13753.cl | 142 +++++++++++++++++++++++++---------------------- 3 files changed, 180 insertions(+), 151 deletions(-) diff --git a/OpenCL/m13751.cl b/OpenCL/m13751.cl index 3892cde..06e3fab 100644 --- a/OpenCL/m13751.cl +++ b/OpenCL/m13751.cl @@ -201,22 +201,27 @@ void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u sha256_transform (w0, w1, w2, w3, digest); - w0[0] = digest[0]; - w0[1] = digest[1]; - w0[2] = digest[2]; - w0[3] = digest[3]; - w1[0] = digest[4]; - w1[1] = digest[5]; - w1[2] = digest[6]; - w1[3] = digest[7]; - w2[0] = 0x80000000; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = (64 + 32) * 8; + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = digest[4]; + t1[1] = digest[5]; + t1[2] = digest[6]; + t1[3] = digest[7]; + t2[0] = 0x80000000; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = (64 + 32) * 8; digest[0] = opad[0]; digest[1] = opad[1]; @@ -227,7 +232,7 @@ void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u digest[6] = opad[6]; digest[7] = opad[7]; - sha256_transform (w0, w1, w2, w3, digest); + sha256_transform (t0, t1, t2, t3, digest); } void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u32 w5[4], u32 w6[4], u32 w7[4], u32 ipad[8], u32 opad[8], u32 digest[8]) @@ -244,22 +249,27 @@ void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u3 sha256_transform (w0, w1, w2, w3, digest); sha256_transform (w4, w5, w6, w7, digest); - w0[0] = digest[0]; - w0[1] = digest[1]; - w0[2] = digest[2]; - w0[3] = digest[3]; - w1[0] = digest[4]; - w1[1] = digest[5]; - w1[2] = digest[6]; - w1[3] = digest[7]; - w2[0] = 0x80000000; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = (64 + 32) * 8; + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = digest[4]; + t1[1] = digest[5]; + t1[2] = digest[6]; + t1[3] = digest[7]; + t2[0] = 0x80000000; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = (64 + 32) * 8; digest[0] = opad[0]; digest[1] = opad[1]; @@ -270,7 +280,7 @@ void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u3 digest[6] = opad[6]; digest[7] = opad[7]; - sha256_transform (w0, w1, w2, w3, digest); + sha256_transform (t0, t1, t2, t3, digest); } u32 u8add (const u32 a, const u32 b) @@ -663,4 +673,3 @@ __kernel void m13751_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf } } } - diff --git a/OpenCL/m13752.cl b/OpenCL/m13752.cl index b392ad2..0613321 100644 --- a/OpenCL/m13752.cl +++ b/OpenCL/m13752.cl @@ -201,22 +201,27 @@ void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u sha256_transform (w0, w1, w2, w3, digest); - w0[0] = digest[0]; - w0[1] = digest[1]; - w0[2] = digest[2]; - w0[3] = digest[3]; - w1[0] = digest[4]; - w1[1] = digest[5]; - w1[2] = digest[6]; - w1[3] = digest[7]; - w2[0] = 0x80000000; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = (64 + 32) * 8; + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = digest[4]; + t1[1] = digest[5]; + t1[2] = digest[6]; + t1[3] = digest[7]; + t2[0] = 0x80000000; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = (64 + 32) * 8; digest[0] = opad[0]; digest[1] = opad[1]; @@ -227,7 +232,7 @@ void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u digest[6] = opad[6]; digest[7] = opad[7]; - sha256_transform (w0, w1, w2, w3, digest); + sha256_transform (t0, t1, t2, t3, digest); } void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u32 w5[4], u32 w6[4], u32 w7[4], u32 ipad[8], u32 opad[8], u32 digest[8]) @@ -244,22 +249,27 @@ void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u3 sha256_transform (w0, w1, w2, w3, digest); sha256_transform (w4, w5, w6, w7, digest); - w0[0] = digest[0]; - w0[1] = digest[1]; - w0[2] = digest[2]; - w0[3] = digest[3]; - w1[0] = digest[4]; - w1[1] = digest[5]; - w1[2] = digest[6]; - w1[3] = digest[7]; - w2[0] = 0x80000000; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = (64 + 32) * 8; + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = digest[4]; + t1[1] = digest[5]; + t1[2] = digest[6]; + t1[3] = digest[7]; + t2[0] = 0x80000000; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = (64 + 32) * 8; digest[0] = opad[0]; digest[1] = opad[1]; @@ -270,7 +280,7 @@ void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u3 digest[6] = opad[6]; digest[7] = opad[7]; - sha256_transform (w0, w1, w2, w3, digest); + sha256_transform (t0, t1, t2, t3, digest); } u32 u8add (const u32 a, const u32 b) @@ -665,25 +675,25 @@ __kernel void m13752_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf u32 ukey3[8]; - ukey3[0] = tmps[gid].out[16]; - ukey3[1] = tmps[gid].out[17]; - ukey3[2] = tmps[gid].out[18]; - ukey3[3] = tmps[gid].out[19]; - ukey3[4] = tmps[gid].out[20]; - ukey3[5] = tmps[gid].out[21]; - ukey3[6] = tmps[gid].out[22]; - ukey3[7] = tmps[gid].out[23]; + ukey3[0] = swap32 (tmps[gid].out[16]); + ukey3[1] = swap32 (tmps[gid].out[17]); + ukey3[2] = swap32 (tmps[gid].out[18]); + ukey3[3] = swap32 (tmps[gid].out[19]); + ukey3[4] = swap32 (tmps[gid].out[20]); + ukey3[5] = swap32 (tmps[gid].out[21]); + ukey3[6] = swap32 (tmps[gid].out[22]); + ukey3[7] = swap32 (tmps[gid].out[23]); u32 ukey4[8]; - ukey4[0] = tmps[gid].out[24]; - ukey4[1] = tmps[gid].out[25]; - ukey4[2] = tmps[gid].out[26]; - ukey4[3] = tmps[gid].out[27]; - ukey4[4] = tmps[gid].out[28]; - ukey4[5] = tmps[gid].out[29]; - ukey4[6] = tmps[gid].out[30]; - ukey4[7] = tmps[gid].out[31]; + ukey4[0] = swap32 (tmps[gid].out[24]); + ukey4[1] = swap32 (tmps[gid].out[25]); + ukey4[2] = swap32 (tmps[gid].out[26]); + ukey4[3] = swap32 (tmps[gid].out[27]); + ukey4[4] = swap32 (tmps[gid].out[28]); + ukey4[5] = swap32 (tmps[gid].out[29]); + ukey4[6] = swap32 (tmps[gid].out[30]); + ukey4[7] = swap32 (tmps[gid].out[31]); { tmp[0] = data[0]; diff --git a/OpenCL/m13753.cl b/OpenCL/m13753.cl index 83f54cd..be1a9e9 100644 --- a/OpenCL/m13753.cl +++ b/OpenCL/m13753.cl @@ -201,22 +201,27 @@ void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u sha256_transform (w0, w1, w2, w3, digest); - w0[0] = digest[0]; - w0[1] = digest[1]; - w0[2] = digest[2]; - w0[3] = digest[3]; - w1[0] = digest[4]; - w1[1] = digest[5]; - w1[2] = digest[6]; - w1[3] = digest[7]; - w2[0] = 0x80000000; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = (64 + 32) * 8; + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = digest[4]; + t1[1] = digest[5]; + t1[2] = digest[6]; + t1[3] = digest[7]; + t2[0] = 0x80000000; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = (64 + 32) * 8; digest[0] = opad[0]; digest[1] = opad[1]; @@ -227,7 +232,7 @@ void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u digest[6] = opad[6]; digest[7] = opad[7]; - sha256_transform (w0, w1, w2, w3, digest); + sha256_transform (t0, t1, t2, t3, digest); } void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u32 w5[4], u32 w6[4], u32 w7[4], u32 ipad[8], u32 opad[8], u32 digest[8]) @@ -244,22 +249,27 @@ void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u3 sha256_transform (w0, w1, w2, w3, digest); sha256_transform (w4, w5, w6, w7, digest); - w0[0] = digest[0]; - w0[1] = digest[1]; - w0[2] = digest[2]; - w0[3] = digest[3]; - w1[0] = digest[4]; - w1[1] = digest[5]; - w1[2] = digest[6]; - w1[3] = digest[7]; - w2[0] = 0x80000000; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = (64 + 32) * 8; + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = digest[4]; + t1[1] = digest[5]; + t1[2] = digest[6]; + t1[3] = digest[7]; + t2[0] = 0x80000000; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = (64 + 32) * 8; digest[0] = opad[0]; digest[1] = opad[1]; @@ -270,7 +280,7 @@ void hmac_sha256_run2 (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u3 digest[6] = opad[6]; digest[7] = opad[7]; - sha256_transform (w0, w1, w2, w3, digest); + sha256_transform (t0, t1, t2, t3, digest); } u32 u8add (const u32 a, const u32 b) @@ -665,25 +675,25 @@ __kernel void m13753_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf u32 ukey3[8]; - ukey3[0] = tmps[gid].out[16]; - ukey3[1] = tmps[gid].out[17]; - ukey3[2] = tmps[gid].out[18]; - ukey3[3] = tmps[gid].out[19]; - ukey3[4] = tmps[gid].out[20]; - ukey3[5] = tmps[gid].out[21]; - ukey3[6] = tmps[gid].out[22]; - ukey3[7] = tmps[gid].out[23]; + ukey3[0] = swap32 (tmps[gid].out[16]); + ukey3[1] = swap32 (tmps[gid].out[17]); + ukey3[2] = swap32 (tmps[gid].out[18]); + ukey3[3] = swap32 (tmps[gid].out[19]); + ukey3[4] = swap32 (tmps[gid].out[20]); + ukey3[5] = swap32 (tmps[gid].out[21]); + ukey3[6] = swap32 (tmps[gid].out[22]); + ukey3[7] = swap32 (tmps[gid].out[23]); u32 ukey4[8]; - ukey4[0] = tmps[gid].out[24]; - ukey4[1] = tmps[gid].out[25]; - ukey4[2] = tmps[gid].out[26]; - ukey4[3] = tmps[gid].out[27]; - ukey4[4] = tmps[gid].out[28]; - ukey4[5] = tmps[gid].out[29]; - ukey4[6] = tmps[gid].out[30]; - ukey4[7] = tmps[gid].out[31]; + ukey4[0] = swap32 (tmps[gid].out[24]); + ukey4[1] = swap32 (tmps[gid].out[25]); + ukey4[2] = swap32 (tmps[gid].out[26]); + ukey4[3] = swap32 (tmps[gid].out[27]); + ukey4[4] = swap32 (tmps[gid].out[28]); + ukey4[5] = swap32 (tmps[gid].out[29]); + ukey4[6] = swap32 (tmps[gid].out[30]); + ukey4[7] = swap32 (tmps[gid].out[31]); { tmp[0] = data[0]; @@ -732,25 +742,25 @@ __kernel void m13753_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf u32 ukey5[8]; - ukey5[0] = tmps[gid].out[32]; - ukey5[1] = tmps[gid].out[33]; - ukey5[2] = tmps[gid].out[34]; - ukey5[3] = tmps[gid].out[35]; - ukey5[4] = tmps[gid].out[36]; - ukey5[5] = tmps[gid].out[37]; - ukey5[6] = tmps[gid].out[38]; - ukey5[7] = tmps[gid].out[39]; + ukey5[0] = swap32 (tmps[gid].out[32]); + ukey5[1] = swap32 (tmps[gid].out[33]); + ukey5[2] = swap32 (tmps[gid].out[34]); + ukey5[3] = swap32 (tmps[gid].out[35]); + ukey5[4] = swap32 (tmps[gid].out[36]); + ukey5[5] = swap32 (tmps[gid].out[37]); + ukey5[6] = swap32 (tmps[gid].out[38]); + ukey5[7] = swap32 (tmps[gid].out[39]); u32 ukey6[8]; - ukey6[0] = tmps[gid].out[40]; - ukey6[1] = tmps[gid].out[41]; - ukey6[2] = tmps[gid].out[42]; - ukey6[3] = tmps[gid].out[43]; - ukey6[4] = tmps[gid].out[44]; - ukey6[5] = tmps[gid].out[45]; - ukey6[6] = tmps[gid].out[46]; - ukey6[7] = tmps[gid].out[47]; + ukey6[0] = swap32 (tmps[gid].out[40]); + ukey6[1] = swap32 (tmps[gid].out[41]); + ukey6[2] = swap32 (tmps[gid].out[42]); + ukey6[3] = swap32 (tmps[gid].out[43]); + ukey6[4] = swap32 (tmps[gid].out[44]); + ukey6[5] = swap32 (tmps[gid].out[45]); + ukey6[6] = swap32 (tmps[gid].out[46]); + ukey6[7] = swap32 (tmps[gid].out[47]); { tmp[0] = data[0]; -- 2.25.1