From 53e2b40bad46dd85fca2474c7560cfedefdeebc5 Mon Sep 17 00:00:00 2001 From: jsteube Date: Fri, 2 Feb 2018 14:02:33 +0100 Subject: [PATCH] Fixed a uninitialized value in OpenCL kernels 9720, 9820 and 10420 leading to absurd benchmark performance --- OpenCL/m09720_a0-optimized.cl | 4 ++-- OpenCL/m09720_a1-optimized.cl | 4 ++-- OpenCL/m09720_a3-optimized.cl | 4 ++-- OpenCL/m09820_a0-optimized.cl | 4 ++-- OpenCL/m09820_a1-optimized.cl | 4 ++-- OpenCL/m09820_a3-optimized.cl | 4 ++-- OpenCL/m10420_a0-optimized.cl | 4 ++-- OpenCL/m10420_a1-optimized.cl | 4 ++-- OpenCL/m10420_a3-optimized.cl | 4 ++-- docs/changes.txt | 1 + 10 files changed, 19 insertions(+), 18 deletions(-) diff --git a/OpenCL/m09720_a0-optimized.cl b/OpenCL/m09720_a0-optimized.cl index 9d08af114..6bef77f66 100644 --- a/OpenCL/m09720_a0-optimized.cl +++ b/OpenCL/m09720_a0-optimized.cl @@ -514,8 +514,8 @@ __kernel void m09720_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m09720_a1-optimized.cl b/OpenCL/m09720_a1-optimized.cl index 60b83be23..fb6baafe9 100644 --- a/OpenCL/m09720_a1-optimized.cl +++ b/OpenCL/m09720_a1-optimized.cl @@ -562,8 +562,8 @@ __kernel void m09720_s04 (__global pw_t *pws, __global const kernel_rule_t *rule { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m09720_a3-optimized.cl b/OpenCL/m09720_a3-optimized.cl index 322ee8e66..cca7e9d85 100644 --- a/OpenCL/m09720_a3-optimized.cl +++ b/OpenCL/m09720_a3-optimized.cl @@ -475,8 +475,8 @@ void m09720s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m09820_a0-optimized.cl b/OpenCL/m09820_a0-optimized.cl index 27059d055..15bfc79ca 100644 --- a/OpenCL/m09820_a0-optimized.cl +++ b/OpenCL/m09820_a0-optimized.cl @@ -198,8 +198,8 @@ __kernel void m09820_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m09820_a1-optimized.cl b/OpenCL/m09820_a1-optimized.cl index 9a9a6b973..04cf549fa 100644 --- a/OpenCL/m09820_a1-optimized.cl +++ b/OpenCL/m09820_a1-optimized.cl @@ -246,8 +246,8 @@ __kernel void m09820_s04 (__global pw_t *pws, __global const kernel_rule_t *rule { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m09820_a3-optimized.cl b/OpenCL/m09820_a3-optimized.cl index 0591d9e56..b1459720d 100644 --- a/OpenCL/m09820_a3-optimized.cl +++ b/OpenCL/m09820_a3-optimized.cl @@ -143,8 +143,8 @@ void m09820s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m10420_a0-optimized.cl b/OpenCL/m10420_a0-optimized.cl index 329673edc..3a424b3ad 100644 --- a/OpenCL/m10420_a0-optimized.cl +++ b/OpenCL/m10420_a0-optimized.cl @@ -250,8 +250,8 @@ __kernel void m10420_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m10420_a1-optimized.cl b/OpenCL/m10420_a1-optimized.cl index cfbfaeb7f..6b352c937 100644 --- a/OpenCL/m10420_a1-optimized.cl +++ b/OpenCL/m10420_a1-optimized.cl @@ -308,8 +308,8 @@ __kernel void m10420_s04 (__global pw_t *pws, __global const kernel_rule_t *rule { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m10420_a3-optimized.cl b/OpenCL/m10420_a3-optimized.cl index 2dff3e821..333035c16 100644 --- a/OpenCL/m10420_a3-optimized.cl +++ b/OpenCL/m10420_a3-optimized.cl @@ -260,8 +260,8 @@ void m10420s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/docs/changes.txt b/docs/changes.txt index 1af053a08..93e2ca86e 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -28,6 +28,7 @@ - Fixed a hash parsing problem when using --show/--left with hashes with long salts that required pure kernels - Fixed a mask-length check issue: Return -1 in case the mask-length is not within the password-length range - Fixed a restore issue leading to "Restore value is greater than keyspace" when mask-files or wordlist-folders were used +- Fixed a uninitialized value in OpenCL kernels 9720, 9820 and 10420 leading to absurd benchmark performance - Fixed the maximum password length check in password-reassembling function - Fixed the output of --show if $HEX[] passwords were present within the potfile