diff --git a/docs/changes.txt b/docs/changes.txt index ddfecb45d..a6e23d0fa 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -145,11 +145,12 @@ - Apple Driver: Updated requirements to use Apple OpenCL API to macOS 13.0 - use - Backend: Added workaround to get rid of internal runtimes memory leaks - Backend: Updated filename chksum format to prevent invalid cache on Apple Silicon when switching arch +- Backend: Updated OpenCL/CUDA/HIP/Metal API's - Backend: Splitting backend_ctx_devices_init into smaller runtime-specific functions - Backend Checks: Describe workaround in error message when detecting more than 64 backend devices - Backend Info: Added --machine-readable format - Brain: Added sanity check and corresponding error message for invalid --brain-port values -- Debug: Added -g to build_options if DEBUG >= 1 and not on Apple +- Debug: Added -g to build_options if DEBUG >= 1 (only with HIP and OpenCL) - Dependencies: Added sse2neon v1.8.0 (commit 658eeac) - Dependencies: Updated LZMA SDK to 24.09 - Dependencies: Updated unrar source to 6.2.7 diff --git a/include/ext_OpenCL.h b/include/ext_OpenCL.h index 7044ac667..5f8339fb9 100644 --- a/include/ext_OpenCL.h +++ b/include/ext_OpenCL.h @@ -126,6 +126,10 @@ const char *val2cstr_cl (cl_int CL_err); int ocl_init (void *hashcat_ctx); void ocl_close (void *hashcat_ctx); +int hc_clReleaseMemObjectPtr (void *hashcat_ctx, cl_mem *mem); +int hc_clReleaseKernelPtr (void *hashcat_ctx, cl_kernel *kernel); +int hc_clReleaseProgramPtr (void *hashcat_ctx, cl_program *program); + int hc_clEnqueueNDRangeKernel (void *hashcat_ctx, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event); int hc_clGetEventInfo (void *hashcat_ctx, cl_event event, cl_event_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret); int hc_clFlush (void *hashcat_ctx, cl_command_queue command_queue); diff --git a/include/ext_cuda.h b/include/ext_cuda.h index a80af5367..28e5f4e2f 100644 --- a/include/ext_cuda.h +++ b/include/ext_cuda.h @@ -1260,6 +1260,11 @@ typedef hc_cuda_lib_t CUDA_PTR; int cuda_init (void *hashcat_ctx); void cuda_close (void *hashcat_ctx); +int hc_cuEventDestroyPtr (void *hashcat_ctx, CUevent *hEvent); +int hc_cuMemFreePtr (void *hashcat_ctx, CUdeviceptr *dptr); +int hc_cuModuleUnloadPtr (void *hashcat_ctx, CUmodule *hmod); +int hc_cuStreamDestroyPtr (void *hashcat_ctx, CUstream *hStream); + int hc_cuCtxCreate (void *hashcat_ctx, CUcontext *pctx, unsigned int flags, CUdevice dev); int hc_cuCtxDestroy (void *hashcat_ctx, CUcontext ctx); int hc_cuCtxSetCurrent (void *hashcat_ctx, CUcontext ctx); diff --git a/include/ext_hip.h b/include/ext_hip.h index 3c0b8433a..569819e76 100644 --- a/include/ext_hip.h +++ b/include/ext_hip.h @@ -650,6 +650,11 @@ typedef hc_hip_lib_t HIP_PTR; int hip_init (void *hashcat_ctx); void hip_close (void *hashcat_ctx); +int hc_hipEventDestroyPtr (void *hashcat_ctx, hipEvent_t *hEvent); +int hc_hipMemFreePtr (void *hashcat_ctx, hipDeviceptr_t *dptr); +int hc_hipModuleUnloadPtr (void *hashcat_ctx, hipModule_t *hmod); +int hc_hipStreamDestroyPtr (void *hashcat_ctx, hipStream_t *hStream); + // deprecated int hc_hipCtxCreate (void *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev); // deprecated diff --git a/include/ext_metal.h b/include/ext_metal.h index afa641871..8e319fd7d 100644 --- a/include/ext_metal.h +++ b/include/ext_metal.h @@ -98,9 +98,9 @@ int hc_mtlMemcpyHtoD (void *hashcat_ctx, mtl_command_queu // read int hc_mtlMemcpyDtoH (void *hashcat_ctx, mtl_command_queue command_queue, void *buf_dst, mtl_mem buf_src, size_t buf_src_off, size_t buf_size); -int hc_mtlReleaseMemObject (void *hashcat_ctx, mtl_mem metal_buffer); -int hc_mtlReleaseFunction (void *hashcat_ctx, mtl_function metal_function); -int hc_mtlReleaseLibrary (void *hashcat_ctx, mtl_function metal_library); +int hc_mtlReleaseMemObject (void *hashcat_ctx, mtl_mem *metal_buffer); +int hc_mtlReleaseFunction (void *hashcat_ctx, mtl_function *metal_function); +int hc_mtlReleaseLibrary (void *hashcat_ctx, mtl_function *metal_library); int hc_mtlReleaseCommandQueue (void *hashcat_ctx, mtl_command_queue command_queue); int hc_mtlReleaseDevice (void *hashcat_ctx, mtl_device_id metal_device); diff --git a/src/backend.c b/src/backend.c index 6d78fb262..1d26df77c 100644 --- a/src/backend.c +++ b/src/backend.c @@ -774,12 +774,12 @@ static bool opencl_test_instruction (hashcat_ctx_t *hashcat_ctx, cl_context cont #endif - hc_clReleaseProgram (hashcat_ctx, program); + hc_clReleaseProgramPtr (hashcat_ctx, &program); return false; } - if (hc_clReleaseProgram (hashcat_ctx, program) == -1) return false; + if (hc_clReleaseProgramPtr (hashcat_ctx, &program) == -1) return false; return true; } @@ -3572,7 +3572,7 @@ int run_kernel_amp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, for (int i = 0; i < tmp_buf_cnt; i++) { - hc_mtlReleaseMemObject (hashcat_ctx, tmp_buf[i]); + hc_mtlReleaseMemObject (hashcat_ctx, &tmp_buf[i]); tmp_buf[i] = NULL; } @@ -8705,7 +8705,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) if (tmp_device[c] != NULL) { - if (hc_mtlReleaseMemObject (hashcat_ctx, tmp_device[c]) == -1) return -1; + if (hc_mtlReleaseMemObject (hashcat_ctx, &tmp_device[c]) == -1) return -1; } } @@ -9011,7 +9011,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) if (tmp_device[c] != NULL) { - if (hc_clReleaseMemObject (hashcat_ctx, tmp_device[c]) == -1) r = -1; + if (hc_clReleaseMemObjectPtr (hashcat_ctx, &tmp_device[c]) == -1) r = -1; } } @@ -9988,7 +9988,7 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p *opencl_program = fin; - hc_clReleaseProgram (hashcat_ctx, p1); + hc_clReleaseProgramPtr (hashcat_ctx, &p1); } if (cache_disable == false) @@ -11051,10 +11051,14 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) char *build_options_buf = (char *) hcmalloc (build_options_sz); - #if !defined (__APPLE__) && defined (DEBUG) && (DEBUG >= 1) - int build_options_len = snprintf(build_options_buf, build_options_sz, "-g -D KERNEL_STATIC "); - #else int build_options_len = snprintf(build_options_buf, build_options_sz, "-D KERNEL_STATIC "); + + #if defined (DEBUG) && (DEBUG >= 1) + // only HIP and OpenCL have '-g' + if (device_param->is_hip == true || device_param->is_opencl == true) + { + build_options_len += snprintf (build_options_buf + build_options_len, build_options_sz - build_options_len, "-g "); + } #endif if ((device_param->is_cuda == true) || (device_param->is_hip == true)) @@ -17004,552 +17008,325 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx) if (device_param->is_cuda == true) { - if (device_param->cuda_d_pws_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_pws_buf); - if (device_param->cuda_d_pws_amp_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_pws_amp_buf); - if (device_param->cuda_d_pws_comp_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_pws_comp_buf); - if (device_param->cuda_d_pws_idx) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_pws_idx); - if (device_param->cuda_d_rules) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_rules); - //if (device_param->cuda_d_rules_c) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_rules_c); - if (device_param->cuda_d_combs) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_combs); - if (device_param->cuda_d_combs_c) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_combs_c); - if (device_param->cuda_d_bfs) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bfs); - //if (device_param->cuda_d_bfs_c) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bfs_c); - if (device_param->cuda_d_bitmap_s1_a) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s1_a); - if (device_param->cuda_d_bitmap_s1_b) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s1_b); - if (device_param->cuda_d_bitmap_s1_c) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s1_c); - if (device_param->cuda_d_bitmap_s1_d) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s1_d); - if (device_param->cuda_d_bitmap_s2_a) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s2_a); - if (device_param->cuda_d_bitmap_s2_b) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s2_b); - if (device_param->cuda_d_bitmap_s2_c) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s2_c); - if (device_param->cuda_d_bitmap_s2_d) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_bitmap_s2_d); - if (device_param->cuda_d_plain_bufs) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_plain_bufs); - if (device_param->cuda_d_digests_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_digests_buf); - if (device_param->cuda_d_digests_shown) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_digests_shown); - if (device_param->cuda_d_salt_bufs) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_salt_bufs); - if (device_param->cuda_d_esalt_bufs) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_esalt_bufs); - if (device_param->cuda_d_tmps) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_tmps); - if (device_param->cuda_d_hooks) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_hooks); - if (device_param->cuda_d_result) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_result); - if (device_param->cuda_d_extra0_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_extra0_buf); - if (device_param->cuda_d_extra1_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_extra1_buf); - if (device_param->cuda_d_extra2_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_extra2_buf); - if (device_param->cuda_d_extra3_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_extra3_buf); - if (device_param->cuda_d_root_css_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_root_css_buf); - if (device_param->cuda_d_markov_css_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_markov_css_buf); - if (device_param->cuda_d_tm_c) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_tm_c); - if (device_param->cuda_d_st_digests_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_st_digests_buf); - if (device_param->cuda_d_st_salts_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_st_salts_buf); - if (device_param->cuda_d_st_esalts_buf) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_st_esalts_buf); - if (device_param->cuda_d_kernel_param) hc_cuMemFree (hashcat_ctx, device_param->cuda_d_kernel_param); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_pws_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_pws_amp_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_pws_comp_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_pws_idx); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_rules); + //hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_rules_c); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_combs); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_combs_c); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bfs); + //hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bfs_c); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s1_a); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s1_b); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s1_c); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s1_d); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s2_a); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s2_b); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s2_c); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_bitmap_s2_d); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_plain_bufs); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_digests_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_digests_shown); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_salt_bufs); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_esalt_bufs); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_tmps); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_hooks); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_result); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_extra0_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_extra1_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_extra2_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_extra3_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_root_css_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_markov_css_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_tm_c); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_st_digests_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_st_salts_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_st_esalts_buf); + hc_cuMemFreePtr (hashcat_ctx, &device_param->cuda_d_kernel_param); - if (device_param->cuda_event1) hc_cuEventDestroy (hashcat_ctx, device_param->cuda_event1); - if (device_param->cuda_event2) hc_cuEventDestroy (hashcat_ctx, device_param->cuda_event2); - if (device_param->cuda_event3) hc_cuEventDestroy (hashcat_ctx, device_param->cuda_event3); + hc_cuEventDestroyPtr (hashcat_ctx, &device_param->cuda_event1); + hc_cuEventDestroyPtr (hashcat_ctx, &device_param->cuda_event2); + hc_cuEventDestroyPtr (hashcat_ctx, &device_param->cuda_event3); - if (device_param->cuda_stream) hc_cuStreamDestroy (hashcat_ctx, device_param->cuda_stream); + hc_cuStreamDestroyPtr (hashcat_ctx, &device_param->cuda_stream); - if (device_param->cuda_module) hc_cuModuleUnload (hashcat_ctx, device_param->cuda_module); - if (device_param->cuda_module_mp) hc_cuModuleUnload (hashcat_ctx, device_param->cuda_module_mp); - if (device_param->cuda_module_amp) hc_cuModuleUnload (hashcat_ctx, device_param->cuda_module_amp); - if (device_param->cuda_module_shared) hc_cuModuleUnload (hashcat_ctx, device_param->cuda_module_shared); + hc_cuModuleUnloadPtr (hashcat_ctx, &device_param->cuda_module); + hc_cuModuleUnloadPtr (hashcat_ctx, &device_param->cuda_module_mp); + hc_cuModuleUnloadPtr (hashcat_ctx, &device_param->cuda_module_amp); + hc_cuModuleUnloadPtr (hashcat_ctx, &device_param->cuda_module_shared); - //if (device_param->cuda_context) hc_cuCtxDestroy (hashcat_ctx, device_param->cuda_context); + device_param->cuda_d_rules_c = 0; + device_param->cuda_d_bfs_c = 0; - device_param->cuda_d_pws_buf = 0; - device_param->cuda_d_pws_amp_buf = 0; - device_param->cuda_d_pws_comp_buf = 0; - device_param->cuda_d_pws_idx = 0; - device_param->cuda_d_rules = 0; - device_param->cuda_d_rules_c = 0; - device_param->cuda_d_combs = 0; - device_param->cuda_d_combs_c = 0; - device_param->cuda_d_bfs = 0; - device_param->cuda_d_bfs_c = 0; - device_param->cuda_d_bitmap_s1_a = 0; - device_param->cuda_d_bitmap_s1_b = 0; - device_param->cuda_d_bitmap_s1_c = 0; - device_param->cuda_d_bitmap_s1_d = 0; - device_param->cuda_d_bitmap_s2_a = 0; - device_param->cuda_d_bitmap_s2_b = 0; - device_param->cuda_d_bitmap_s2_c = 0; - device_param->cuda_d_bitmap_s2_d = 0; - device_param->cuda_d_plain_bufs = 0; - device_param->cuda_d_digests_buf = 0; - device_param->cuda_d_digests_shown = 0; - device_param->cuda_d_salt_bufs = 0; - device_param->cuda_d_esalt_bufs = 0; - device_param->cuda_d_tmps = 0; - device_param->cuda_d_hooks = 0; - device_param->cuda_d_result = 0; - device_param->cuda_d_extra0_buf = 0; - device_param->cuda_d_extra1_buf = 0; - device_param->cuda_d_extra2_buf = 0; - device_param->cuda_d_extra3_buf = 0; - device_param->cuda_d_root_css_buf = 0; - device_param->cuda_d_markov_css_buf = 0; - device_param->cuda_d_tm_c = 0; - device_param->cuda_d_st_digests_buf = 0; - device_param->cuda_d_st_salts_buf = 0; - device_param->cuda_d_st_esalts_buf = 0; - device_param->cuda_d_kernel_param = 0; - - device_param->cuda_function1 = NULL; - device_param->cuda_function12 = NULL; - device_param->cuda_function2p = NULL; - device_param->cuda_function2 = NULL; - device_param->cuda_function2e = NULL; - device_param->cuda_function23 = NULL; - device_param->cuda_function3 = NULL; - device_param->cuda_function4 = NULL; - device_param->cuda_function_init2 = NULL; - device_param->cuda_function_loop2p = NULL; - device_param->cuda_function_loop2 = NULL; - device_param->cuda_function_mp = NULL; - device_param->cuda_function_mp_l = NULL; - device_param->cuda_function_mp_r = NULL; - device_param->cuda_function_tm = NULL; - device_param->cuda_function_amp = NULL; - device_param->cuda_function_memset = NULL; - device_param->cuda_function_bzero = NULL; - device_param->cuda_function_atinit = NULL; + device_param->cuda_function1 = NULL; + device_param->cuda_function12 = NULL; + device_param->cuda_function2p = NULL; + device_param->cuda_function2 = NULL; + device_param->cuda_function2e = NULL; + device_param->cuda_function23 = NULL; + device_param->cuda_function3 = NULL; + device_param->cuda_function4 = NULL; + device_param->cuda_function_init2 = NULL; + device_param->cuda_function_loop2p = NULL; + device_param->cuda_function_loop2 = NULL; + device_param->cuda_function_mp = NULL; + device_param->cuda_function_mp_l = NULL; + device_param->cuda_function_mp_r = NULL; + device_param->cuda_function_tm = NULL; + device_param->cuda_function_amp = NULL; + device_param->cuda_function_memset = NULL; + device_param->cuda_function_bzero = NULL; + device_param->cuda_function_atinit = NULL; device_param->cuda_function_utf8toutf16le = NULL; - device_param->cuda_function_decompress = NULL; - device_param->cuda_function_aux1 = NULL; - device_param->cuda_function_aux2 = NULL; - device_param->cuda_function_aux3 = NULL; - device_param->cuda_function_aux4 = NULL; - - device_param->cuda_event1 = NULL; - device_param->cuda_event2 = NULL; - device_param->cuda_event3 = NULL; - - device_param->cuda_stream = NULL; - - device_param->cuda_module = NULL; - device_param->cuda_module_mp = NULL; - device_param->cuda_module_amp = NULL; - device_param->cuda_module_shared = NULL; + device_param->cuda_function_decompress = NULL; + device_param->cuda_function_aux1 = NULL; + device_param->cuda_function_aux2 = NULL; + device_param->cuda_function_aux3 = NULL; + device_param->cuda_function_aux4 = NULL; + //if (device_param->cuda_context) hc_cuCtxDestroy (hashcat_ctx, device_param->cuda_context); //device_param->cuda_context = NULL; } if (device_param->is_hip == true) { - if (device_param->hip_d_pws_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_buf); - if (device_param->hip_d_pws_amp_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_amp_buf); - if (device_param->hip_d_pws_comp_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_comp_buf); - if (device_param->hip_d_pws_idx) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_idx); - if (device_param->hip_d_rules) hc_hipMemFree (hashcat_ctx, device_param->hip_d_rules); - //if (device_param->hip_d_rules_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_rules_c); - if (device_param->hip_d_combs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_combs); - if (device_param->hip_d_combs_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_combs_c); - if (device_param->hip_d_bfs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bfs); - //if (device_param->hip_d_bfs_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bfs_c); - if (device_param->hip_d_bitmap_s1_a) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_a); - if (device_param->hip_d_bitmap_s1_b) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_b); - if (device_param->hip_d_bitmap_s1_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_c); - if (device_param->hip_d_bitmap_s1_d) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_d); - if (device_param->hip_d_bitmap_s2_a) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_a); - if (device_param->hip_d_bitmap_s2_b) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_b); - if (device_param->hip_d_bitmap_s2_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_c); - if (device_param->hip_d_bitmap_s2_d) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_d); - if (device_param->hip_d_plain_bufs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_plain_bufs); - if (device_param->hip_d_digests_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_digests_buf); - if (device_param->hip_d_digests_shown) hc_hipMemFree (hashcat_ctx, device_param->hip_d_digests_shown); - if (device_param->hip_d_salt_bufs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_salt_bufs); - if (device_param->hip_d_esalt_bufs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_esalt_bufs); - if (device_param->hip_d_tmps) hc_hipMemFree (hashcat_ctx, device_param->hip_d_tmps); - if (device_param->hip_d_hooks) hc_hipMemFree (hashcat_ctx, device_param->hip_d_hooks); - if (device_param->hip_d_result) hc_hipMemFree (hashcat_ctx, device_param->hip_d_result); - if (device_param->hip_d_extra0_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra0_buf); - if (device_param->hip_d_extra1_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra1_buf); - if (device_param->hip_d_extra2_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra2_buf); - if (device_param->hip_d_extra3_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra3_buf); - if (device_param->hip_d_root_css_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_root_css_buf); - if (device_param->hip_d_markov_css_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_markov_css_buf); - if (device_param->hip_d_tm_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_tm_c); - if (device_param->hip_d_st_digests_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_st_digests_buf); - if (device_param->hip_d_st_salts_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_st_salts_buf); - if (device_param->hip_d_st_esalts_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_st_esalts_buf); - if (device_param->hip_d_kernel_param) hc_hipMemFree (hashcat_ctx, device_param->hip_d_kernel_param); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_pws_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_pws_amp_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_pws_comp_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_pws_idx); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_rules); + //hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_rules_c); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_combs); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_combs_c); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bfs); + //hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bfs_c); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s1_a); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s1_b); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s1_c); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s1_d); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s2_a); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s2_b); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s2_c); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_bitmap_s2_d); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_plain_bufs); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_digests_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_digests_shown); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_salt_bufs); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_esalt_bufs); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_tmps); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_hooks); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_result); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_extra0_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_extra1_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_extra2_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_extra3_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_root_css_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_markov_css_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_tm_c); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_st_digests_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_st_salts_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_st_esalts_buf); + hc_hipMemFreePtr (hashcat_ctx, &device_param->hip_d_kernel_param); - if (device_param->hip_event1) hc_hipEventDestroy (hashcat_ctx, device_param->hip_event1); - if (device_param->hip_event2) hc_hipEventDestroy (hashcat_ctx, device_param->hip_event2); - if (device_param->hip_event3) hc_hipEventDestroy (hashcat_ctx, device_param->hip_event3); + hc_hipEventDestroyPtr (hashcat_ctx, &device_param->hip_event1); + hc_hipEventDestroyPtr (hashcat_ctx, &device_param->hip_event2); + hc_hipEventDestroyPtr (hashcat_ctx, &device_param->hip_event3); - if (device_param->hip_stream) hc_hipStreamDestroy (hashcat_ctx, device_param->hip_stream); + hc_hipStreamDestroyPtr (hashcat_ctx, &device_param->hip_stream); - if (device_param->hip_module) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module); - if (device_param->hip_module_mp) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module_mp); - if (device_param->hip_module_amp) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module_amp); - if (device_param->hip_module_shared) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module_shared); + hc_hipModuleUnloadPtr (hashcat_ctx, &device_param->hip_module); + hc_hipModuleUnloadPtr (hashcat_ctx, &device_param->hip_module_mp); + hc_hipModuleUnloadPtr (hashcat_ctx, &device_param->hip_module_amp); + hc_hipModuleUnloadPtr (hashcat_ctx, &device_param->hip_module_shared); - device_param->hip_d_pws_buf = 0; - device_param->hip_d_pws_amp_buf = 0; - device_param->hip_d_pws_comp_buf = 0; - device_param->hip_d_pws_idx = 0; - device_param->hip_d_rules = 0; - device_param->hip_d_rules_c = 0; - device_param->hip_d_combs = 0; - device_param->hip_d_combs_c = 0; - device_param->hip_d_bfs = 0; - device_param->hip_d_bfs_c = 0; - device_param->hip_d_bitmap_s1_a = 0; - device_param->hip_d_bitmap_s1_b = 0; - device_param->hip_d_bitmap_s1_c = 0; - device_param->hip_d_bitmap_s1_d = 0; - device_param->hip_d_bitmap_s2_a = 0; - device_param->hip_d_bitmap_s2_b = 0; - device_param->hip_d_bitmap_s2_c = 0; - device_param->hip_d_bitmap_s2_d = 0; - device_param->hip_d_plain_bufs = 0; - device_param->hip_d_digests_buf = 0; - device_param->hip_d_digests_shown = 0; - device_param->hip_d_salt_bufs = 0; - device_param->hip_d_esalt_bufs = 0; - device_param->hip_d_tmps = 0; - device_param->hip_d_hooks = 0; - device_param->hip_d_result = 0; - device_param->hip_d_extra0_buf = 0; - device_param->hip_d_extra1_buf = 0; - device_param->hip_d_extra2_buf = 0; - device_param->hip_d_extra3_buf = 0; - device_param->hip_d_root_css_buf = 0; - device_param->hip_d_markov_css_buf = 0; - device_param->hip_d_tm_c = 0; - device_param->hip_d_st_digests_buf = 0; - device_param->hip_d_st_salts_buf = 0; - device_param->hip_d_st_esalts_buf = 0; - device_param->hip_d_kernel_param = 0; + device_param->hip_d_rules_c = 0; + device_param->hip_d_bfs_c = 0; - device_param->hip_function1 = NULL; - device_param->hip_function12 = NULL; - device_param->hip_function2p = NULL; - device_param->hip_function2 = NULL; - device_param->hip_function2e = NULL; - device_param->hip_function23 = NULL; - device_param->hip_function3 = NULL; - device_param->hip_function4 = NULL; - device_param->hip_function_init2 = NULL; - device_param->hip_function_loop2p = NULL; - device_param->hip_function_loop2 = NULL; - device_param->hip_function_mp = NULL; - device_param->hip_function_mp_l = NULL; - device_param->hip_function_mp_r = NULL; - device_param->hip_function_tm = NULL; - device_param->hip_function_amp = NULL; - device_param->hip_function_memset = NULL; - device_param->hip_function_bzero = NULL; - device_param->hip_function_atinit = NULL; + device_param->hip_function1 = NULL; + device_param->hip_function12 = NULL; + device_param->hip_function2p = NULL; + device_param->hip_function2 = NULL; + device_param->hip_function2e = NULL; + device_param->hip_function23 = NULL; + device_param->hip_function3 = NULL; + device_param->hip_function4 = NULL; + device_param->hip_function_init2 = NULL; + device_param->hip_function_loop2p = NULL; + device_param->hip_function_loop2 = NULL; + device_param->hip_function_mp = NULL; + device_param->hip_function_mp_l = NULL; + device_param->hip_function_mp_r = NULL; + device_param->hip_function_tm = NULL; + device_param->hip_function_amp = NULL; + device_param->hip_function_memset = NULL; + device_param->hip_function_bzero = NULL; + device_param->hip_function_atinit = NULL; device_param->hip_function_utf8toutf16le = NULL; - device_param->hip_function_decompress = NULL; - device_param->hip_function_aux1 = NULL; - device_param->hip_function_aux2 = NULL; - device_param->hip_function_aux3 = NULL; - device_param->hip_function_aux4 = NULL; - - device_param->hip_event1 = NULL; - device_param->hip_event2 = NULL; - device_param->hip_event3 = NULL; - - device_param->hip_stream = NULL; - - device_param->hip_module = NULL; - device_param->hip_module_mp = NULL; - device_param->hip_module_amp = NULL; - device_param->hip_module_shared = NULL; + device_param->hip_function_decompress = NULL; + device_param->hip_function_aux1 = NULL; + device_param->hip_function_aux2 = NULL; + device_param->hip_function_aux3 = NULL; + device_param->hip_function_aux4 = NULL; } #if defined (__APPLE__) if (device_param->is_metal == true) { - if (device_param->metal_d_pws_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_pws_buf); - if (device_param->metal_d_pws_amp_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_pws_amp_buf); - if (device_param->metal_d_pws_comp_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_pws_comp_buf); - if (device_param->metal_d_pws_idx) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_pws_idx); - if (device_param->metal_d_rules) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_rules); - if (device_param->metal_d_rules_c) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_rules_c); - if (device_param->metal_d_combs) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_combs); - if (device_param->metal_d_combs_c) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_combs_c); - if (device_param->metal_d_bfs) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bfs); - if (device_param->metal_d_bfs_c) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bfs_c); - if (device_param->metal_d_bitmap_s1_a) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s1_a); - if (device_param->metal_d_bitmap_s1_b) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s1_b); - if (device_param->metal_d_bitmap_s1_c) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s1_c); - if (device_param->metal_d_bitmap_s1_d) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s1_d); - if (device_param->metal_d_bitmap_s2_a) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s2_a); - if (device_param->metal_d_bitmap_s2_b) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s2_b); - if (device_param->metal_d_bitmap_s2_c) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s2_c); - if (device_param->metal_d_bitmap_s2_d) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_bitmap_s2_d); - if (device_param->metal_d_plain_bufs) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_plain_bufs); - if (device_param->metal_d_digests_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_digests_buf); - if (device_param->metal_d_digests_shown) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_digests_shown); - if (device_param->metal_d_salt_bufs) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_salt_bufs); - if (device_param->metal_d_esalt_bufs) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_esalt_bufs); - if (device_param->metal_d_tmps) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_tmps); - if (device_param->metal_d_hooks) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_hooks); - if (device_param->metal_d_result) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_result); - if (device_param->metal_d_extra0_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_extra0_buf); - if (device_param->metal_d_extra1_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_extra1_buf); - if (device_param->metal_d_extra2_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_extra2_buf); - if (device_param->metal_d_extra3_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_extra3_buf); - if (device_param->metal_d_root_css_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_root_css_buf); - if (device_param->metal_d_markov_css_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_markov_css_buf); - if (device_param->metal_d_tm_c) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_tm_c); - if (device_param->metal_d_st_digests_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_st_digests_buf); - if (device_param->metal_d_st_salts_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_st_salts_buf); - if (device_param->metal_d_st_esalts_buf) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_st_esalts_buf); - if (device_param->metal_d_kernel_param) hc_mtlReleaseMemObject (hashcat_ctx, device_param->metal_d_kernel_param); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_pws_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_pws_amp_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_pws_comp_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_pws_idx); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_rules); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_rules_c); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_combs); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_combs_c); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bfs); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bfs_c); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s1_a); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s1_b); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s1_c); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s1_d); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s2_a); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s2_b); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s2_c); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_bitmap_s2_d); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_plain_bufs); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_digests_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_digests_shown); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_salt_bufs); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_esalt_bufs); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_tmps); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_hooks); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_result); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_extra0_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_extra1_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_extra2_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_extra3_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_root_css_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_markov_css_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_tm_c); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_st_digests_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_st_salts_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_st_esalts_buf); + hc_mtlReleaseMemObject (hashcat_ctx, &device_param->metal_d_kernel_param); - if (device_param->metal_function1) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function1); - if (device_param->metal_function12) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function12); - if (device_param->metal_function2p) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function2p); - if (device_param->metal_function2) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function2); - if (device_param->metal_function2e) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function2e); - if (device_param->metal_function23) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function23); - if (device_param->metal_function3) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function3); - if (device_param->metal_function4) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function4); - if (device_param->metal_function_init2) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_init2); - if (device_param->metal_function_loop2p) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_loop2p); - if (device_param->metal_function_loop2) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_loop2); - if (device_param->metal_function_mp) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_mp); - if (device_param->metal_function_mp_l) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_mp_l); - if (device_param->metal_function_mp_r) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_mp_r); - if (device_param->metal_function_tm) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_tm); - if (device_param->metal_function_amp) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_amp); - if (device_param->metal_function_memset) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_memset); - if (device_param->metal_function_bzero) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_bzero); - if (device_param->metal_function_atinit) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_atinit); - if (device_param->metal_function_utf8toutf16le) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_utf8toutf16le); - if (device_param->metal_function_decompress) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_decompress); - if (device_param->metal_function_aux1) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_aux1); - if (device_param->metal_function_aux2) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_aux2); - if (device_param->metal_function_aux3) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_aux3); - if (device_param->metal_function_aux4) hc_mtlReleaseFunction (hashcat_ctx, device_param->metal_function_aux4); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function1); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function12); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function2p); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function2); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function2e); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function23); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function3); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function4); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_init2); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_loop2p); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_loop2); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_mp); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_mp_l); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_mp_r); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_tm); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_amp); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_memset); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_bzero); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_atinit); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_utf8toutf16le); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_decompress); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_aux1); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_aux2); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_aux3); + hc_mtlReleaseFunction (hashcat_ctx, &device_param->metal_function_aux4); - if (device_param->metal_library) hc_mtlReleaseLibrary (hashcat_ctx, device_param->metal_library); - if (device_param->metal_library_mp) hc_mtlReleaseLibrary (hashcat_ctx, device_param->metal_library_mp); - if (device_param->metal_library_amp) hc_mtlReleaseLibrary (hashcat_ctx, device_param->metal_library_amp); - if (device_param->metal_library_shared) hc_mtlReleaseLibrary (hashcat_ctx, device_param->metal_library_shared); + hc_mtlReleaseLibrary (hashcat_ctx, &device_param->metal_library); + hc_mtlReleaseLibrary (hashcat_ctx, &device_param->metal_library_mp); + hc_mtlReleaseLibrary (hashcat_ctx, &device_param->metal_library_amp); + hc_mtlReleaseLibrary (hashcat_ctx, &device_param->metal_library_shared); - //if (device_param->metal_command_queue) hc_mtlReleaseCommandQueue (hashcat_ctx, device_param->metal_command_queue); + //if (device_param->metal_command_queue) hc_mtlReleaseCommandQueue (hashcat_ctx, device_param->metal_command_queue); + //if (device_param->metal_device) hc_mtlReleaseDevice (hashcat_ctx, device_param->metal_device); - //if (device_param->metal_device) hc_mtlReleaseDevice (hashcat_ctx, device_param->metal_device); - - device_param->metal_d_pws_buf = NULL; - device_param->metal_d_pws_amp_buf = NULL; - device_param->metal_d_pws_comp_buf = NULL; - device_param->metal_d_pws_idx = NULL; - device_param->metal_d_rules = NULL; - device_param->metal_d_rules_c = NULL; - device_param->metal_d_combs = NULL; - device_param->metal_d_combs_c = NULL; - device_param->metal_d_bfs = NULL; - device_param->metal_d_bfs_c = NULL; - device_param->metal_d_bitmap_s1_a = NULL; - device_param->metal_d_bitmap_s1_b = NULL; - device_param->metal_d_bitmap_s1_c = NULL; - device_param->metal_d_bitmap_s1_d = NULL; - device_param->metal_d_bitmap_s2_a = NULL; - device_param->metal_d_bitmap_s2_b = NULL; - device_param->metal_d_bitmap_s2_c = NULL; - device_param->metal_d_bitmap_s2_d = NULL; - device_param->metal_d_plain_bufs = NULL; - device_param->metal_d_digests_buf = NULL; - device_param->metal_d_digests_shown = NULL; - device_param->metal_d_salt_bufs = NULL; - device_param->metal_d_esalt_bufs = NULL; - device_param->metal_d_tmps = NULL; - device_param->metal_d_hooks = NULL; - device_param->metal_d_result = NULL; - device_param->metal_d_extra0_buf = NULL; - device_param->metal_d_extra1_buf = NULL; - device_param->metal_d_extra2_buf = NULL; - device_param->metal_d_extra3_buf = NULL; - device_param->metal_d_root_css_buf = NULL; - device_param->metal_d_markov_css_buf = NULL; - device_param->metal_d_tm_c = NULL; - device_param->metal_d_st_digests_buf = NULL; - device_param->metal_d_st_salts_buf = NULL; - device_param->metal_d_st_esalts_buf = NULL; - device_param->metal_d_kernel_param = NULL; - device_param->metal_function1 = NULL; - device_param->metal_function12 = NULL; - device_param->metal_function2p = NULL; - device_param->metal_function2 = NULL; - device_param->metal_function2e = NULL; - device_param->metal_function23 = NULL; - device_param->metal_function3 = NULL; - device_param->metal_function4 = NULL; - device_param->metal_function_init2 = NULL; - device_param->metal_function_loop2p = NULL; - device_param->metal_function_loop2 = NULL; - device_param->metal_function_mp = NULL; - device_param->metal_function_mp_l = NULL; - device_param->metal_function_mp_r = NULL; - device_param->metal_function_tm = NULL; - device_param->metal_function_amp = NULL; - device_param->metal_function_memset = NULL; - device_param->metal_function_bzero = NULL; - device_param->metal_function_atinit = NULL; - device_param->metal_function_utf8toutf16le = NULL; - device_param->metal_function_decompress = NULL; - device_param->metal_function_aux1 = NULL; - device_param->metal_function_aux2 = NULL; - device_param->metal_function_aux3 = NULL; - device_param->metal_function_aux4 = NULL; - device_param->metal_library = NULL; - device_param->metal_library_mp = NULL; - device_param->metal_library_amp = NULL; - device_param->metal_library_shared = NULL; - //device_param->metal_command_queue = NULL; - //device_param->metal_device = NULL; + //device_param->metal_command_queue = NULL; + //device_param->metal_device = NULL; } #endif // __APPLE__ if (device_param->is_opencl == true) { - if (device_param->opencl_d_pws_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_pws_buf); - if (device_param->opencl_d_pws_amp_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_pws_amp_buf); - if (device_param->opencl_d_pws_comp_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_pws_comp_buf); - if (device_param->opencl_d_pws_idx) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_pws_idx); - if (device_param->opencl_d_rules) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_rules); - if (device_param->opencl_d_rules_c) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_rules_c); - if (device_param->opencl_d_combs) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_combs); - if (device_param->opencl_d_combs_c) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_combs_c); - if (device_param->opencl_d_bfs) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bfs); - if (device_param->opencl_d_bfs_c) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bfs_c); - if (device_param->opencl_d_bitmap_s1_a) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s1_a); - if (device_param->opencl_d_bitmap_s1_b) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s1_b); - if (device_param->opencl_d_bitmap_s1_c) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s1_c); - if (device_param->opencl_d_bitmap_s1_d) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s1_d); - if (device_param->opencl_d_bitmap_s2_a) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s2_a); - if (device_param->opencl_d_bitmap_s2_b) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s2_b); - if (device_param->opencl_d_bitmap_s2_c) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s2_c); - if (device_param->opencl_d_bitmap_s2_d) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_bitmap_s2_d); - if (device_param->opencl_d_plain_bufs) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_plain_bufs); - if (device_param->opencl_d_digests_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_digests_buf); - if (device_param->opencl_d_digests_shown) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_digests_shown); - if (device_param->opencl_d_salt_bufs) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_salt_bufs); - if (device_param->opencl_d_esalt_bufs) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_esalt_bufs); - if (device_param->opencl_d_tmps) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_tmps); - if (device_param->opencl_d_hooks) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_hooks); - if (device_param->opencl_d_result) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_result); - if (device_param->opencl_d_extra0_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_extra0_buf); - if (device_param->opencl_d_extra1_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_extra1_buf); - if (device_param->opencl_d_extra2_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_extra2_buf); - if (device_param->opencl_d_extra3_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_extra3_buf); - if (device_param->opencl_d_root_css_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_root_css_buf); - if (device_param->opencl_d_markov_css_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_markov_css_buf); - if (device_param->opencl_d_tm_c) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_tm_c); - if (device_param->opencl_d_st_digests_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_st_digests_buf); - if (device_param->opencl_d_st_salts_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_st_salts_buf); - if (device_param->opencl_d_st_esalts_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_st_esalts_buf); - if (device_param->opencl_d_kernel_param) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_kernel_param); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_pws_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_pws_amp_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_pws_comp_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_pws_idx); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_rules); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_rules_c); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_combs); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_combs_c); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bfs); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bfs_c); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s1_a); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s1_b); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s1_c); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s1_d); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s2_a); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s2_b); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s2_c); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_bitmap_s2_d); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_plain_bufs); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_digests_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_digests_shown); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_salt_bufs); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_esalt_bufs); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_tmps); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_hooks); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_result); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_extra0_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_extra1_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_extra2_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_extra3_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_root_css_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_markov_css_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_tm_c); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_st_digests_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_st_salts_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_st_esalts_buf); + hc_clReleaseMemObjectPtr (hashcat_ctx, &device_param->opencl_d_kernel_param); - if (device_param->opencl_kernel1) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel1); - if (device_param->opencl_kernel12) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel12); - if (device_param->opencl_kernel2p) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel2p); - if (device_param->opencl_kernel2) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel2); - if (device_param->opencl_kernel2e) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel2e); - if (device_param->opencl_kernel23) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel23); - if (device_param->opencl_kernel3) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel3); - if (device_param->opencl_kernel4) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel4); - if (device_param->opencl_kernel_init2) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_init2); - if (device_param->opencl_kernel_loop2p) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_loop2p); - if (device_param->opencl_kernel_loop2) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_loop2); - if (device_param->opencl_kernel_mp) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_mp); - if (device_param->opencl_kernel_mp_l) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_mp_l); - if (device_param->opencl_kernel_mp_r) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_mp_r); - if (device_param->opencl_kernel_tm) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_tm); - if (device_param->opencl_kernel_amp) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_amp); - if (device_param->opencl_kernel_memset) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_memset); - if (device_param->opencl_kernel_bzero) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_bzero); - if (device_param->opencl_kernel_atinit) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_atinit); - if (device_param->opencl_kernel_utf8toutf16le) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_utf8toutf16le); - if (device_param->opencl_kernel_decompress)hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_decompress); - if (device_param->opencl_kernel_aux1) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_aux1); - if (device_param->opencl_kernel_aux2) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_aux2); - if (device_param->opencl_kernel_aux3) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_aux3); - if (device_param->opencl_kernel_aux4) hc_clReleaseKernel (hashcat_ctx, device_param->opencl_kernel_aux4); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel1); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel12); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel2p); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel2); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel2e); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel23); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel3); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel4); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_init2); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_loop2p); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_loop2); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_mp); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_mp_l); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_mp_r); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_tm); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_amp); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_memset); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_bzero); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_atinit); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_utf8toutf16le); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_decompress); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_aux1); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_aux2); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_aux3); + hc_clReleaseKernelPtr (hashcat_ctx, &device_param->opencl_kernel_aux4); - if (device_param->opencl_program) hc_clReleaseProgram (hashcat_ctx, device_param->opencl_program); - if (device_param->opencl_program_mp) hc_clReleaseProgram (hashcat_ctx, device_param->opencl_program_mp); - if (device_param->opencl_program_amp) hc_clReleaseProgram (hashcat_ctx, device_param->opencl_program_amp); - if (device_param->opencl_program_shared) hc_clReleaseProgram (hashcat_ctx, device_param->opencl_program_shared); + hc_clReleaseProgramPtr (hashcat_ctx, &device_param->opencl_program); + hc_clReleaseProgramPtr (hashcat_ctx, &device_param->opencl_program_mp); + hc_clReleaseProgramPtr (hashcat_ctx, &device_param->opencl_program_amp); + hc_clReleaseProgramPtr (hashcat_ctx, &device_param->opencl_program_shared); - //if (device_param->opencl_command_queue) hc_clReleaseCommandQueue (hashcat_ctx, device_param->opencl_command_queue); + //if (device_param->opencl_command_queue) hc_clReleaseCommandQueue (hashcat_ctx, device_param->opencl_command_queue); + //if (device_param->opencl_context) hc_clReleaseContext (hashcat_ctx, device_param->opencl_context); - //if (device_param->opencl_context) hc_clReleaseContext (hashcat_ctx, device_param->opencl_context); - - device_param->opencl_d_pws_buf = NULL; - device_param->opencl_d_pws_amp_buf = NULL; - device_param->opencl_d_pws_comp_buf = NULL; - device_param->opencl_d_pws_idx = NULL; - device_param->opencl_d_rules = NULL; - device_param->opencl_d_rules_c = NULL; - device_param->opencl_d_combs = NULL; - device_param->opencl_d_combs_c = NULL; - device_param->opencl_d_bfs = NULL; - device_param->opencl_d_bfs_c = NULL; - device_param->opencl_d_bitmap_s1_a = NULL; - device_param->opencl_d_bitmap_s1_b = NULL; - device_param->opencl_d_bitmap_s1_c = NULL; - device_param->opencl_d_bitmap_s1_d = NULL; - device_param->opencl_d_bitmap_s2_a = NULL; - device_param->opencl_d_bitmap_s2_b = NULL; - device_param->opencl_d_bitmap_s2_c = NULL; - device_param->opencl_d_bitmap_s2_d = NULL; - device_param->opencl_d_plain_bufs = NULL; - device_param->opencl_d_digests_buf = NULL; - device_param->opencl_d_digests_shown = NULL; - device_param->opencl_d_salt_bufs = NULL; - device_param->opencl_d_esalt_bufs = NULL; - device_param->opencl_d_tmps = NULL; - device_param->opencl_d_hooks = NULL; - device_param->opencl_d_result = NULL; - device_param->opencl_d_extra0_buf = NULL; - device_param->opencl_d_extra1_buf = NULL; - device_param->opencl_d_extra2_buf = NULL; - device_param->opencl_d_extra3_buf = NULL; - device_param->opencl_d_root_css_buf = NULL; - device_param->opencl_d_markov_css_buf = NULL; - device_param->opencl_d_tm_c = NULL; - device_param->opencl_d_st_digests_buf = NULL; - device_param->opencl_d_st_salts_buf = NULL; - device_param->opencl_d_st_esalts_buf = NULL; - device_param->opencl_d_kernel_param = NULL; - device_param->opencl_kernel1 = NULL; - device_param->opencl_kernel12 = NULL; - device_param->opencl_kernel2p = NULL; - device_param->opencl_kernel2 = NULL; - device_param->opencl_kernel2e = NULL; - device_param->opencl_kernel23 = NULL; - device_param->opencl_kernel3 = NULL; - device_param->opencl_kernel4 = NULL; - device_param->opencl_kernel_init2 = NULL; - device_param->opencl_kernel_loop2p = NULL; - device_param->opencl_kernel_loop2 = NULL; - device_param->opencl_kernel_mp = NULL; - device_param->opencl_kernel_mp_l = NULL; - device_param->opencl_kernel_mp_r = NULL; - device_param->opencl_kernel_tm = NULL; - device_param->opencl_kernel_amp = NULL; - device_param->opencl_kernel_memset = NULL; - device_param->opencl_kernel_bzero = NULL; - device_param->opencl_kernel_atinit = NULL; - device_param->opencl_kernel_utf8toutf16le = NULL; - device_param->opencl_kernel_decompress = NULL; - device_param->opencl_kernel_aux1 = NULL; - device_param->opencl_kernel_aux2 = NULL; - device_param->opencl_kernel_aux3 = NULL; - device_param->opencl_kernel_aux4 = NULL; - device_param->opencl_program = NULL; - device_param->opencl_program_mp = NULL; - device_param->opencl_program_amp = NULL; - device_param->opencl_program_shared = NULL; - //device_param->opencl_command_queue = NULL; - //device_param->opencl_context = NULL; + //device_param->opencl_command_queue = NULL; + //device_param->opencl_context = NULL; } device_param->h_tmps = NULL; diff --git a/src/ext_OpenCL.c b/src/ext_OpenCL.c index b935c33aa..25413a611 100644 --- a/src/ext_OpenCL.c +++ b/src/ext_OpenCL.c @@ -249,6 +249,45 @@ void ocl_close (void *hashcat_ctx) } } +int hc_clReleaseMemObjectPtr (void *hashcat_ctx, cl_mem *mem) +{ + int rc = -1; + + if (mem == NULL || *mem == NULL) return rc; + + rc = hc_clReleaseMemObject (hashcat_ctx, *mem); + + *mem = NULL; + + return rc; +} + +int hc_clReleaseKernelPtr (void *hashcat_ctx, cl_kernel *kernel) +{ + int rc = -1; + + if (kernel == NULL || *kernel == NULL) return rc; + + rc = hc_clReleaseKernel (hashcat_ctx, *kernel); + + *kernel = NULL; + + return rc; +} + +int hc_clReleaseProgramPtr (void *hashcat_ctx, cl_program *program) +{ + int rc = -1; + + if (program == NULL || *program == NULL) return rc; + + rc = hc_clReleaseProgram (hashcat_ctx, *program); + + *program = NULL; + + return rc; +} + int hc_clEnqueueNDRangeKernel (void *hashcat_ctx, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; diff --git a/src/ext_cuda.c b/src/ext_cuda.c index 878977ade..9240ad455 100644 --- a/src/ext_cuda.c +++ b/src/ext_cuda.c @@ -144,6 +144,58 @@ void cuda_close (void *hashcat_ctx) } } +int hc_cuEventDestroyPtr (void *hashcat_ctx, CUevent *hEvent) +{ + int rc = -1; + + if (hEvent == NULL || *hEvent == NULL) return rc; + + rc = hc_cuEventDestroy (hashcat_ctx, *hEvent); + + *hEvent = NULL; + + return rc; +} + +int hc_cuMemFreePtr (void *hashcat_ctx, CUdeviceptr *dptr) +{ + int rc = -1; + + if (dptr == NULL || *dptr == 0) return rc; + + rc = hc_cuMemFree (hashcat_ctx, *dptr); + + *dptr = 0; + + return rc; +} + +int hc_cuModuleUnloadPtr (void *hashcat_ctx, CUmodule *hmod) +{ + int rc = -1; + + if (hmod == NULL || *hmod == NULL) return rc; + + rc = hc_cuModuleUnload (hashcat_ctx, *hmod); + + *hmod = NULL; + + return rc; +} + +int hc_cuStreamDestroyPtr (void *hashcat_ctx, CUstream *hStream) +{ + int rc = -1; + + if (hStream == NULL || *hStream == NULL) return rc; + + rc = hc_cuStreamDestroy (hashcat_ctx, *hStream); + + *hStream = NULL; + + return rc; +} + int hc_cuInit (void *hashcat_ctx, unsigned int Flags) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; diff --git a/src/ext_hip.c b/src/ext_hip.c index 9e7eb00ab..a6539dc09 100644 --- a/src/ext_hip.c +++ b/src/ext_hip.c @@ -187,6 +187,58 @@ void hip_close (void *hashcat_ctx) } } +int hc_hipEventDestroyPtr (void *hashcat_ctx, hipEvent_t *hEvent) +{ + int rc = -1; + + if (hEvent == NULL || *hEvent == NULL) return rc; + + rc = hc_hipEventDestroy (hashcat_ctx, *hEvent); + + *hEvent = NULL; + + return rc; +} + +int hc_hipMemFreePtr (void *hashcat_ctx, hipDeviceptr_t *dptr) +{ + int rc = -1; + + if (dptr == NULL || *dptr == NULL) return rc; + + rc = hc_hipMemFree (hashcat_ctx, *dptr); + + *dptr = 0; + + return rc; +} + +int hc_hipModuleUnloadPtr (void *hashcat_ctx, hipModule_t *hmod) +{ + int rc = -1; + + if (hmod == NULL || *hmod == NULL) return rc; + + rc = hc_hipModuleUnload (hashcat_ctx, *hmod); + + *hmod = NULL; + + return rc; +} + +int hc_hipStreamDestroyPtr (void *hashcat_ctx, hipStream_t *hStream) +{ + int rc = -1; + + if (hStream == NULL || *hStream == NULL) return rc; + + rc = hc_hipStreamDestroy (hashcat_ctx, *hStream); + + *hStream = NULL; + + return rc; +} + int hc_hipCtxCreate (void *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; diff --git a/src/ext_metal.m b/src/ext_metal.m index 091b13e3c..1bf348798 100644 --- a/src/ext_metal.m +++ b/src/ext_metal.m @@ -867,7 +867,7 @@ int hc_mtlCreateBuffer (void *hashcat_ctx, mtl_device_id metal_device, size_t si return 0; } -int hc_mtlReleaseMemObject (void *hashcat_ctx, mtl_mem metal_buffer) +int hc_mtlReleaseMemObject (void *hashcat_ctx, mtl_mem *metal_buffer) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; @@ -875,44 +875,46 @@ int hc_mtlReleaseMemObject (void *hashcat_ctx, mtl_mem metal_buffer) if (mtl == NULL) return -1; - if (metal_buffer == nil) - { - event_log_error (hashcat_ctx, "%s(): invalid metal buffer", __func__); + if (metal_buffer == NULL || *metal_buffer == nil) return -1; - return -1; - } + [*metal_buffer setPurgeableState:MTLPurgeableStateEmpty]; + [*metal_buffer release]; - [metal_buffer setPurgeableState:MTLPurgeableStateEmpty]; - [metal_buffer release]; + *metal_buffer = nil; return 0; } -int hc_mtlReleaseFunction (void *hashcat_ctx, mtl_function metal_function) +int hc_mtlReleaseFunction (void *hashcat_ctx, mtl_function *metal_function) { - if (metal_function == nil) - { - event_log_error (hashcat_ctx, "%s(): invalid metal function", __func__); + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; - return -1; - } + MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl; - [metal_function release]; + if (mtl == NULL) return -1; + + if (metal_function == NULL || *metal_function == nil) return -1; + + [*metal_function release]; + + *metal_function = nil; return 0; } -int hc_mtlReleaseLibrary (void *hashcat_ctx, mtl_library metal_library) +int hc_mtlReleaseLibrary (void *hashcat_ctx, mtl_library *metal_library) { - if (metal_library == nil) - { - event_log_error (hashcat_ctx, "%s(): invalid metal library", __func__); + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; - return -1; - } + MTL_PTR *mtl = (MTL_PTR *) backend_ctx->mtl; - [metal_library release]; - metal_library = nil; + if (mtl == NULL) return -1; + + if (metal_library == NULL || *metal_library == nil) return -1; + + [*metal_library release]; + + *metal_library = nil; return 0; } @@ -927,6 +929,7 @@ int hc_mtlReleaseCommandQueue (void *hashcat_ctx, mtl_command_queue command_queu } [command_queue release]; + command_queue = nil; return 0; @@ -942,6 +945,7 @@ int hc_mtlReleaseDevice (void *hashcat_ctx, mtl_device_id metal_device) } [metal_device release]; + metal_device = nil; return 0; @@ -995,14 +999,14 @@ int hc_mtlMemcpyDtoD (void *hashcat_ctx, mtl_command_queue command_queue, mtl_me if (buf_src_off + buf_size > [buf_src length]) { - event_log_error(hashcat_ctx, "%s(): src buffer offset + size out of bounds", __func__); + event_log_error (hashcat_ctx, "%s(): src buffer offset + size out of bounds", __func__); return -1; } if (buf_dst_off + buf_size > [buf_dst length]) { - event_log_error(hashcat_ctx, "%s(): dst buffer offset + size out of bounds", __func__); + event_log_error (hashcat_ctx, "%s(): dst buffer offset + size out of bounds", __func__); return -1; } @@ -1081,7 +1085,7 @@ int hc_mtlMemcpyHtoD (void *hashcat_ctx, mtl_command_queue command_queue, mtl_me if (buf_dst_off + buf_size > [buf_dst length]) { - event_log_error(hashcat_ctx, "%s(): buffer offset + size out of bounds", __func__); + event_log_error (hashcat_ctx, "%s(): buffer offset + size out of bounds", __func__); return -1; }