arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-22 08:08:10 +00:00
Code Issues Releases Wiki Activity

Prepare native CUDA hybrid integration

Browse Source
This commit is contained in:
Jens Steube 2019-04-25 14:45:17 +02:00
parent c02083281f
commit 4b986de5fb
29 changed files with 1825 additions and 510 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

31
OpenCL/inc_common.cl
View File

@ -5,6 +5,7 @@
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.h"
#include "inc_common.h"
/**
@ -1415,8 +1416,8 @@ DECLSPEC int is_valid_hex_8 (const u8 v)
{
// direct lookup table is slower thanks to CMOV
if ((v >= '0') && (v <= '9')) return 1;
if ((v >= 'a') && (v <= 'f')) return 1;
if ((v >= (u8) '0') && (v <= (u8) '9')) return 1;
if ((v >= (u8) 'a') && (v <= (u8) 'f')) return 1;
return 0;
}
@ -1433,10 +1434,10 @@ DECLSPEC int is_valid_hex_32 (const u32 v)
DECLSPEC int is_valid_base58_8 (const u8 v)
{
if (v > 'z') return 0;
if (v < '1') return 0;
if ((v > '9') && (v < 'A')) return 0;
if ((v > 'Z') && (v < 'a')) return 0;
if (v > (u8) 'z') return 0;
if (v < (u8) '1') return 0;
if ((v > (u8) '9') && (v < (u8) 'A')) return 0;
if ((v > (u8) 'Z') && (v < (u8) 'a')) return 0;
return 1;
}
@ -60860,7 +60861,23 @@ KERNEL_FQ void gpu_memset (GLOBAL_AS uint4 *buf, const u32 value, const u64 gid_
if (gid >= gid_max) return;
buf[gid] = (uint4) (value);
uint4 r;
#if defined IS_NATIVE
r = value;
#elif defined IS_OPENCL
r.s0 = value;
r.s1 = value;
r.s2 = value;
r.s3 = value;
#elif defined IS_CUDA
r.x = value;
r.y = value;
r.z = value;
r.w = value;
#endif
buf[gid] = r;
}
KERNEL_FQ void gpu_atinit (GLOBAL_AS pw_t *buf, const u64 gid_max)

44
OpenCL/inc_platform.cl Normal file
View File

@ -0,0 +1,44 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.h"
#ifdef IS_NATIVE
#endif
#ifdef IS_CUDA
DECLSPEC u32 atomic_dec (u32 *p)
{
return atomicSub (p, 1);
}
DECLSPEC u32 atomic_inc (u32 *p)
{
return atomicAdd (p, 1);
}
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
{
return blockDim.x * blockIdx.x + threadIdx.x;
}
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)))
{
return threadIdx.x;
}
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
{
// verify
return blockDim.x;
}
#endif
#ifdef IS_OPENCL
#endif

24
OpenCL/inc_platform.h Normal file
View File

@ -0,0 +1,24 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#ifndef _INC_PLATFORM_H
#ifdef IS_CUDA
DECLSPEC u32 atomic_dec (u32 *p);
DECLSPEC u32 atomic_inc (u32 *p);
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)));
DECLSPEC uint4 uint4_init (const u32 a);
DECLSPEC uint4 uint4_init (const u32 a, const u32 b, const u32 c, const u32 d);
DECLSPEC __inline__ u8 rotate (const u8 v, const int i);
DECLSPEC __inline__ u32 rotate (const u32 v, const int i);
DECLSPEC __inline__ u64 rotate (const u64 v, const int i);
#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
#define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a))))
#endif
#endif // _INC_PLATFORM_H

4
OpenCL/inc_types.h
View File

@ -6,6 +6,10 @@
#ifndef _INC_TYPES_H
#define _INC_TYPES_H
#ifdef IS_CUDA
typedef unsigned char uchar;
#endif
#ifdef KERNEL_STATIC
typedef uchar u8;
typedef ushort u16;

21
OpenCL/inc_vendor.h
View File

@ -6,12 +6,25 @@
#ifndef _INC_VENDOR_H
#define _INC_VENDOR_H
#ifdef _CPU_OPENCL_EMU_H
#if defined _CPU_OPENCL_EMU_H
#define IS_NATIVE
#elif defined __CUDACC__
#define IS_CUDA
#else
#define IS_OPENCL
#endif
#if defined IS_NATIVE
#define CONSTANT_AS
#define GLOBAL_AS
#define LOCAL_AS
#define KERNEL_FQ
#else
#elif defined IS_CUDA
#define CONSTANT_AS
#define GLOBAL_AS
#define LOCAL_AS
#define KERNEL_FQ __global__
#elif defined IS_OPENCL
#define CONSTANT_AS __constant
#define GLOBAL_AS __global
#define LOCAL_AS __local
@ -90,11 +103,15 @@
#if defined IS_CPU
#define DECLSPEC inline
#elif defined IS_GPU
#if defined IS_CUDA
#define DECLSPEC __device__
#else
#if defined IS_AMD
#define DECLSPEC inline static
#else
#define DECLSPEC
#endif
#endif
#else
#define DECLSPEC
#endif

1
OpenCL/m00000_a3-optimized.cl
View File

@ -8,6 +8,7 @@
#ifdef KERNEL_STATIC
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md5.cl"

4
OpenCL/m02500-pure.cl
View File

@ -679,7 +679,11 @@ KERNEL_FQ void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)
s_te4[i] = te4[i];
}
#ifdef IS_CUDA
__syncthreads();
#else
barrier (CLK_LOCAL_MEM_FENCE);
#endif
#else

39
include/opencl.h → include/backend.h
View File

@ -3,8 +3,8 @@
* License.....: MIT
*/
#ifndef _OPENCL_H
#define _OPENCL_H
#ifndef _BACKEND_H
#define _BACKEND_H
#include <stdio.h>
#include <errno.h>
@ -22,8 +22,11 @@ static const char CL_VENDOR_MESA[] = "Mesa";
static const char CL_VENDOR_NV[] = "NVIDIA Corporation";
static const char CL_VENDOR_POCL[] = "The pocl project";
int ocl_init (hashcat_ctx_t *hashcat_ctx);
void ocl_close (hashcat_ctx_t *hashcat_ctx);
int cuda_init (hashcat_ctx_t *hashcat_ctx);
void cuda_close (hashcat_ctx_t *hashcat_ctx);
int ocl_init (hashcat_ctx_t *hashcat_ctx);
void ocl_close (hashcat_ctx_t *hashcat_ctx);
int hc_clBuildProgram (hashcat_ctx_t *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data);
int hc_clCreateBuffer (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_mem *mem);
@ -82,20 +85,20 @@ void generate_cached_kernel_mp_filename (const u32 opti_type, const u64 opts_ty
void generate_source_kernel_amp_filename (const u32 attack_kern, char *shared_dir, char *source_file);
void generate_cached_kernel_amp_filename (const u32 attack_kern, char *profile_dir, const char *device_name_chksum, char *cached_file);
int opencl_ctx_init (hashcat_ctx_t *hashcat_ctx);
void opencl_ctx_destroy (hashcat_ctx_t *hashcat_ctx);
int backend_ctx_init (hashcat_ctx_t *hashcat_ctx);
void backend_ctx_destroy (hashcat_ctx_t *hashcat_ctx);
int opencl_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime);
void opencl_ctx_devices_destroy (hashcat_ctx_t *hashcat_ctx);
void opencl_ctx_devices_sync_tuning (hashcat_ctx_t *hashcat_ctx);
void opencl_ctx_devices_update_power (hashcat_ctx_t *hashcat_ctx);
void opencl_ctx_devices_kernel_loops (hashcat_ctx_t *hashcat_ctx);
int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime);
void backend_ctx_devices_destroy (hashcat_ctx_t *hashcat_ctx);
void backend_ctx_devices_sync_tuning (hashcat_ctx_t *hashcat_ctx);
void backend_ctx_devices_update_power (hashcat_ctx_t *hashcat_ctx);
void backend_ctx_devices_kernel_loops (hashcat_ctx_t *hashcat_ctx);
int opencl_session_begin (hashcat_ctx_t *hashcat_ctx);
void opencl_session_destroy (hashcat_ctx_t *hashcat_ctx);
void opencl_session_reset (hashcat_ctx_t *hashcat_ctx);
int opencl_session_update_combinator (hashcat_ctx_t *hashcat_ctx);
int opencl_session_update_mp (hashcat_ctx_t *hashcat_ctx);
int opencl_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_l, const u32 css_cnt_r);
int backend_session_begin (hashcat_ctx_t *hashcat_ctx);
void backend_session_destroy (hashcat_ctx_t *hashcat_ctx);
void backend_session_reset (hashcat_ctx_t *hashcat_ctx);
int backend_session_update_combinator (hashcat_ctx_t *hashcat_ctx);
int backend_session_update_mp (hashcat_ctx_t *hashcat_ctx);
int backend_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_l, const u32 css_cnt_r);
#endif // _OPENCL_H
#endif // _BACKEND_H

