From 2dadae4e9ac6f96ffb905bf6d46cb2cf5b963a3c Mon Sep 17 00:00:00 2001 From: philsmd Date: Sun, 17 Sep 2017 15:28:24 +0200 Subject: [PATCH] fixed incorrect use of the esalt_bufs for -m 600 = BLAKE2-512 --- OpenCL/m00600_a0-optimized.cl | 48 +++++++++++++++++------------------ OpenCL/m00600_a1-optimized.cl | 48 +++++++++++++++++------------------ OpenCL/m00600_a3-optimized.cl | 48 +++++++++++++++++------------------ docs/changes.txt | 1 + 4 files changed, 73 insertions(+), 72 deletions(-) diff --git a/OpenCL/m00600_a0-optimized.cl b/OpenCL/m00600_a0-optimized.cl index 588f2fcaa..6b7009f7d 100644 --- a/OpenCL/m00600_a0-optimized.cl +++ b/OpenCL/m00600_a0-optimized.cl @@ -148,19 +148,19 @@ __kernel void m00600_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru u64 tmp_t[2]; u64 tmp_f[2]; - tmp_h[0] = esalt_bufs->h[0]; - tmp_h[1] = esalt_bufs->h[1]; - tmp_h[2] = esalt_bufs->h[2]; - tmp_h[3] = esalt_bufs->h[3]; - tmp_h[4] = esalt_bufs->h[4]; - tmp_h[5] = esalt_bufs->h[5]; - tmp_h[6] = esalt_bufs->h[6]; - tmp_h[7] = esalt_bufs->h[7]; + tmp_h[0] = esalt_bufs[digests_offset].h[0]; + tmp_h[1] = esalt_bufs[digests_offset].h[1]; + tmp_h[2] = esalt_bufs[digests_offset].h[2]; + tmp_h[3] = esalt_bufs[digests_offset].h[3]; + tmp_h[4] = esalt_bufs[digests_offset].h[4]; + tmp_h[5] = esalt_bufs[digests_offset].h[5]; + tmp_h[6] = esalt_bufs[digests_offset].h[6]; + tmp_h[7] = esalt_bufs[digests_offset].h[7]; - tmp_t[0] = esalt_bufs->t[0]; - tmp_t[1] = esalt_bufs->t[1]; - tmp_f[0] = esalt_bufs->f[0]; - tmp_f[1] = esalt_bufs->f[1]; + tmp_t[0] = esalt_bufs[digests_offset].t[0]; + tmp_t[1] = esalt_bufs[digests_offset].t[1]; + tmp_f[0] = esalt_bufs[digests_offset].f[0]; + tmp_f[1] = esalt_bufs[digests_offset].f[1]; /** * loop @@ -256,19 +256,19 @@ __kernel void m00600_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru u64 tmp_t[2]; u64 tmp_f[2]; - tmp_h[0] = esalt_bufs->h[0]; - tmp_h[1] = esalt_bufs->h[1]; - tmp_h[2] = esalt_bufs->h[2]; - tmp_h[3] = esalt_bufs->h[3]; - tmp_h[4] = esalt_bufs->h[4]; - tmp_h[5] = esalt_bufs->h[5]; - tmp_h[6] = esalt_bufs->h[6]; - tmp_h[7] = esalt_bufs->h[7]; + tmp_h[0] = esalt_bufs[digests_offset].h[0]; + tmp_h[1] = esalt_bufs[digests_offset].h[1]; + tmp_h[2] = esalt_bufs[digests_offset].h[2]; + tmp_h[3] = esalt_bufs[digests_offset].h[3]; + tmp_h[4] = esalt_bufs[digests_offset].h[4]; + tmp_h[5] = esalt_bufs[digests_offset].h[5]; + tmp_h[6] = esalt_bufs[digests_offset].h[6]; + tmp_h[7] = esalt_bufs[digests_offset].h[7]; - tmp_t[0] = esalt_bufs->t[0]; - tmp_t[1] = esalt_bufs->t[1]; - tmp_f[0] = esalt_bufs->f[0]; - tmp_f[1] = esalt_bufs->f[1]; + tmp_t[0] = esalt_bufs[digests_offset].t[0]; + tmp_t[1] = esalt_bufs[digests_offset].t[1]; + tmp_f[0] = esalt_bufs[digests_offset].f[0]; + tmp_f[1] = esalt_bufs[digests_offset].f[1]; /** * digest diff --git a/OpenCL/m00600_a1-optimized.cl b/OpenCL/m00600_a1-optimized.cl index 09b12ac42..829093288 100644 --- a/OpenCL/m00600_a1-optimized.cl +++ b/OpenCL/m00600_a1-optimized.cl @@ -148,19 +148,19 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule u64 tmp_t[2]; u64 tmp_f[2]; - tmp_h[0] = esalt_bufs->h[0]; - tmp_h[1] = esalt_bufs->h[1]; - tmp_h[2] = esalt_bufs->h[2]; - tmp_h[3] = esalt_bufs->h[3]; - tmp_h[4] = esalt_bufs->h[4]; - tmp_h[5] = esalt_bufs->h[5]; - tmp_h[6] = esalt_bufs->h[6]; - tmp_h[7] = esalt_bufs->h[7]; + tmp_h[0] = esalt_bufs[digests_offset].h[0]; + tmp_h[1] = esalt_bufs[digests_offset].h[1]; + tmp_h[2] = esalt_bufs[digests_offset].h[2]; + tmp_h[3] = esalt_bufs[digests_offset].h[3]; + tmp_h[4] = esalt_bufs[digests_offset].h[4]; + tmp_h[5] = esalt_bufs[digests_offset].h[5]; + tmp_h[6] = esalt_bufs[digests_offset].h[6]; + tmp_h[7] = esalt_bufs[digests_offset].h[7]; - tmp_t[0] = esalt_bufs->t[0]; - tmp_t[1] = esalt_bufs->t[1]; - tmp_f[0] = esalt_bufs->f[0]; - tmp_f[1] = esalt_bufs->f[1]; + tmp_t[0] = esalt_bufs[digests_offset].t[0]; + tmp_t[1] = esalt_bufs[digests_offset].t[1]; + tmp_f[0] = esalt_bufs[digests_offset].f[0]; + tmp_f[1] = esalt_bufs[digests_offset].f[1]; /** * loop @@ -315,19 +315,19 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule u64 tmp_t[2]; u64 tmp_f[2]; - tmp_h[0] = esalt_bufs->h[0]; - tmp_h[1] = esalt_bufs->h[1]; - tmp_h[2] = esalt_bufs->h[2]; - tmp_h[3] = esalt_bufs->h[3]; - tmp_h[4] = esalt_bufs->h[4]; - tmp_h[5] = esalt_bufs->h[5]; - tmp_h[6] = esalt_bufs->h[6]; - tmp_h[7] = esalt_bufs->h[7]; + tmp_h[0] = esalt_bufs[digests_offset].h[0]; + tmp_h[1] = esalt_bufs[digests_offset].h[1]; + tmp_h[2] = esalt_bufs[digests_offset].h[2]; + tmp_h[3] = esalt_bufs[digests_offset].h[3]; + tmp_h[4] = esalt_bufs[digests_offset].h[4]; + tmp_h[5] = esalt_bufs[digests_offset].h[5]; + tmp_h[6] = esalt_bufs[digests_offset].h[6]; + tmp_h[7] = esalt_bufs[digests_offset].h[7]; - tmp_t[0] = esalt_bufs->t[0]; - tmp_t[1] = esalt_bufs->t[1]; - tmp_f[0] = esalt_bufs->f[0]; - tmp_f[1] = esalt_bufs->f[1]; + tmp_t[0] = esalt_bufs[digests_offset].t[0]; + tmp_t[1] = esalt_bufs[digests_offset].t[1]; + tmp_f[0] = esalt_bufs[digests_offset].f[0]; + tmp_f[1] = esalt_bufs[digests_offset].f[1]; /** * digest diff --git a/OpenCL/m00600_a3-optimized.cl b/OpenCL/m00600_a3-optimized.cl index e172d1a77..ca4da7144 100644 --- a/OpenCL/m00600_a3-optimized.cl +++ b/OpenCL/m00600_a3-optimized.cl @@ -132,19 +132,19 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule u64 tmp_t[2]; u64 tmp_f[2]; - tmp_h[0] = esalt_bufs->h[0]; - tmp_h[1] = esalt_bufs->h[1]; - tmp_h[2] = esalt_bufs->h[2]; - tmp_h[3] = esalt_bufs->h[3]; - tmp_h[4] = esalt_bufs->h[4]; - tmp_h[5] = esalt_bufs->h[5]; - tmp_h[6] = esalt_bufs->h[6]; - tmp_h[7] = esalt_bufs->h[7]; + tmp_h[0] = esalt_bufs[digests_offset].h[0]; + tmp_h[1] = esalt_bufs[digests_offset].h[1]; + tmp_h[2] = esalt_bufs[digests_offset].h[2]; + tmp_h[3] = esalt_bufs[digests_offset].h[3]; + tmp_h[4] = esalt_bufs[digests_offset].h[4]; + tmp_h[5] = esalt_bufs[digests_offset].h[5]; + tmp_h[6] = esalt_bufs[digests_offset].h[6]; + tmp_h[7] = esalt_bufs[digests_offset].h[7]; - tmp_t[0] = esalt_bufs->t[0]; - tmp_t[1] = esalt_bufs->t[1]; - tmp_f[0] = esalt_bufs->f[0]; - tmp_f[1] = esalt_bufs->f[1]; + tmp_t[0] = esalt_bufs[digests_offset].t[0]; + tmp_t[1] = esalt_bufs[digests_offset].t[1]; + tmp_f[0] = esalt_bufs[digests_offset].f[0]; + tmp_f[1] = esalt_bufs[digests_offset].f[1]; /** * loop @@ -244,19 +244,19 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule u64 tmp_t[2]; u64 tmp_f[2]; - tmp_h[0] = esalt_bufs->h[0]; - tmp_h[1] = esalt_bufs->h[1]; - tmp_h[2] = esalt_bufs->h[2]; - tmp_h[3] = esalt_bufs->h[3]; - tmp_h[4] = esalt_bufs->h[4]; - tmp_h[5] = esalt_bufs->h[5]; - tmp_h[6] = esalt_bufs->h[6]; - tmp_h[7] = esalt_bufs->h[7]; + tmp_h[0] = esalt_bufs[digests_offset].h[0]; + tmp_h[1] = esalt_bufs[digests_offset].h[1]; + tmp_h[2] = esalt_bufs[digests_offset].h[2]; + tmp_h[3] = esalt_bufs[digests_offset].h[3]; + tmp_h[4] = esalt_bufs[digests_offset].h[4]; + tmp_h[5] = esalt_bufs[digests_offset].h[5]; + tmp_h[6] = esalt_bufs[digests_offset].h[6]; + tmp_h[7] = esalt_bufs[digests_offset].h[7]; - tmp_t[0] = esalt_bufs->t[0]; - tmp_t[1] = esalt_bufs->t[1]; - tmp_f[0] = esalt_bufs->f[0]; - tmp_f[1] = esalt_bufs->f[1]; + tmp_t[0] = esalt_bufs[digests_offset].t[0]; + tmp_t[1] = esalt_bufs[digests_offset].t[1]; + tmp_f[0] = esalt_bufs[digests_offset].f[0]; + tmp_f[1] = esalt_bufs[digests_offset].f[1]; /** * digest diff --git a/docs/changes.txt b/docs/changes.txt index 8d6d74031..9ba7aab22 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -689,6 +689,7 @@ It combines all features of all hashcat projects in one project. - Fixed a bug in implementation of DCC2: forced default iteration count for hashes to 10240 - Fixed a bug in implementation of WPA/WPA2: MAC and nonce stay one their original position as in the hccap file - Fixed a bug in implementation of GOST R 34.11-94: zero length passwords were not cracked +- Fixed a bug in implementation of BLAKE2-512 kernels: incorrect access of the esalt buffer ## ## Technical