arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-23 07:08:19 +00:00
Code Issues Releases Wiki Activity

Simplify kernel thread calculation for SCRYPT based algorithms

Browse Source
This commit is contained in:
Jens Steube 2017-08-29 22:43:26 +02:00
parent 12295dcd90
commit 8fb583f0a7
2 changed files with 20 additions and 53 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

43
src/interface.c
View File

@ -24899,7 +24899,9 @@ u32 hashconfig_forced_kernel_threads (hashcat_ctx_t *hashcat_ctx)
if (hashconfig->hash_mode == 3200) kernel_threads = 8; // Blowfish
if (hashconfig->hash_mode == 7500) kernel_threads = 64; // RC4
if (hashconfig->hash_mode == 8500) kernel_threads = 64; // DES
if (hashconfig->hash_mode == 8900) kernel_threads = 16; // SCRYPT
if (hashconfig->hash_mode == 9000) kernel_threads = 8; // Blowfish
if (hashconfig->hash_mode == 9300) kernel_threads = 8; // SCRYPT
if (hashconfig->hash_mode == 9700) kernel_threads = 64; // RC4
if (hashconfig->hash_mode == 9710) kernel_threads = 64; // RC4
if (hashconfig->hash_mode == 9800) kernel_threads = 64; // RC4
@ -24910,44 +24912,37 @@ u32 hashconfig_forced_kernel_threads (hashcat_ctx_t *hashcat_ctx)
if (hashconfig->hash_mode == 13100) kernel_threads = 64; // RC4
if (hashconfig->hash_mode == 14000) kernel_threads = 64; // DES
if (hashconfig->hash_mode == 14100) kernel_threads = 64; // DES
if (hashconfig->hash_mode == 15700) kernel_threads = 1; // SCRYPT
return kernel_threads;
}
u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param_t *device_param)
{
const hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
// kernel_threads is 256 at this point
u32 kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size);
u32 kernel_threads = KERNEL_THREADS_MAX;
if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
kernel_threads = 64;
}
if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300) || (hashconfig->hash_mode == 15700))
{
const hashes_t *hashes = hashcat_ctx->hashes;
const u32 scrypt_r = hashes->salts_buf[0].scrypt_r;
const u32 scrypt_p = hashes->salts_buf[0].scrypt_p;
const u32 scrypt_l = scrypt_r * scrypt_p;
if (scrypt_l)
{
kernel_threads = 256 / scrypt_l;
}
else
{
kernel_threads = 256;
}
}
// a CPU doesn't have "opencl threads", this sets it to 1
if (device_param->device_type & CL_DEVICE_TYPE_CPU)
{
kernel_threads = KERNEL_THREADS_MAX_CPU;
}
// but 64 for AMD GPU, they prefer it
if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
kernel_threads = MIN (kernel_threads, 64);
}
// a kernel can force an even lower value
kernel_threads = MIN (kernel_threads, device_param->device_maxworkgroup_size);
// a kernel can also force a fixed value
const u32 forced_kernel_threads = hashconfig_forced_kernel_threads (hashcat_ctx);
if (forced_kernel_threads) kernel_threads = forced_kernel_threads;

30
src/opencl.c
View File

@ -3976,7 +3976,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
hashconfig->tmp_size = scrypt_tmp_size;
u32 tmto_start = 0;
u32 tmto_start = 1;
u32 tmto_stop = 6;
if (user_options->scrypt_tmto)
@ -3984,34 +3984,6 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
tmto_start = user_options->scrypt_tmto;
tmto_stop = user_options->scrypt_tmto;
}
else
{
// in case the user did not specify the tmto manually
// use some values known to run best (tested on 290x for AMD and GTX1080 for NV)
if (hashconfig->hash_mode == 8900)
{
if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
tmto_start = 4;
}
else if (device_param->device_vendor_id == VENDOR_ID_NV)
{
tmto_start = 3;
}
}
else if (hashconfig->hash_mode == 9300)
{
if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
tmto_start = 5;
}
else if (device_param->device_vendor_id == VENDOR_ID_NV)
{
tmto_start = 3;
}
}
}
const u32 kernel_power_max = device_param->hardware_power * device_param->kernel_accel_max;