1
mirror of https://github.com/hashcat/hashcat synced 2024-11-24 14:27:14 +01:00

Fixed thread count maximum for pure kernels in straight attack mode

This commit is contained in:
Jens Steube 2018-11-20 15:29:24 +01:00
parent 2a6444c05a
commit c9da60c73a
3 changed files with 29 additions and 11 deletions

View File

@ -59,17 +59,19 @@
#define RP_PASSWORD_SIZE 256
#ifdef IS_NV
#define COPY_PW(x) \
pw_t pw = (x);
#ifdef REAL_SHM
#define COPY_PW(x) \
__local pw_t s_pws[64]; \
s_pws[get_local_id(0)] = (x); \
s_pws[get_local_id(0)].pw_len &= 255;
#else
#define COPY_PW(x) \
__local pw_t s_pws[64]; \
s_pws[get_local_id(0)] = (x);
#define COPY_PW(x) \
pw_t pw = (x); \
pw.pw_len &= 255;
#endif
#ifdef IS_NV
#define PASTE_PW pw;
#else
#ifdef REAL_SHM
#define PASTE_PW s_pws[get_local_id(0)];
#else
#define PASTE_PW pw;
#endif

View File

@ -48,8 +48,9 @@
- Fixed missing call to WSACleanup() to cleanly shutdown windows sockets system
- Fixed missing call to WSAStartup() and client indexing in order to start the brain server on Windows
- Fixed out-of-bounds write in short-term memory of the brain server
- Fixed status output of progress value when -S and -l are used in combination
- Fixed output of --speed-only and --progress-only when fast hashes are used in combination with --slow-candidates
- Fixed status output of progress value when -S and -l are used in combination
- Fixed thread count maximum for pure kernels in straight attack mode
##
## Technical

View File

@ -6482,7 +6482,22 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
* now everything that depends on threads and accel, basically dynamic workload
*/
const u32 kernel_threads = hashconfig_get_kernel_threads (hashcat_ctx, device_param);
u32 kernel_threads = hashconfig_get_kernel_threads (hashcat_ctx, device_param);
// this is required because inside the kernels there is this:
// __local pw_t s_pws[64];
if (user_options->attack_mode == ATTACK_MODE_STRAIGHT)
{
if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
{
// not required
}
else
{
kernel_threads = MIN (kernel_threads, 64);
}
}
device_param->kernel_threads = kernel_threads;