projects
/
hashcat.git
/ blobdiff
commit
grep
author
committer
pickaxe
?
search:
re
summary
|
shortlog
|
log
|
commit
|
commitdiff
|
tree
raw
|
inline
| side by side
Workaround memory allocation limit from OpenCL by using multiple buffers for scrypt
[hashcat.git]
/
include
/
types.h
diff --git
a/include/types.h
b/include/types.h
index
3306aee
..
02df8ea
100644
(file)
--- a/
include/types.h
+++ b/
include/types.h
@@
-6,6
+6,12
@@
#ifndef TYPES_H
#define TYPES_H
#ifndef TYPES_H
#define TYPES_H
+#ifdef _WIN
+#define EOL "\r\n"
+#else
+#define EOL "\n"
+#endif
+
typedef struct
{
uint salt_buf[16];
typedef struct
{
uint salt_buf[16];
@@
-26,8
+32,6
@@
typedef struct
uint scrypt_N;
uint scrypt_r;
uint scrypt_p;
uint scrypt_N;
uint scrypt_r;
uint scrypt_p;
- uint scrypt_tmto;
- uint scrypt_phy;
} salt_t;
} salt_t;
@@
-156,7
+160,7
@@
typedef struct
/* specific to version 1 */
u32 contents_len;
/* specific to version 1 */
u32 contents_len;
- u32 contents[12500];
+ u32 contents[75000];
/* specific to version 2 */
u32 expected_bytes[8];
/* specific to version 2 */
u32 expected_bytes[8];
@@
-168,6
+172,7
@@
typedef struct
uint salt_buf[16];
uint data_buf[112];
uint keyfile_buf[16];
uint salt_buf[16];
uint data_buf[112];
uint keyfile_buf[16];
+ uint signature;
} tc_t;
} tc_t;
@@
-258,9
+263,35
@@
typedef struct
typedef struct
{
typedef struct
{
- uint P[256];
+ u32 salt_buf[128];
+ u32 salt_len;
+
+ u32 pc_digest[5];
+ u32 pc_offset;
+
+} pstoken_t;
+
+typedef struct
+{
+ u32 type;
+ u32 mode;
+ u32 magic;
+ u32 salt_len;
+ u32 salt_buf[4];
+ u32 verify_bytes;
+ u32 compress_length;
+ u32 data_len;
+ u32 data_buf[2048];
+ u32 auth_len;
+ u32 auth_buf[4];
-} scrypt_tmp_t;
+} zip2_t;
+
+typedef struct
+{
+ uint salt_buf[32];
+
+} win8phone_t;
typedef struct
{
typedef struct
{
@@
-616,8
+647,7
@@
typedef struct
typedef struct
{
typedef struct
{
- u32 KEK[5];
-
+ u32 KEK[4];
u32 lsb[4];
u32 cipher[4];
u32 lsb[4];
u32 cipher[4];
@@
-832,6
+862,9
@@
typedef struct
typedef struct
{
typedef struct
{
+ uint salt_pos;
+ uint digest_pos;
+ uint hash_pos;
uint gidvid;
uint il_pos;
uint gidvid;
uint il_pos;
@@
-884,7
+917,7
@@
typedef struct
#define PW_MAX1 (PW_MAX + 1)
#define PW_DICTMAX 31
#define PW_DICTMAX1 (PW_DICTMAX + 1)
#define PW_MAX1 (PW_MAX + 1)
#define PW_DICTMAX 31
#define PW_DICTMAX1 (PW_DICTMAX + 1)
-#define PARAMCNT 32
+#define PARAMCNT 64
struct __hc_device_param
{
struct __hc_device_param
{
@@
-901,10
+934,10
@@
struct __hc_device_param
uint kernel_exec_timeout;
uint device_processors;
uint kernel_exec_timeout;
uint device_processors;
- uint device_processor_cores;
u64 device_maxmem_alloc;
u64 device_global_mem;
u32 device_maxclock_frequency;
u64 device_maxmem_alloc;
u64 device_global_mem;
u32 device_maxclock_frequency;
+ size_t device_maxworkgroup_size;
uint vector_width;
uint vector_width;
@@
-916,22
+949,22
@@
struct __hc_device_param
uint kernel_accel_min;
uint kernel_accel_max;
uint kernel_power;
uint kernel_accel_min;
uint kernel_accel_max;
uint kernel_power;
- uint kernel_power_user;
-
- uint size_pws;
- uint size_tmps;
- uint size_hooks;
- uint size_bfs;
- uint size_combs;
- uint size_rules;
- uint size_rules_c;
- uint size_root_css;
- uint size_markov_css;
- uint size_digests;
- uint size_salts;
- uint size_shown;
- uint size_results;
- uint size_plains;
+ uint hardware_power;
+
+ size_t size_pws;
+ size_t size_tmps;
+ size_t size_hooks;
+ size_t size_bfs;
+ size_t size_combs;
+ size_t size_rules;
+ size_t size_rules_c;
+ size_t size_root_css;
+ size_t size_markov_css;
+ size_t size_digests;
+ size_t size_salts;
+ size_t size_shown;
+ size_t size_results;
+ size_t size_plains;
FILE *combs_fp;
comb_t *combs_buf;
FILE *combs_fp;
comb_t *combs_buf;
@@
-944,8
+977,6
@@
struct __hc_device_param
u64 words_off;
u64 words_done;
u64 words_off;
u64 words_done;
- uint *result;
-
uint outerloop_pos;
uint outerloop_left;
uint outerloop_pos;
uint outerloop_left;
@@
-955,24
+986,34
@@
struct __hc_device_param
uint exec_pos;
double exec_ms[EXEC_CACHE];
uint exec_pos;
double exec_ms[EXEC_CACHE];
+ // workaround cpu spinning
+
+ double exec_us_prev1[EXPECTED_ITERATIONS];
+ double exec_us_prev2[EXPECTED_ITERATIONS];
+ double exec_us_prev3[EXPECTED_ITERATIONS];
+
+ // this is "current" speed
+
uint speed_pos;
u64 speed_cnt[SPEED_CACHE];
uint speed_pos;
u64 speed_cnt[SPEED_CACHE];
- float speed_ms[SPEED_CACHE];
-
- hc_timer_t speed_rec[SPEED_CACHE];
+ double speed_ms[SPEED_CACHE];
hc_timer_t timer_speed;
// device specific attributes starting
char *device_name;
hc_timer_t timer_speed;
// device specific attributes starting
char *device_name;
+ char *device_vendor;
char *device_name_chksum;
char *device_version;
char *driver_version;
bool opencl_v12;
char *device_name_chksum;
char *device_version;
char *driver_version;
bool opencl_v12;
- cl_uint vendor_id;
+ double nvidia_spin_damp;
+
+ cl_uint device_vendor_id;
+ cl_uint platform_vendor_id;
cl_kernel kernel1;
cl_kernel kernel12;
cl_kernel kernel1;
cl_kernel kernel12;
@@
-985,13
+1026,13
@@
struct __hc_device_param
cl_kernel kernel_amp;
cl_kernel kernel_tm;
cl_kernel kernel_weak;
cl_kernel kernel_amp;
cl_kernel kernel_tm;
cl_kernel kernel_weak;
+ cl_kernel kernel_memset;
cl_context context;
cl_program program;
cl_program program_mp;
cl_program program_amp;
cl_context context;
cl_program program;
cl_program program_mp;
cl_program program_amp;
- cl_program program_weak;
cl_command_queue command_queue;
cl_command_queue command_queue;
@@
-1023,7
+1064,10
@@
struct __hc_device_param
cl_mem d_tmps;
cl_mem d_hooks;
cl_mem d_result;
cl_mem d_tmps;
cl_mem d_hooks;
cl_mem d_result;
- cl_mem d_scryptV_buf;
+ cl_mem d_scryptV0_buf;
+ cl_mem d_scryptV1_buf;
+ cl_mem d_scryptV2_buf;
+ cl_mem d_scryptV3_buf;
cl_mem d_root_css_buf;
cl_mem d_markov_css_buf;
cl_mem d_root_css_buf;
cl_mem d_markov_css_buf;
@@
-1033,6
+1077,7
@@
struct __hc_device_param
void *kernel_params_mp_l[PARAMCNT];
void *kernel_params_amp[PARAMCNT];
void *kernel_params_tm[PARAMCNT];
void *kernel_params_mp_l[PARAMCNT];
void *kernel_params_amp[PARAMCNT];
void *kernel_params_tm[PARAMCNT];
+ void *kernel_params_memset[PARAMCNT];
u32 kernel_params_buf32[PARAMCNT];
u32 kernel_params_buf32[PARAMCNT];
@@
-1046,6
+1091,7
@@
struct __hc_device_param
u64 kernel_params_mp_l_buf64[PARAMCNT];
u32 kernel_params_amp_buf32[PARAMCNT];
u64 kernel_params_mp_l_buf64[PARAMCNT];
u32 kernel_params_amp_buf32[PARAMCNT];
+ u32 kernel_params_memset_buf32[PARAMCNT];
};
typedef struct __hc_device_param hc_device_param_t;
};
typedef struct __hc_device_param hc_device_param_t;
@@
-1053,23
+1099,15
@@
typedef struct __hc_device_param hc_device_param_t;
#ifdef HAVE_HWMON
typedef struct
{
#ifdef HAVE_HWMON
typedef struct
{
- union
- {
- #ifdef HAVE_ADL
- HM_ADAPTER_AMD amd;
- #endif
-
- #if defined(HAVE_NVML) || defined(HAVE_NVAPI)
- HM_ADAPTER_NV nv;
- #endif
-
- } adapter_index;
+ HM_ADAPTER_ADL adl;
+ HM_ADAPTER_NVML nvml;
+ HM_ADAPTER_NVAPI nvapi;
+ HM_ADAPTER_XNVCTRL xnvctrl;
- int od_version;
- int fan_supported;
+ int od_version;
- // int busid; // used for CL_DEVICE_TOPOLOGY_AMD but broken for dual GPUs
- // int devid; // used for CL_DEVICE_TOPOLOGY_AMD but broken for dual GPUs
+ int fan_get_supported;
+ int fan_set_supported;
} hm_attrs_t;
#endif // HAVE_HWMON
} hm_attrs_t;
#endif // HAVE_HWMON
@@
-1086,12
+1124,16
@@
typedef struct
hc_device_param_t *devices_param;
hc_device_param_t *devices_param;
+ uint shutdown_inner;
+ uint shutdown_outer;
+
/**
* workload specific
*/
/**
* workload specific
*/
+ uint hardware_power_all;
uint kernel_power_all;
uint kernel_power_all;
- float kernel_power_div;
+ u64 kernel_power_final; // we save that so that all divisions are done from the same base
/**
* attack specific
/**
* attack specific
@@
-1136,8
+1178,10
@@
typedef struct
*/
#ifdef HAVE_HWMON
*/
#ifdef HAVE_HWMON
- void *hm_nv;
- void *hm_amd;
+ void *hm_adl;
+ void *hm_nvml;
+ void *hm_nvapi;
+ void *hm_xnvctrl;
hm_attrs_t hm_device[DEVICES_MAX];
#endif
hm_attrs_t hm_device[DEVICES_MAX];
#endif
@@
-1161,6
+1205,9
@@
typedef struct
void *esalts_buf;
void *esalts_buf;
+ uint scrypt_tmp_size;
+ uint scrypt_tmto_final;
+
/**
* logging
*/
/**
* logging
*/
@@
-1211,11
+1258,10
@@
typedef struct
uint restore_disable;
uint status;
uint status_timer;
uint restore_disable;
uint status;
uint status_timer;
- uint status_automat;
+ uint machine_readable;
uint quiet;
uint force;
uint benchmark;
uint quiet;
uint force;
uint benchmark;
- uint benchmark_repeats;
uint runtime;
uint remove;
uint remove_timer;
uint runtime;
uint remove;
uint remove_timer;
@@
-1230,6
+1276,8
@@
typedef struct
uint scrypt_tmto;
uint segment_size;
char *truecrypt_keyfiles;
uint scrypt_tmto;
uint segment_size;
char *truecrypt_keyfiles;
+ char *veracrypt_keyfiles;
+ uint veracrypt_pim;
uint workload_profile;
uint hash_mode;
uint workload_profile;
uint hash_mode;
@@
-1291,7
+1339,7
@@
typedef struct
hc_timer_t timer_running; // timer on current dict
hc_timer_t timer_paused; // timer on current dict
hc_timer_t timer_running; // timer on current dict
hc_timer_t timer_paused; // timer on current dict
- float ms_paused; // timer on current dict
+ double ms_paused; // timer on current dict
/**
* hash_info and username
/**
* hash_info and username
@@
-1309,3
+1357,4
@@
typedef struct
extern hc_global_data_t data;
#endif
extern hc_global_data_t data;
#endif
+