From 58d1dedd1e27aa3be5e0c6148b21099601253cf6 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 19 Aug 2017 19:25:37 +0200 Subject: [PATCH] OpenCL Buffers: Do not allocate memory for amplifiers for fast hashes, it's simply not needed --- docs/changes.txt | 1 + include/types.h | 1 + src/opencl.c | 7 +++++-- 3 files changed, 7 insertions(+), 2 deletions(-) diff --git a/docs/changes.txt b/docs/changes.txt index aafc642a2..edee965e8 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -38,6 +38,7 @@ ## - Charset: Add additional DES charsets with corrected parity +- OpenCL Buffers: Do not allocate memory for amplifiers for fast hashes, it's simply not needed - OpenCL Kernels: Improved performance of SHA-3 Kernel (keccak) by hardcoding the 0x80 stopbit - OpenCL Kernels: Move from ld.global.v4.u32 to ld.const.v4.u32 in _a3 kernels - OpenCL Kernels: Replace bitwise swaps with rotate() versions for AMD diff --git a/include/types.h b/include/types.h index 0191095ec..f66ca7f36 100644 --- a/include/types.h +++ b/include/types.h @@ -962,6 +962,7 @@ typedef struct hc_device_param u32 hardware_power; size_t size_pws; + size_t size_pws_amp; size_t size_tmps; size_t size_hooks; size_t size_bfs; diff --git a/src/opencl.c b/src/opencl.c index 052485133..ef7354570 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -4158,11 +4158,14 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) } */ + const size_t size_pws_amp = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) ? size_pws : size_pws; + device_param->size_bfs = size_bfs; device_param->size_combs = size_combs; device_param->size_rules = size_rules; device_param->size_rules_c = size_rules_c; device_param->size_pws = size_pws; + device_param->size_pws_amp = size_pws_amp; device_param->size_tmps = size_tmps; device_param->size_hooks = size_hooks; @@ -4785,7 +4788,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) */ CL_rc = hc_clCreateBuffer (hashcat_ctx, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_buf); if (CL_rc == -1) return -1; - CL_rc = hc_clCreateBuffer (hashcat_ctx, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_amp_buf); if (CL_rc == -1) return -1; + CL_rc = hc_clCreateBuffer (hashcat_ctx, device_param->context, CL_MEM_READ_ONLY, size_pws_amp, NULL, &device_param->d_pws_amp_buf); if (CL_rc == -1) return -1; CL_rc = hc_clCreateBuffer (hashcat_ctx, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL, &device_param->d_tmps); if (CL_rc == -1) return -1; CL_rc = hc_clCreateBuffer (hashcat_ctx, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL, &device_param->d_hooks); if (CL_rc == -1) return -1; CL_rc = hc_clCreateBuffer (hashcat_ctx, device_param->context, CL_MEM_READ_ONLY, bitmap_ctx->bitmap_size, NULL, &device_param->d_bitmap_s1_a); if (CL_rc == -1) return -1; @@ -5374,7 +5377,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) // zero some data buffers CL_rc = run_kernel_bzero (hashcat_ctx, device_param, device_param->d_pws_buf, device_param->size_pws); if (CL_rc == -1) return -1; - CL_rc = run_kernel_bzero (hashcat_ctx, device_param, device_param->d_pws_amp_buf, device_param->size_pws); if (CL_rc == -1) return -1; + CL_rc = run_kernel_bzero (hashcat_ctx, device_param, device_param->d_pws_amp_buf, device_param->size_pws_amp); if (CL_rc == -1) return -1; CL_rc = run_kernel_bzero (hashcat_ctx, device_param, device_param->d_tmps, device_param->size_tmps); if (CL_rc == -1) return -1; CL_rc = run_kernel_bzero (hashcat_ctx, device_param, device_param->d_hooks, device_param->size_hooks); if (CL_rc == -1) return -1; CL_rc = run_kernel_bzero (hashcat_ctx, device_param, device_param->d_plain_bufs, device_param->size_plains); if (CL_rc == -1) return -1;