From 0f906dcfeb1d8e34124b339cc326a764c6771738 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Sat, 26 Dec 2015 17:10:03 +0100 Subject: [PATCH] Fix a bug in the implementation of GOST R 34.11-94, zero length passwords were not cracked --- OpenCL/m06900_a0.cl | 38 ++++++++++++++++++++++---------------- OpenCL/m06900_a1.cl | 38 ++++++++++++++++++++++---------------- OpenCL/m06900_a3.cl | 38 ++++++++++++++++++++++---------------- docs/changes.txt | 4 ++++ 4 files changed, 70 insertions(+), 48 deletions(-) diff --git a/OpenCL/m06900_a0.cl b/OpenCL/m06900_a0.cl index 31a5b6f..43a4fcd 100644 --- a/OpenCL/m06900_a0.cl +++ b/OpenCL/m06900_a0.cl @@ -850,14 +850,17 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06900_m04 (__glo u32 tmp[8]; - PASS0 (state, tmp, state_m, data_m, s_tables); - PASS2 (state, tmp, state_m, data_m, s_tables); - PASS4 (state, tmp, state_m, data_m, s_tables); - PASS6 (state, tmp, state_m, data_m, s_tables); - - SHIFT12 (state_m, data, tmp); - SHIFT16 (state, data_m, state_m); - SHIFT61 (state, data_m); + if (pw_len > 0) + { + PASS0 (state, tmp, state_m, data_m, s_tables); + PASS2 (state, tmp, state_m, data_m, s_tables); + PASS4 (state, tmp, state_m, data_m, s_tables); + PASS6 (state, tmp, state_m, data_m, s_tables); + + SHIFT12 (state_m, data, tmp); + SHIFT16 (state, data_m, state_m); + SHIFT61 (state, data_m); + } data[0] = w14; data[1] = 0; @@ -1121,14 +1124,17 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06900_s04 (__glo u32 tmp[8]; - PASS0 (state, tmp, state_m, data_m, s_tables); - PASS2 (state, tmp, state_m, data_m, s_tables); - PASS4 (state, tmp, state_m, data_m, s_tables); - PASS6 (state, tmp, state_m, data_m, s_tables); - - SHIFT12 (state_m, data, tmp); - SHIFT16 (state, data_m, state_m); - SHIFT61 (state, data_m); + if (pw_len > 0) + { + PASS0 (state, tmp, state_m, data_m, s_tables); + PASS2 (state, tmp, state_m, data_m, s_tables); + PASS4 (state, tmp, state_m, data_m, s_tables); + PASS6 (state, tmp, state_m, data_m, s_tables); + + SHIFT12 (state_m, data, tmp); + SHIFT16 (state, data_m, state_m); + SHIFT61 (state, data_m); + } data[0] = w14; data[1] = 0; diff --git a/OpenCL/m06900_a1.cl b/OpenCL/m06900_a1.cl index 19d2ace..28d1a8d 100644 --- a/OpenCL/m06900_a1.cl +++ b/OpenCL/m06900_a1.cl @@ -902,14 +902,17 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06900_m04 (__glo u32 tmp[8]; - PASS0 (state, tmp, state_m, data_m, s_tables); - PASS2 (state, tmp, state_m, data_m, s_tables); - PASS4 (state, tmp, state_m, data_m, s_tables); - PASS6 (state, tmp, state_m, data_m, s_tables); - - SHIFT12 (state_m, data, tmp); - SHIFT16 (state, data_m, state_m); - SHIFT61 (state, data_m); + if (pw_len > 0) + { + PASS0 (state, tmp, state_m, data_m, s_tables); + PASS2 (state, tmp, state_m, data_m, s_tables); + PASS4 (state, tmp, state_m, data_m, s_tables); + PASS6 (state, tmp, state_m, data_m, s_tables); + + SHIFT12 (state_m, data, tmp); + SHIFT16 (state, data_m, state_m); + SHIFT61 (state, data_m); + } data[0] = w14; data[1] = 0; @@ -1227,14 +1230,17 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06900_s04 (__glo u32 tmp[8]; - PASS0 (state, tmp, state_m, data_m, s_tables); - PASS2 (state, tmp, state_m, data_m, s_tables); - PASS4 (state, tmp, state_m, data_m, s_tables); - PASS6 (state, tmp, state_m, data_m, s_tables); - - SHIFT12 (state_m, data, tmp); - SHIFT16 (state, data_m, state_m); - SHIFT61 (state, data_m); + if (pw_len > 0) + { + PASS0 (state, tmp, state_m, data_m, s_tables); + PASS2 (state, tmp, state_m, data_m, s_tables); + PASS4 (state, tmp, state_m, data_m, s_tables); + PASS6 (state, tmp, state_m, data_m, s_tables); + + SHIFT12 (state_m, data, tmp); + SHIFT16 (state, data_m, state_m); + SHIFT61 (state, data_m); + } data[0] = w14; data[1] = 0; diff --git a/OpenCL/m06900_a3.cl b/OpenCL/m06900_a3.cl index 1c4f3b3..fdb442b 100644 --- a/OpenCL/m06900_a3.cl +++ b/OpenCL/m06900_a3.cl @@ -775,14 +775,17 @@ static void m06900m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le u32 tmp[8]; - PASS0 (state, tmp, state_m, data_m, s_tables); - PASS2 (state, tmp, state_m, data_m, s_tables); - PASS4 (state, tmp, state_m, data_m, s_tables); - PASS6 (state, tmp, state_m, data_m, s_tables); - - SHIFT12 (state_m, data, tmp); - SHIFT16 (state, data_m, state_m); - SHIFT61 (state, data_m); + if (pw_len > 0) + { + PASS0 (state, tmp, state_m, data_m, s_tables); + PASS2 (state, tmp, state_m, data_m, s_tables); + PASS4 (state, tmp, state_m, data_m, s_tables); + PASS6 (state, tmp, state_m, data_m, s_tables); + + SHIFT12 (state_m, data, tmp); + SHIFT16 (state, data_m, state_m); + SHIFT61 (state, data_m); + } data[0] = w14; data[1] = 0; @@ -964,14 +967,17 @@ static void m06900s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le u32 tmp[8]; - PASS0 (state, tmp, state_m, data_m, s_tables); - PASS2 (state, tmp, state_m, data_m, s_tables); - PASS4 (state, tmp, state_m, data_m, s_tables); - PASS6 (state, tmp, state_m, data_m, s_tables); - - SHIFT12 (state_m, data, tmp); - SHIFT16 (state, data_m, state_m); - SHIFT61 (state, data_m); + if (pw_len > 0) + { + PASS0 (state, tmp, state_m, data_m, s_tables); + PASS2 (state, tmp, state_m, data_m, s_tables); + PASS4 (state, tmp, state_m, data_m, s_tables); + PASS6 (state, tmp, state_m, data_m, s_tables); + + SHIFT12 (state_m, data, tmp); + SHIFT16 (state, data_m, state_m); + SHIFT61 (state, data_m); + } data[0] = w14; data[1] = 0; diff --git a/docs/changes.txt b/docs/changes.txt index 1b8abca..7c79538 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -19,6 +19,10 @@ File.: Host Desc.: Fixed a bug in combination of --restore and a user immediately aborting the session after restart Trac.: 684 +Type.: Bug +File.: Kernel +Desc.: Fix a bug in the implementation of GOST R 34.11-94, zero length passwords were not cracked + * changes v2.00 -> v2.01: Type.: Bug -- 2.43.0