OpenCL Buffers: Do not allocate memory for amplifiers for fast hashes, it's simply not needed

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

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;

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;