1042
include/ext_cuda.h Normal file
View File

File diff suppressed because it is too large Load Diff

8
include/types.h
View File

@ -989,6 +989,7 @@ typedef struct link_speed
} link_speed_t;
#include "ext_cuda.h"
#include "ext_OpenCL.h"
typedef struct hc_device_param
@ -1328,11 +1329,12 @@ typedef struct hc_device_param
} hc_device_param_t;
typedef struct opencl_ctx
typedef struct backend_ctx
{
bool enabled;
void *ocl;
void *cuda;
cl_uint platforms_cnt;
cl_platform_id *platforms;
@ -1369,7 +1371,7 @@ typedef struct opencl_ctx
int force_jit_compilation;
} opencl_ctx_t;
} backend_ctx_t;
typedef enum kernel_workload
{
@ -2299,7 +2301,7 @@ typedef struct hashcat_ctx
loopback_ctx_t *loopback_ctx;
mask_ctx_t *mask_ctx;
module_ctx_t *module_ctx;
opencl_ctx_t *opencl_ctx;
backend_ctx_t *backend_ctx;
outcheck_ctx_t *outcheck_ctx;
outfile_ctx_t *outfile_ctx;
pidfile_ctx_t *pidfile_ctx;

4
src/Makefile
View File

@ -274,13 +274,13 @@ endif # MSYS2
## Objects
##
EMU_OBJS_ALL := emu_general emu_inc_common emu_inc_scalar emu_inc_simd
EMU_OBJS_ALL := emu_general emu_inc_common emu_inc_platform emu_inc_scalar emu_inc_simd
EMU_OBJS_ALL += emu_inc_rp emu_inc_rp_optimized
EMU_OBJS_ALL += emu_inc_truecrypt_crc32 emu_inc_truecrypt_keyfile emu_inc_truecrypt_xts emu_inc_veracrypt_xts
EMU_OBJS_ALL += emu_inc_hash_md4 emu_inc_hash_md5 emu_inc_hash_ripemd160 emu_inc_hash_sha1 emu_inc_hash_sha256 emu_inc_hash_sha384 emu_inc_hash_sha512 emu_inc_hash_streebog256 emu_inc_hash_streebog512
EMU_OBJS_ALL += emu_inc_cipher_aes emu_inc_cipher_camellia emu_inc_cipher_des emu_inc_cipher_kuznyechik emu_inc_cipher_serpent emu_inc_cipher_twofish
OBJS_ALL := affinity autotune benchmark bitmap bitops combinator common convert cpt cpu_crc32 debugfile dictstat dispatch dynloader event ext_ADL ext_nvapi ext_nvml ext_OpenCL ext_sysfs ext_lzma filehandling folder hashcat hashes hlfmt hwmon induct interface keyboard_layout locking logfile loopback memory monitor mpsp opencl outfile_check outfile pidfile potfile restore rp rp_cpu selftest slow_candidates shared status stdout straight terminal thread timer tuningdb usage user_options wordlist $(EMU_OBJS_ALL)
OBJS_ALL := affinity autotune backend benchmark bitmap bitops combinator common convert cpt cpu_crc32 debugfile dictstat dispatch dynloader event ext_ADL ext_cuda ext_nvapi ext_nvml ext_OpenCL ext_sysfs ext_lzma filehandling folder hashcat hashes hlfmt hwmon induct interface keyboard_layout locking logfile loopback memory monitor mpsp outfile_check outfile pidfile potfile restore rp rp_cpu selftest slow_candidates shared status stdout straight terminal thread timer tuningdb usage user_options wordlist $(EMU_OBJS_ALL)
ifeq ($(ENABLE_BRAIN),1)
OBJS_ALL += brain

12
src/autotune.c
View File

@ -6,7 +6,7 @@
#include "common.h"
#include "types.h"
#include "event.h"
#include "opencl.h"
#include "backend.h"
#include "status.h"
#include "autotune.h"
@ -50,11 +50,11 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
{
const hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
const straight_ctx_t *straight_ctx = hashcat_ctx->straight_ctx;
const user_options_t *user_options = hashcat_ctx->user_options;
const double target_msec = opencl_ctx->target_msec;
const double target_msec = backend_ctx->target_msec;
const u32 kernel_accel_min = device_param->kernel_accel_min;
const u32 kernel_accel_max = device_param->kernel_accel_max;
@ -283,11 +283,11 @@ HC_API_CALL void *thread_autotune (void *p)
hashcat_ctx_t *hashcat_ctx = thread_param->hashcat_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (opencl_ctx->enabled == false) return NULL;
if (backend_ctx->enabled == false) return NULL;
hc_device_param_t *device_param = opencl_ctx->devices_param + thread_param->tid;
hc_device_param_t *device_param = backend_ctx->devices_param + thread_param->tid;
if (device_param->skipped == true) return NULL;

596
src/opencl.c → src/backend.c
View File

File diff suppressed because it is too large Load Diff

52
src/dispatch.c
View File

@ -7,7 +7,7 @@
#include "types.h"
#include "event.h"
#include "memory.h"
#include "opencl.h"
#include "backend.h"
#include "wordlist.h"
#include "shared.h"
#include "thread.h"
@ -23,13 +23,13 @@
static u64 get_highest_words_done (const hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
u64 words_cur = 0;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;
@ -45,13 +45,13 @@ static u64 get_highest_words_done (const hashcat_ctx_t *hashcat_ctx)
static u64 get_lowest_words_done (const hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
u64 words_cur = 0xffffffffffffffff;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;
@ -76,20 +76,20 @@ static int set_kernel_power_final (hashcat_ctx_t *hashcat_ctx, const u64 kernel_
{
EVENT (EVENT_SET_KERNEL_POWER_FINAL);
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
opencl_ctx->kernel_power_final = kernel_power_final;
backend_ctx->kernel_power_final = kernel_power_final;
return 0;
}
static u64 get_power (opencl_ctx_t *opencl_ctx, hc_device_param_t *device_param)
static u64 get_power (backend_ctx_t *backend_ctx, hc_device_param_t *device_param)
{
const u64 kernel_power_final = opencl_ctx->kernel_power_final;
const u64 kernel_power_final = backend_ctx->kernel_power_final;
if (kernel_power_final)
{
const double device_factor = (double) device_param->hardware_power / opencl_ctx->hardware_power_all;
const double device_factor = (double) device_param->hardware_power / backend_ctx->hardware_power_all;
const u64 words_left_device = (u64) CEIL (kernel_power_final * device_factor);
@ -109,7 +109,7 @@ static u64 get_power (opencl_ctx_t *opencl_ctx, hc_device_param_t *device_param)
static u64 get_work (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 max)
{
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
status_ctx_t *status_ctx = hashcat_ctx->status_ctx;
user_options_t *user_options = hashcat_ctx->user_options;
@ -120,19 +120,19 @@ static u64 get_work (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
device_param->words_off = words_off;
const u64 kernel_power_all = opencl_ctx->kernel_power_all;
const u64 kernel_power_all = backend_ctx->kernel_power_all;
const u64 words_left = words_base - words_off;
if (words_left < kernel_power_all)
{
if (opencl_ctx->kernel_power_final == 0)
if (backend_ctx->kernel_power_final == 0)
{
set_kernel_power_final (hashcat_ctx, words_left);
}
}
const u64 kernel_power = get_power (opencl_ctx, device_param);
const u64 kernel_power = get_power (backend_ctx, device_param);
u64 work = MIN (words_left, kernel_power);
@ -339,11 +339,11 @@ HC_API_CALL void *thread_calc_stdin (void *p)
hashcat_ctx_t *hashcat_ctx = thread_param->hashcat_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (opencl_ctx->enabled == false) return NULL;
if (backend_ctx->enabled == false) return NULL;
hc_device_param_t *device_param = opencl_ctx->devices_param + thread_param->tid;
hc_device_param_t *device_param = backend_ctx->devices_param + thread_param->tid;
if (device_param->skipped) return NULL;
@ -370,7 +370,7 @@ static int calc (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx;
straight_ctx_t *straight_ctx = hashcat_ctx->straight_ctx;
combinator_ctx_t *combinator_ctx = hashcat_ctx->combinator_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
status_ctx_t *status_ctx = hashcat_ctx->status_ctx;
const u32 attack_mode = user_options->attack_mode;
@ -468,7 +468,7 @@ static int calc (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
// this greatly reduces spam on hashcat console
const u64 pre_rejects_ignore = get_power (opencl_ctx, device_param) / 2;
const u64 pre_rejects_ignore = get_power (backend_ctx, device_param) / 2;
while (pre_rejects > pre_rejects_ignore)
{
@ -801,7 +801,7 @@ static int calc (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
// this greatly reduces spam on hashcat console
const u64 pre_rejects_ignore = get_power (opencl_ctx, device_param) / 2;
const u64 pre_rejects_ignore = get_power (backend_ctx, device_param) / 2;
while (pre_rejects > pre_rejects_ignore)
{
@ -1082,7 +1082,7 @@ static int calc (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
// this greatly reduces spam on hashcat console
const u64 pre_rejects_ignore = get_power (opencl_ctx, device_param) / 2;
const u64 pre_rejects_ignore = get_power (backend_ctx, device_param) / 2;
while (pre_rejects > pre_rejects_ignore)
{
@ -1658,11 +1658,11 @@ HC_API_CALL void *thread_calc (void *p)
hashcat_ctx_t *hashcat_ctx = thread_param->hashcat_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (opencl_ctx->enabled == false) return NULL;
if (backend_ctx->enabled == false) return NULL;
hc_device_param_t *device_param = opencl_ctx->devices_param + thread_param->tid;
hc_device_param_t *device_param = backend_ctx->devices_param + thread_param->tid;
if (device_param->skipped) return NULL;

11
src/emu_inc_platform.c Normal file
View File

@ -0,0 +1,11 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#include "common.h"
#include "types.h"
#include "bitops.h"
#include "emu_general.h"
#include "inc_platform.cl"

8
src/ext_cuda.c Normal file
View File

@ -0,0 +1,8 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#include "common.h"
#include "types.h"
#include "ext_cuda.h"

94
src/hashcat.c
View File

@ -34,7 +34,7 @@
#include "loopback.h"
#include "monitor.h"
#include "mpsp.h"
#include "opencl.h"
#include "backend.h"
#include "outfile_check.h"
#include "outfile.h"
#include "pidfile.h"
@ -59,7 +59,7 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
hashes_t *hashes = hashcat_ctx->hashes;
induct_ctx_t *induct_ctx = hashcat_ctx->induct_ctx;
logfile_ctx_t *logfile_ctx = hashcat_ctx->logfile_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
restore_ctx_t *restore_ctx = hashcat_ctx->restore_ctx;
status_ctx_t *status_ctx = hashcat_ctx->status_ctx;
user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra;
@ -109,7 +109,7 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
user_options->skip = 0;
}
opencl_session_reset (hashcat_ctx);
backend_session_reset (hashcat_ctx);
cpt_ctx_reset (hashcat_ctx);
@ -174,15 +174,15 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
* this is required for autotune
*/
opencl_ctx_devices_kernel_loops (hashcat_ctx);
backend_ctx_devices_kernel_loops (hashcat_ctx);
/**
* prepare thread buffers
*/
thread_param_t *threads_param = (thread_param_t *) hccalloc (opencl_ctx->devices_cnt, sizeof (thread_param_t));
thread_param_t *threads_param = (thread_param_t *) hccalloc (backend_ctx->devices_cnt, sizeof (thread_param_t));
hc_thread_t *c_threads = (hc_thread_t *) hccalloc (opencl_ctx->devices_cnt, sizeof (hc_thread_t));
hc_thread_t *c_threads = (hc_thread_t *) hccalloc (backend_ctx->devices_cnt, sizeof (hc_thread_t));
/**
* create autotune threads
@ -192,7 +192,7 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
status_ctx->devices_status = STATUS_AUTOTUNE;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
thread_param_t *thread_param = threads_param + device_id;
@ -202,7 +202,7 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
hc_thread_create (c_threads[device_id], thread_autotune, thread_param);
}
hc_thread_wait (opencl_ctx->devices_cnt, c_threads);
hc_thread_wait (backend_ctx->devices_cnt, c_threads);
EVENT (EVENT_AUTOTUNE_FINISHED);
@ -210,13 +210,13 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
* find same opencl devices and equal results
*/
opencl_ctx_devices_sync_tuning (hashcat_ctx);
backend_ctx_devices_sync_tuning (hashcat_ctx);
/**
* autotune modified kernel_accel, which modifies opencl_ctx->kernel_power_all
* autotune modified kernel_accel, which modifies backend_ctx->kernel_power_all
*/
opencl_ctx_devices_update_power (hashcat_ctx);
backend_ctx_devices_update_power (hashcat_ctx);
/**
* Begin loopback recording
@ -249,7 +249,7 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
status_ctx->accessible = true;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
thread_param_t *thread_param = threads_param + device_id;
@ -266,7 +266,7 @@ static int inner2_loop (hashcat_ctx_t *hashcat_ctx)
}
}
hc_thread_wait (opencl_ctx->devices_cnt, c_threads);
hc_thread_wait (backend_ctx->devices_cnt, c_threads);
hcfree (c_threads);
@ -438,7 +438,7 @@ static int outer_loop (hashcat_ctx_t *hashcat_ctx)
hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
hashes_t *hashes = hashcat_ctx->hashes;
mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
outcheck_ctx_t *outcheck_ctx = hashcat_ctx->outcheck_ctx;
restore_ctx_t *restore_ctx = hashcat_ctx->restore_ctx;
status_ctx_t *status_ctx = hashcat_ctx->status_ctx;
@ -722,7 +722,7 @@ static int outer_loop (hashcat_ctx_t *hashcat_ctx)
EVENT (EVENT_OPENCL_SESSION_PRE);
const int rc_session_begin = opencl_session_begin (hashcat_ctx);
const int rc_session_begin = backend_session_begin (hashcat_ctx);
if (rc_session_begin == -1) return -1;
@ -736,13 +736,13 @@ static int outer_loop (hashcat_ctx_t *hashcat_ctx)
{
EVENT (EVENT_SELFTEST_STARTING);
thread_param_t *threads_param = (thread_param_t *) hccalloc (opencl_ctx->devices_cnt, sizeof (thread_param_t));
thread_param_t *threads_param = (thread_param_t *) hccalloc (backend_ctx->devices_cnt, sizeof (thread_param_t));
hc_thread_t *selftest_threads = (hc_thread_t *) hccalloc (opencl_ctx->devices_cnt, sizeof (hc_thread_t));
hc_thread_t *selftest_threads = (hc_thread_t *) hccalloc (backend_ctx->devices_cnt, sizeof (hc_thread_t));
status_ctx->devices_status = STATUS_SELFTEST;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
thread_param_t *thread_param = threads_param + device_id;
@ -752,7 +752,7 @@ static int outer_loop (hashcat_ctx_t *hashcat_ctx)
hc_thread_create (selftest_threads[device_id], thread_selftest, thread_param);
}
hc_thread_wait (opencl_ctx->devices_cnt, selftest_threads);
hc_thread_wait (backend_ctx->devices_cnt, selftest_threads);
hcfree (threads_param);
@ -760,11 +760,11 @@ static int outer_loop (hashcat_ctx_t *hashcat_ctx)
// check for any selftest failures
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
if (opencl_ctx->enabled == false) continue;
if (backend_ctx->enabled == false) continue;
hc_device_param_t *device_param = opencl_ctx->devices_param + device_id;
hc_device_param_t *device_param = backend_ctx->devices_param + device_id;
if (device_param->skipped == true) continue;
@ -881,7 +881,7 @@ static int outer_loop (hashcat_ctx_t *hashcat_ctx)
// finalize opencl session
opencl_session_destroy (hashcat_ctx);
backend_session_destroy (hashcat_ctx);
// clean up
@ -930,7 +930,7 @@ int hashcat_init (hashcat_ctx_t *hashcat_ctx, void (*event) (const u32, struct h
hashcat_ctx->loopback_ctx = (loopback_ctx_t *) hcmalloc (sizeof (loopback_ctx_t));
hashcat_ctx->mask_ctx = (mask_ctx_t *) hcmalloc (sizeof (mask_ctx_t));
hashcat_ctx->module_ctx = (module_ctx_t *) hcmalloc (sizeof (module_ctx_t));
hashcat_ctx->opencl_ctx = (opencl_ctx_t *) hcmalloc (sizeof (opencl_ctx_t));
hashcat_ctx->backend_ctx = (backend_ctx_t *) hcmalloc (sizeof (backend_ctx_t));
hashcat_ctx->outcheck_ctx = (outcheck_ctx_t *) hcmalloc (sizeof (outcheck_ctx_t));
hashcat_ctx->outfile_ctx = (outfile_ctx_t *) hcmalloc (sizeof (outfile_ctx_t));
hashcat_ctx->pidfile_ctx = (pidfile_ctx_t *) hcmalloc (sizeof (pidfile_ctx_t));
@ -964,7 +964,7 @@ void hashcat_destroy (hashcat_ctx_t *hashcat_ctx)
hcfree (hashcat_ctx->loopback_ctx);
hcfree (hashcat_ctx->mask_ctx);
hcfree (hashcat_ctx->module_ctx);
hcfree (hashcat_ctx->opencl_ctx);
hcfree (hashcat_ctx->backend_ctx);
hcfree (hashcat_ctx->outcheck_ctx);
hcfree (hashcat_ctx->outfile_ctx);
hcfree (hashcat_ctx->pidfile_ctx);
@ -1172,15 +1172,15 @@ int hashcat_session_init (hashcat_ctx_t *hashcat_ctx, const char *install_folder
* Init OpenCL library loader
*/
const int rc_opencl_init = opencl_ctx_init (hashcat_ctx);
const int rc_backend_init = backend_ctx_init (hashcat_ctx);
if (rc_opencl_init == -1) return -1;
if (rc_backend_init == -1) return -1;
/**
* Init OpenCL devices
*/
const int rc_devices_init = opencl_ctx_devices_init (hashcat_ctx, comptime);
const int rc_devices_init = backend_ctx_devices_init (hashcat_ctx, comptime);
if (rc_devices_init == -1) return -1;
@ -1341,25 +1341,25 @@ int hashcat_session_destroy (hashcat_ctx_t *hashcat_ctx)
#endif
#endif
debugfile_destroy (hashcat_ctx);
dictstat_destroy (hashcat_ctx);
folder_config_destroy (hashcat_ctx);
hwmon_ctx_destroy (hashcat_ctx);
induct_ctx_destroy (hashcat_ctx);
logfile_destroy (hashcat_ctx);
loopback_destroy (hashcat_ctx);
opencl_ctx_devices_destroy (hashcat_ctx);
opencl_ctx_destroy (hashcat_ctx);
outcheck_ctx_destroy (hashcat_ctx);
outfile_destroy (hashcat_ctx);
pidfile_ctx_destroy (hashcat_ctx);
potfile_destroy (hashcat_ctx);
restore_ctx_destroy (hashcat_ctx);
tuning_db_destroy (hashcat_ctx);
user_options_destroy (hashcat_ctx);
user_options_extra_destroy (hashcat_ctx);
status_ctx_destroy (hashcat_ctx);
event_ctx_destroy (hashcat_ctx);
debugfile_destroy (hashcat_ctx);
dictstat_destroy (hashcat_ctx);
folder_config_destroy (hashcat_ctx);
hwmon_ctx_destroy (hashcat_ctx);
induct_ctx_destroy (hashcat_ctx);
logfile_destroy (hashcat_ctx);
loopback_destroy (hashcat_ctx);
backend_ctx_devices_destroy (hashcat_ctx);
backend_ctx_destroy (hashcat_ctx);
outcheck_ctx_destroy (hashcat_ctx);
outfile_destroy (hashcat_ctx);
pidfile_ctx_destroy (hashcat_ctx);
potfile_destroy (hashcat_ctx);
restore_ctx_destroy (hashcat_ctx);
tuning_db_destroy (hashcat_ctx);
user_options_destroy (hashcat_ctx);
user_options_extra_destroy (hashcat_ctx);
status_ctx_destroy (hashcat_ctx);
event_ctx_destroy (hashcat_ctx);
return 0;
}

2
src/hashes.c
View File

@ -14,7 +14,7 @@
#include "terminal.h"
#include "logfile.h"
#include "loopback.h"
#include "opencl.h"
#include "backend.h"
#include "outfile.h"
#include "potfile.h"
#include "rp.h"

138
src/hwmon.c
View File

@ -47,9 +47,9 @@ static void sysfs_close (hashcat_ctx_t *hashcat_ctx)
static char *hm_SYSFS_get_syspath_device (hashcat_ctx_t *hashcat_ctx, const int device_id)
{
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
char *syspath;
@ -1344,16 +1344,16 @@ static int hm_get_adapter_index_nvml (hashcat_ctx_t *hashcat_ctx, HM_ADAPTER_NVM
int hm_get_threshold_slowdown_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].threshold_slowdown_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1380,7 +1380,7 @@ int hm_get_threshold_slowdown_with_device_id (hashcat_ctx_t *hashcat_ctx, const
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1404,16 +1404,16 @@ int hm_get_threshold_slowdown_with_device_id (hashcat_ctx_t *hashcat_ctx, const
int hm_get_threshold_shutdown_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].threshold_shutdown_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1428,7 +1428,7 @@ int hm_get_threshold_shutdown_with_device_id (hashcat_ctx_t *hashcat_ctx, const
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1452,16 +1452,16 @@ int hm_get_threshold_shutdown_with_device_id (hashcat_ctx_t *hashcat_ctx, const
int hm_get_temperature_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].temperature_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1511,7 +1511,7 @@ int hm_get_temperature_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 dev
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1535,16 +1535,16 @@ int hm_get_temperature_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 dev
int hm_get_fanpolicy_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].fanpolicy_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1580,7 +1580,7 @@ int hm_get_fanpolicy_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 devic
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
return 1;
}
@ -1593,16 +1593,16 @@ int hm_get_fanpolicy_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 devic
int hm_get_fanspeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].fanspeed_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1658,7 +1658,7 @@ int hm_get_fanspeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1682,16 +1682,16 @@ int hm_get_fanspeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device
int hm_get_buslanes_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].buslanes_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1724,7 +1724,7 @@ int hm_get_buslanes_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1748,16 +1748,16 @@ int hm_get_buslanes_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device
int hm_get_utilization_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].utilization_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1776,7 +1776,7 @@ int hm_get_utilization_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 dev
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1800,16 +1800,16 @@ int hm_get_utilization_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 dev
int hm_get_memoryspeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].memoryspeed_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1842,7 +1842,7 @@ int hm_get_memoryspeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 dev
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1866,16 +1866,16 @@ int hm_get_memoryspeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 dev
int hm_get_corespeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].corespeed_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
if (hwmon_ctx->hm_adl)
{
@ -1908,7 +1908,7 @@ int hm_get_corespeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 devic
}
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1932,20 +1932,20 @@ int hm_get_corespeed_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 devic
int hm_get_throttle_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device_id)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (hwmon_ctx->enabled == false) return -1;
if (hwmon_ctx->hm_device[device_id].throttle_get_supported == false) return -1;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) return -1;
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_AMD)
{
}
if (opencl_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
if (backend_ctx->devices_param[device_id].device_vendor_id == VENDOR_ID_NV)
{
if (hwmon_ctx->hm_nvml)
{
@ -1961,7 +1961,7 @@ int hm_get_throttle_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device
clocksThrottleReasons &= ~nvmlClocksThrottleReasonApplicationsClocksSetting;
clocksThrottleReasons &= ~nvmlClocksThrottleReasonUnknown;
if (opencl_ctx->kernel_power_final)
if (backend_ctx->kernel_power_final)
{
clocksThrottleReasons &= ~nvmlClocksThrottleReasonHwSlowdown;
}
@ -1999,7 +1999,7 @@ int hm_get_throttle_with_device_id (hashcat_ctx_t *hashcat_ctx, const u32 device
int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
{
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
user_options_t *user_options = hashcat_ctx->user_options;
hwmon_ctx->enabled = false;
@ -2037,7 +2037,7 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
hcfree (hm_adapters_sysfs); \
}
if (opencl_ctx->need_nvml == true)
if (backend_ctx->need_nvml == true)
{
hwmon_ctx->hm_nvml = (NVML_PTR *) hcmalloc (sizeof (NVML_PTR));
@ -2049,7 +2049,7 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
}
}
if ((opencl_ctx->need_nvapi == true) && (hwmon_ctx->hm_nvml)) // nvapi can't work alone, we need nvml, too
if ((backend_ctx->need_nvapi == true) && (hwmon_ctx->hm_nvml)) // nvapi can't work alone, we need nvml, too
{
hwmon_ctx->hm_nvapi = (NVAPI_PTR *) hcmalloc (sizeof (NVAPI_PTR));
@ -2061,7 +2061,7 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
}
}
if (opencl_ctx->need_adl == true)
if (backend_ctx->need_adl == true)
{
hwmon_ctx->hm_adl = (ADL_PTR *) hcmalloc (sizeof (ADL_PTR));
@ -2073,7 +2073,7 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
}
}
if (opencl_ctx->need_sysfs == true)
if (backend_ctx->need_sysfs == true)
{
hwmon_ctx->hm_sysfs = (SYSFS_PTR *) hcmalloc (sizeof (SYSFS_PTR));
@ -2102,9 +2102,9 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
int tmp_in = hm_get_adapter_index_nvml (hashcat_ctx, nvmlGPUHandle);
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;
@ -2152,9 +2152,9 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
int tmp_in = hm_get_adapter_index_nvapi (hashcat_ctx, nvGPUHandle);
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;
@ -2221,9 +2221,9 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
return -1;
}
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;
@ -2271,9 +2271,9 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
{
int hm_adapters_id = 0;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if ((device_param->device_type & CL_DEVICE_TYPE_GPU) == 0) continue;
@ -2308,15 +2308,15 @@ int hwmon_ctx_init (hashcat_ctx_t *hashcat_ctx)
* save buffer required for later restores
*/
hwmon_ctx->od_clock_mem_status = (ADLOD6MemClockState *) hccalloc (opencl_ctx->devices_cnt, sizeof (ADLOD6MemClockState));
hwmon_ctx->od_clock_mem_status = (ADLOD6MemClockState *) hccalloc (backend_ctx->devices_cnt, sizeof (ADLOD6MemClockState));
/**
* HM devices: copy
*/
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;

2
src/interface.c
View File

@ -8,7 +8,7 @@
#include "memory.h"
#include "event.h"
#include "shared.h"
#include "opencl.h"
#include "backend.h"
#include "modules.h"
#include "dynloader.h"
#include "interface.h"

8
src/main.c
View File

@ -549,7 +549,7 @@ static void main_outerloop_mainscreen (MAYBE_UNUSED hashcat_ctx_t *hashcat_ctx,
event_log_info (hashcat_ctx, NULL);
}
static void main_opencl_session_pre (MAYBE_UNUSED hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED const void *buf, MAYBE_UNUSED const size_t len)
static void main_backend_session_pre (MAYBE_UNUSED hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED const void *buf, MAYBE_UNUSED const size_t len)
{
const user_options_t *user_options = hashcat_ctx->user_options;
@ -558,7 +558,7 @@ static void main_opencl_session_pre (MAYBE_UNUSED hashcat_ctx_t *hashcat_ctx, MA
event_log_info_nn (hashcat_ctx, "Initializing device kernels and memory...");
}
static void main_opencl_session_post (MAYBE_UNUSED hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED const void *buf, MAYBE_UNUSED const size_t len)
static void main_backend_session_post (MAYBE_UNUSED hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED const void *buf, MAYBE_UNUSED const size_t len)
{
const user_options_t *user_options = hashcat_ctx->user_options;
@ -1022,8 +1022,8 @@ static void event (const u32 id, hashcat_ctx_t *hashcat_ctx, const void *buf, co
case EVENT_MONITOR_PERFORMANCE_HINT: main_monitor_performance_hint (hashcat_ctx, buf, len); break;
case EVENT_MONITOR_NOINPUT_HINT: main_monitor_noinput_hint (hashcat_ctx, buf, len); break;
case EVENT_MONITOR_NOINPUT_ABORT: main_monitor_noinput_abort (hashcat_ctx, buf, len); break;
case EVENT_OPENCL_SESSION_POST: main_opencl_session_post (hashcat_ctx, buf, len); break;
case EVENT_OPENCL_SESSION_PRE: main_opencl_session_pre (hashcat_ctx, buf, len); break;
case EVENT_OPENCL_SESSION_POST: main_backend_session_post (hashcat_ctx, buf, len); break;
case EVENT_OPENCL_SESSION_PRE: main_backend_session_pre (hashcat_ctx, buf, len); break;
case EVENT_OPENCL_DEVICE_INIT_POST: main_opencl_device_init_post (hashcat_ctx, buf, len); break;
case EVENT_OPENCL_DEVICE_INIT_PRE: main_opencl_device_init_pre (hashcat_ctx, buf, len); break;
case EVENT_OUTERLOOP_FINISHED: main_outerloop_finished (hashcat_ctx, buf, len); break;

16
src/monitor.c
View File

@ -44,7 +44,7 @@ static int monitor (hashcat_ctx_t *hashcat_ctx)
{
hashes_t *hashes = hashcat_ctx->hashes;
hwmon_ctx_t *hwmon_ctx = hashcat_ctx->hwmon_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
restore_ctx_t *restore_ctx = hashcat_ctx->restore_ctx;
status_ctx_t *status_ctx = hashcat_ctx->status_ctx;
user_options_t *user_options = hashcat_ctx->user_options;
@ -114,13 +114,13 @@ static int monitor (hashcat_ctx_t *hashcat_ctx)
{
hc_thread_mutex_lock (status_ctx->mux_hwmon);
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;
if ((opencl_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) continue;
if ((backend_ctx->devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) continue;
const int temperature = hm_get_temperature_with_device_id (hashcat_ctx, device_id);
@ -132,9 +132,9 @@ static int monitor (hashcat_ctx_t *hashcat_ctx)
}
}
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;
@ -232,9 +232,9 @@ static int monitor (hashcat_ctx_t *hashcat_ctx)
hc_thread_mutex_lock (status_ctx->mux_hwmon);
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) continue;

10
src/mpsp.c
View File

@ -11,7 +11,7 @@
#include "logfile.h"
#include "convert.h"
#include "filehandling.h"
#include "opencl.h"
#include "backend.h"
#include "shared.h"
#include "ext_lzma.h"
#include "mpsp.h"
@ -1224,7 +1224,7 @@ int mask_ctx_update_loop (hashcat_ctx_t *hashcat_ctx)
return -1;
}
const int rc_update_mp = opencl_session_update_mp (hashcat_ctx);
const int rc_update_mp = backend_session_update_mp (hashcat_ctx);
if (rc_update_mp == -1) return -1;
}
@ -1257,13 +1257,13 @@ int mask_ctx_update_loop (hashcat_ctx_t *hashcat_ctx)
return -1;
}
const int rc_update_mp = opencl_session_update_mp (hashcat_ctx);
const int rc_update_mp = backend_session_update_mp (hashcat_ctx);
if (rc_update_mp == -1) return -1;
}
}
const int rc_update_combinator = opencl_session_update_combinator (hashcat_ctx);
const int rc_update_combinator = backend_session_update_combinator (hashcat_ctx);
if (rc_update_combinator == -1) return -1;
}
@ -1378,7 +1378,7 @@ int mask_ctx_update_loop (hashcat_ctx_t *hashcat_ctx)
return -1;
}
const int rc_update_mp_rl = opencl_session_update_mp_rl (hashcat_ctx, css_cnt_lr[0], css_cnt_lr[1]);
const int rc_update_mp_rl = backend_session_update_mp_rl (hashcat_ctx, css_cnt_lr[0], css_cnt_lr[1]);
if (rc_update_mp_rl == -1) return -1;
}

2
src/outfile.c
View File

@ -11,7 +11,7 @@
#include "rp.h"
#include "emu_inc_rp.h"
#include "emu_inc_rp_optimized.h"
#include "opencl.h"
#include "backend.h"
#include "shared.h"
#include "locking.h"
#include "outfile.h"

8
src/selftest.c
View File

@ -8,7 +8,7 @@
#include "event.h"
#include "bitops.h"
#include "convert.h"
#include "opencl.h"
#include "backend.h"
#include "thread.h"
#include "selftest.h"
@ -572,15 +572,15 @@ HC_API_CALL void *thread_selftest (void *p)
hashcat_ctx_t *hashcat_ctx = thread_param->hashcat_ctx;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
if (opencl_ctx->enabled == false) return NULL;
if (backend_ctx->enabled == false) return NULL;
user_options_t *user_options = hashcat_ctx->user_options;
if (user_options->self_test_disable == true) return NULL;
hc_device_param_t *device_param = opencl_ctx->devices_param + thread_param->tid;
hc_device_param_t *device_param = backend_ctx->devices_param + thread_param->tid;
if (device_param->skipped == true) return NULL;

120
src/status.c
View File

@ -200,32 +200,32 @@ double get_avg_exec_time (hc_device_param_t *device_param, const int last_num_en
int status_get_device_info_cnt (const hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
return opencl_ctx->devices_cnt;
return backend_ctx->devices_cnt;
}
int status_get_device_info_active (const hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
return opencl_ctx->devices_active;
return backend_ctx->devices_active;
}
bool status_get_skipped_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
return device_param->skipped;
}
bool status_get_skipped_warning_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
return device_param->skipped_warning;
}
@ -836,13 +836,13 @@ int status_get_guess_mask_length (const hashcat_ctx_t *hashcat_ctx)
char *status_get_guess_candidates_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
const status_ctx_t *status_ctx = hashcat_ctx->status_ctx;
const user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra;
if (status_ctx->accessible == false) return NULL;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
char *display = (char *) hcmalloc (HCBUFSIZ_TINY);
@ -1410,11 +1410,11 @@ u64 status_get_progress_end_relative_skip (const hashcat_ctx_t *hashcat_ctx)
double status_get_hashes_msec_all (const hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
double hashes_all_msec = 0;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
hashes_all_msec += status_get_hashes_msec_dev (hashcat_ctx, device_id);
}
@ -1424,12 +1424,12 @@ double status_get_hashes_msec_all (const hashcat_ctx_t *hashcat_ctx)
double status_get_hashes_msec_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
u64 speed_cnt = 0;
double speed_msec = 0;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if ((device_param->skipped == false) && (device_param->skipped_warning == false))
{
@ -1459,12 +1459,12 @@ double status_get_hashes_msec_dev_benchmark (const hashcat_ctx_t *hashcat_ctx, c
{
// this function increases accuracy for benchmark modes
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
u64 speed_cnt = 0;
double speed_msec = 0;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if ((device_param->skipped == false) && (device_param->skipped_warning == false))
{
@ -1486,11 +1486,11 @@ double status_get_hashes_msec_dev_benchmark (const hashcat_ctx_t *hashcat_ctx, c
double status_get_exec_msec_all (const hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
double exec_all_msec = 0;
for (u32 device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++)
for (u32 device_id = 0; device_id < backend_ctx->devices_cnt; device_id++)
{
exec_all_msec += status_get_exec_msec_dev (hashcat_ctx, device_id);
}
@ -1500,9 +1500,9 @@ double status_get_exec_msec_all (const hashcat_ctx_t *hashcat_ctx)
double status_get_exec_msec_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
double exec_dev_msec = 0;
@ -1700,9 +1700,9 @@ char *status_get_cpt (const hashcat_ctx_t *hashcat_ctx)
int status_get_salt_pos_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
int salt_pos = 0;
@ -1716,9 +1716,9 @@ int status_get_salt_pos_dev (const hashcat_ctx_t *hashcat_ctx, const int device_
int status_get_innerloop_pos_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
int innerloop_pos = 0;
@ -1732,9 +1732,9 @@ int status_get_innerloop_pos_dev (const hashcat_ctx_t *hashcat_ctx, const int de
int status_get_innerloop_left_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
int innerloop_left = 0;
@ -1748,9 +1748,9 @@ int status_get_innerloop_left_dev (const hashcat_ctx_t *hashcat_ctx, const int d
int status_get_iteration_pos_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
int iteration_pos = 0;
@ -1764,9 +1764,9 @@ int status_get_iteration_pos_dev (const hashcat_ctx_t *hashcat_ctx, const int de
int status_get_iteration_left_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
int iteration_left = 0;
@ -1781,9 +1781,9 @@ int status_get_iteration_left_dev (const hashcat_ctx_t *hashcat_ctx, const int d
#ifdef WITH_BRAIN
int status_get_brain_link_client_id_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
int brain_client_id = -1;
@ -1797,9 +1797,9 @@ int status_get_brain_link_client_id_dev (const hashcat_ctx_t *hashcat_ctx, const
int status_get_brain_link_status_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
int brain_link_status_dev = 0;
@ -1815,9 +1815,9 @@ int status_get_brain_link_status_dev (const hashcat_ctx_t *hashcat_ctx, const in
char *status_get_brain_link_recv_bytes_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
u64 brain_link_recv_bytes = 0;
@ -1835,9 +1835,9 @@ char *status_get_brain_link_recv_bytes_dev (const hashcat_ctx_t *hashcat_ctx, co
char *status_get_brain_link_send_bytes_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
u64 brain_link_send_bytes = 0;
@ -1855,9 +1855,9 @@ char *status_get_brain_link_send_bytes_dev (const hashcat_ctx_t *hashcat_ctx, co
char *status_get_brain_link_recv_bytes_sec_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
u64 brain_link_recv_bytes = 0;
@ -1882,9 +1882,9 @@ char *status_get_brain_link_recv_bytes_sec_dev (const hashcat_ctx_t *hashcat_ctx
char *status_get_brain_link_send_bytes_sec_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
u64 brain_link_send_bytes = 0;
@ -1910,9 +1910,9 @@ char *status_get_brain_link_send_bytes_sec_dev (const hashcat_ctx_t *hashcat_ctx
char *status_get_hwmon_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
char *output_buf = (char *) hcmalloc (HCBUFSIZ_TINY);
@ -1983,9 +1983,9 @@ char *status_get_hwmon_dev (const hashcat_ctx_t *hashcat_ctx, const int device_i
int status_get_corespeed_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return -1;
@ -2004,9 +2004,9 @@ int status_get_corespeed_dev (const hashcat_ctx_t *hashcat_ctx, const int device
int status_get_memoryspeed_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return -1;
@ -2025,9 +2025,9 @@ int status_get_memoryspeed_dev (const hashcat_ctx_t *hashcat_ctx, const int devi
u64 status_get_progress_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return 0;
@ -2038,9 +2038,9 @@ u64 status_get_progress_dev (const hashcat_ctx_t *hashcat_ctx, const int device_
double status_get_runtime_msec_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return 0;
@ -2051,9 +2051,9 @@ double status_get_runtime_msec_dev (const hashcat_ctx_t *hashcat_ctx, const int
int status_get_kernel_accel_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return 0;
@ -2066,9 +2066,9 @@ int status_get_kernel_accel_dev (const hashcat_ctx_t *hashcat_ctx, const int dev
int status_get_kernel_loops_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return 0;
@ -2081,9 +2081,9 @@ int status_get_kernel_loops_dev (const hashcat_ctx_t *hashcat_ctx, const int dev
int status_get_kernel_threads_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return 0;
@ -2094,9 +2094,9 @@ int status_get_kernel_threads_dev (const hashcat_ctx_t *hashcat_ctx, const int d
int status_get_vector_width_dev (const hashcat_ctx_t *hashcat_ctx, const int device_id)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id];
hc_device_param_t *device_param = &backend_ctx->devices_param[device_id];
if (device_param->skipped == true) return 0;

2
src/stdout.c
View File

@ -10,7 +10,7 @@
#include "emu_inc_rp.h"
#include "emu_inc_rp_optimized.h"
#include "mpsp.h"
#include "opencl.h"
#include "backend.h"
#include "shared.h"
#include "stdout.h"

30
src/terminal.c
View File

@ -656,17 +656,17 @@ void example_hashes (hashcat_ctx_t *hashcat_ctx)
void opencl_info (hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
event_log_info (hashcat_ctx, "OpenCL Info:");
event_log_info (hashcat_ctx, NULL);
cl_uint platforms_cnt = opencl_ctx->platforms_cnt;
cl_platform_id *platforms = opencl_ctx->platforms;
char **platforms_vendor = opencl_ctx->platforms_vendor;
char **platforms_name = opencl_ctx->platforms_name;
char **platforms_version = opencl_ctx->platforms_version;
cl_uint devices_cnt = opencl_ctx->devices_cnt;
cl_uint platforms_cnt = backend_ctx->platforms_cnt;
cl_platform_id *platforms = backend_ctx->platforms;
char **platforms_vendor = backend_ctx->platforms_vendor;
char **platforms_name = backend_ctx->platforms_name;
char **platforms_version = backend_ctx->platforms_version;
cl_uint devices_cnt = backend_ctx->devices_cnt;
for (cl_uint platforms_idx = 0; platforms_idx < platforms_cnt; platforms_idx++)
{
@ -683,7 +683,7 @@ void opencl_info (hashcat_ctx_t *hashcat_ctx)
for (cl_uint devices_idx = 0; devices_idx < devices_cnt; devices_idx++)
{
const hc_device_param_t *device_param = opencl_ctx->devices_param + devices_idx;
const hc_device_param_t *device_param = backend_ctx->devices_param + devices_idx;
if (device_param->platform != platform_id) continue;
@ -717,18 +717,18 @@ void opencl_info (hashcat_ctx_t *hashcat_ctx)
void opencl_info_compact (hashcat_ctx_t *hashcat_ctx)
{
const opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
const backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
const user_options_t *user_options = hashcat_ctx->user_options;
if (user_options->quiet == true) return;
if (user_options->machine_readable == true) return;
if (user_options->status_json == true) return;
cl_uint platforms_cnt = opencl_ctx->platforms_cnt;
cl_platform_id *platforms = opencl_ctx->platforms;
char **platforms_vendor = opencl_ctx->platforms_vendor;
bool *platforms_skipped = opencl_ctx->platforms_skipped;
cl_uint devices_cnt = opencl_ctx->devices_cnt;
cl_uint platforms_cnt = backend_ctx->platforms_cnt;
cl_platform_id *platforms = backend_ctx->platforms;
char **platforms_vendor = backend_ctx->platforms_vendor;
bool *platforms_skipped = backend_ctx->platforms_skipped;
cl_uint devices_cnt = backend_ctx->devices_cnt;
for (cl_uint platforms_idx = 0; platforms_idx < platforms_cnt; platforms_idx++)
{
@ -755,7 +755,7 @@ void opencl_info_compact (hashcat_ctx_t *hashcat_ctx)
for (cl_uint devices_idx = 0; devices_idx < devices_cnt; devices_idx++)
{
const hc_device_param_t *device_param = opencl_ctx->devices_param + devices_idx;
const hc_device_param_t *device_param = backend_ctx->devices_param + devices_idx;
if (device_param->platform != platform_id) continue;

2
src/user_options.c
View File

@ -11,7 +11,7 @@
#include "interface.h"
#include "shared.h"
#include "usage.h"
#include "opencl.h"
#include "backend.h"
#include "user_options.h"
#ifdef WITH_BRAIN