From d9434894049ef6ad41920a7e391e6210029be0fa Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sat, 19 Jul 2025 01:03:43 +0200 Subject: [PATCH 1/4] Update core in the following functions I changed the type for the parameter used to specify the target of the operation: - hc_clReleaseMemObject - hc_clReleaseKernel - hc_clReleaseProgram - hc_cuModuleUnload - hc_cuMemFree - hc_cuStreamDestroy - hc_cuEventDestroy - hc_hipEventDestroy - hc_hipMemFree - hc_hipModuleUnload - hc_hipStreamDestroy - hc_mtlReleaseMemObject - hc_mtlReleaseFunction - hc_mtlReleaseLibrary With this change, it was possible to remove several lines of code from backend.c, making it more readable. --- include/ext_OpenCL.h | 6 +- include/ext_cuda.h | 8 +- include/ext_hip.h | 8 +- include/ext_metal.h | 6 +- src/backend.c | 761 ++++++++++++++----------------------------- src/ext_OpenCL.c | 24 +- src/ext_cuda.c | 32 +- src/ext_hip.c | 32 +- src/ext_metal.m | 50 +-- src/memory.c | 2 + 10 files changed, 347 insertions(+), 582 deletions(-) diff --git a/include/ext_OpenCL.h b/include/ext_OpenCL.h index 7044ac667..36e31af4e 100644 --- a/include/ext_OpenCL.h +++ b/include/ext_OpenCL.h @@ -148,9 +148,9 @@ int hc_clBuildProgram (void *hashcat_ctx, cl_program program, cl_uint int hc_clCompileProgram (void *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, cl_uint num_input_headers, const cl_program *input_headers, const char **header_include_names, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data); int hc_clLinkProgram (void *hashcat_ctx, cl_context context, cl_uint num_devices, const cl_device_id *device_list, const char *options, cl_uint num_input_programs, const cl_program *input_programs, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data, cl_program *program); int hc_clCreateKernel (void *hashcat_ctx, cl_program program, const char *kernel_name, cl_kernel *kernel); -int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem mem); -int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel kernel); -int hc_clReleaseProgram (void *hashcat_ctx, cl_program program); +int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem *mem); +int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel *kernel); +int hc_clReleaseProgram (void *hashcat_ctx, cl_program *program); int hc_clReleaseCommandQueue (void *hashcat_ctx, cl_command_queue command_queue); int hc_clReleaseContext (void *hashcat_ctx, cl_context context); int hc_clEnqueueMapBuffer (void *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_map, cl_map_flags map_flags, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event, void **buf); diff --git a/include/ext_cuda.h b/include/ext_cuda.h index a80af5367..51cf3540f 100644 --- a/include/ext_cuda.h +++ b/include/ext_cuda.h @@ -1272,7 +1272,7 @@ int hc_cuDeviceGetName (void *hashcat_ctx, char *name, int len, CUdevice int hc_cuDeviceTotalMem (void *hashcat_ctx, size_t *bytes, CUdevice dev); int hc_cuDriverGetVersion (void *hashcat_ctx, int *driverVersion); int hc_cuEventCreate (void *hashcat_ctx, CUevent *phEvent, unsigned int Flags); -int hc_cuEventDestroy (void *hashcat_ctx, CUevent hEvent); +int hc_cuEventDestroy (void *hashcat_ctx, CUevent *hEvent); int hc_cuEventElapsedTime (void *hashcat_ctx, float *pMilliseconds, CUevent hStart, CUevent hEnd); int hc_cuEventQuery (void *hashcat_ctx, CUevent hEvent); int hc_cuEventRecord (void *hashcat_ctx, CUevent hEvent, CUstream hStream); @@ -1292,14 +1292,14 @@ int hc_cuMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, CUdeviceptr sr int hc_cuMemcpyHtoDAsync (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream); int hc_cuMemsetD32Async (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N, CUstream hStream); int hc_cuMemsetD8Async (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N, CUstream hStream); -int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr); +int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr *dptr); int hc_cuMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); int hc_cuModuleGetFunction (void *hashcat_ctx, CUfunction *hfunc, CUmodule hmod, const char *name); int hc_cuModuleGetGlobal (void *hashcat_ctx, CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name); int hc_cuModuleLoadDataEx (void *hashcat_ctx, CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues); -int hc_cuModuleUnload (void *hashcat_ctx, CUmodule hmod); +int hc_cuModuleUnload (void *hashcat_ctx, CUmodule *hmod); int hc_cuStreamCreate (void *hashcat_ctx, CUstream *phStream, unsigned int Flags); -int hc_cuStreamDestroy (void *hashcat_ctx, CUstream hStream); +int hc_cuStreamDestroy (void *hashcat_ctx, CUstream *hStream); int hc_cuStreamSynchronize (void *hashcat_ctx, CUstream hStream); int hc_cuCtxPushCurrent (void *hashcat_ctx, CUcontext ctx); int hc_cuCtxPopCurrent (void *hashcat_ctx, CUcontext *pctx); diff --git a/include/ext_hip.h b/include/ext_hip.h index 3c0b8433a..4b8bd678b 100644 --- a/include/ext_hip.h +++ b/include/ext_hip.h @@ -670,7 +670,7 @@ int hc_hipDeviceTotalMem (void *hashcat_ctx, size_t *bytes, hipDevice_t d int hc_hipDriverGetVersion (void *hashcat_ctx, int *driverVersion); int hc_hipEventCreate (void *hashcat_ctx, hipEvent_t *phEvent); int hc_hipEventCreateWithFlags (void *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags); -int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t hEvent); +int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t *hEvent); int hc_hipEventElapsedTime (void *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd); int hc_hipEventQuery (void *hashcat_ctx, hipEvent_t hEvent); int hc_hipEventRecord (void *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream); @@ -679,7 +679,7 @@ int hc_hipFuncGetAttribute (void *hashcat_ctx, int *pi, hipFunction_attribu int hc_hipInit (void *hashcat_ctx, unsigned int Flags); int hc_hipLaunchKernel (void *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra); int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize); -int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t dptr); +int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t *dptr); int hc_hipMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); int hc_hipMemcpyDtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount); int hc_hipMemcpyDtoH (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount); @@ -694,13 +694,13 @@ int hc_hipMemsetD8Async (void *hashcat_ctx, hipDeviceptr_t dstDevice, un int hc_hipModuleGetFunction (void *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name); int hc_hipModuleGetGlobal (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name); int hc_hipModuleLoadDataEx (void *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues); -int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t hmod); +int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t *hmod); int hc_hipRuntimeGetVersion (void *hashcat_ctx, int *runtimeVersion); int hc_hipSetDevice (void *hashcat_ctx, hipDevice_t dev); int hc_hipSetDeviceFlags (void *hashcat_ctx, unsigned int flags); int hc_hipStreamCreate (void *hashcat_ctx, hipStream_t *phStream); int hc_hipStreamCreateWithFlags (void *hashcat_ctx, hipStream_t *phStream, unsigned int flags); -int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t hStream); +int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t *hStream); int hc_hipStreamSynchronize (void *hashcat_ctx, hipStream_t hStream); int hc_hipGetDeviceProperties (void *hashcat_ctx, hipDeviceProp_t *prop, hipDevice_t dev); int hc_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor (void *hashcat_ctx, int *numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk); 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..45c1df35a 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_clReleaseProgram (hashcat_ctx, &program); return false; } - if (hc_clReleaseProgram (hashcat_ctx, program) == -1) return false; + if (hc_clReleaseProgram (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_clReleaseMemObject (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_clReleaseProgram (hashcat_ctx, &p1); } if (cache_disable == false) @@ -17004,552 +17004,267 @@ 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_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_amp_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_comp_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_idx); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_rules); + //hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_rules_c); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_combs); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_combs_c); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bfs); + //hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bfs_c); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_a); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_b); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_c); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_d); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_a); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_b); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_c); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_d); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_plain_bufs); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_digests_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_digests_shown); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_salt_bufs); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_esalt_bufs); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_tmps); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_hooks); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_result); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra0_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra1_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra2_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra3_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_root_css_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_markov_css_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_tm_c); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_st_digests_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_st_salts_buf); + hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_st_esalts_buf); + hc_cuMemFree (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_cuEventDestroy (hashcat_ctx, &device_param->cuda_event1); + hc_cuEventDestroy (hashcat_ctx, &device_param->cuda_event2); + hc_cuEventDestroy (hashcat_ctx, &device_param->cuda_event3); - if (device_param->cuda_stream) hc_cuStreamDestroy (hashcat_ctx, device_param->cuda_stream); + hc_cuStreamDestroy (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_cuModuleUnload (hashcat_ctx, &device_param->cuda_module); + hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module_mp); + hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module_amp); + hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module_shared); - //if (device_param->cuda_context) hc_cuCtxDestroy (hashcat_ctx, device_param->cuda_context); - - 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_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_context = 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_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_amp_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_comp_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_idx); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_rules); + //hc_hipMemFree (hashcat_ctx, &device_param->hip_d_rules_c); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_combs); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_combs_c); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bfs); + //hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bfs_c); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_a); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_b); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_c); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_d); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_a); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_b); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_c); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_d); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_plain_bufs); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_digests_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_digests_shown); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_salt_bufs); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_esalt_bufs); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_tmps); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_hooks); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_result); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra0_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra1_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra2_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra3_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_root_css_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_markov_css_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_tm_c); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_st_digests_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_st_salts_buf); + hc_hipMemFree (hashcat_ctx, &device_param->hip_d_st_esalts_buf); + hc_hipMemFree (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_hipEventDestroy (hashcat_ctx, &device_param->hip_event1); + hc_hipEventDestroy (hashcat_ctx, &device_param->hip_event2); + hc_hipEventDestroy (hashcat_ctx, &device_param->hip_event3); - if (device_param->hip_stream) hc_hipStreamDestroy (hashcat_ctx, device_param->hip_stream); + hc_hipStreamDestroy (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); - - 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_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; + hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module); + hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_mp); + hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_amp); + hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_shared); } #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_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_amp_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_comp_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_idx); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_rules); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_rules_c); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_combs); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_combs_c); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bfs); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bfs_c); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_a); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_b); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_c); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_d); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_a); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_b); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_c); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_d); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_plain_bufs); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_digests_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_digests_shown); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_salt_bufs); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_esalt_bufs); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_tmps); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_hooks); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_result); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra0_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra1_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra2_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra3_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_root_css_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_markov_css_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_tm_c); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_st_digests_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_st_salts_buf); + hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_st_esalts_buf); + hc_clReleaseMemObject (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_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel1); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel12); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel2p); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel2); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel2e); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel23); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel3); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel4); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_init2); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_loop2p); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_loop2); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_mp); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_mp_l); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_mp_r); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_tm); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_amp); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_memset); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_bzero); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_atinit); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_utf8toutf16le); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_decompress); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_aux1); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_aux2); + hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_aux3); + hc_clReleaseKernel (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_clReleaseProgram (hashcat_ctx, &device_param->opencl_program); + hc_clReleaseProgram (hashcat_ctx, &device_param->opencl_program_mp); + hc_clReleaseProgram (hashcat_ctx, &device_param->opencl_program_amp); + hc_clReleaseProgram (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..d86a5b1ba 100644 --- a/src/ext_OpenCL.c +++ b/src/ext_OpenCL.c @@ -663,13 +663,17 @@ int hc_clCreateKernel (void *hashcat_ctx, cl_program program, const char *kernel return 0; } -int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem mem) +int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem *mem) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; OCL_PTR *ocl = (OCL_PTR *) backend_ctx->ocl; - const cl_int CL_err = ocl->clReleaseMemObject (mem); + if (mem == NULL || *mem == NULL) return -1; + + const cl_int CL_err = ocl->clReleaseMemObject (*mem); + + *mem = NULL; if (CL_err != CL_SUCCESS) { @@ -681,13 +685,17 @@ int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem mem) return 0; } -int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel kernel) +int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel *kernel) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; OCL_PTR *ocl = (OCL_PTR *) backend_ctx->ocl; - const cl_int CL_err = ocl->clReleaseKernel (kernel); + if (kernel == NULL || *kernel == NULL) return -1; + + const cl_int CL_err = ocl->clReleaseKernel (*kernel); + + *kernel = NULL; if (CL_err != CL_SUCCESS) { @@ -699,13 +707,17 @@ int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel kernel) return 0; } -int hc_clReleaseProgram (void *hashcat_ctx, cl_program program) +int hc_clReleaseProgram (void *hashcat_ctx, cl_program *program) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; OCL_PTR *ocl = (OCL_PTR *) backend_ctx->ocl; - const cl_int CL_err = ocl->clReleaseProgram (program); + if (program == NULL || *program == NULL) return -1; + + const cl_int CL_err = ocl->clReleaseProgram (*program); + + *program = NULL; if (CL_err != CL_SUCCESS) { diff --git a/src/ext_cuda.c b/src/ext_cuda.c index 878977ade..88ab2ee1d 100644 --- a/src/ext_cuda.c +++ b/src/ext_cuda.c @@ -414,13 +414,17 @@ int hc_cuModuleLoadDataEx (void *hashcat_ctx, CUmodule *module, const void *imag return 0; } -int hc_cuModuleUnload (void *hashcat_ctx, CUmodule hmod) +int hc_cuModuleUnload (void *hashcat_ctx, CUmodule *hmod) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - const CUresult CU_err = cuda->cuModuleUnload (hmod); + if (hmod == NULL || *hmod == NULL) return -1; + + const CUresult CU_err = cuda->cuModuleUnload (*hmod); + + *hmod = NULL; if (CU_err != CUDA_SUCCESS) { @@ -495,13 +499,17 @@ int hc_cuMemAlloc (void *hashcat_ctx, CUdeviceptr *dptr, size_t bytesize) return 0; } -int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr) +int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr *dptr) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - const CUresult CU_err = cuda->cuMemFree (dptr); + if (dptr == NULL || *dptr == 0) return -1; + + const CUresult CU_err = cuda->cuMemFree (*dptr); + + *dptr = 0; if (CU_err != CUDA_SUCCESS) { @@ -955,13 +963,17 @@ int hc_cuStreamCreate (void *hashcat_ctx, CUstream *phStream, unsigned int Flags return 0; } -int hc_cuStreamDestroy (void *hashcat_ctx, CUstream hStream) +int hc_cuStreamDestroy (void *hashcat_ctx, CUstream *hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - const CUresult CU_err = cuda->cuStreamDestroy (hStream); + if (hStream == NULL || *hStream == NULL) return -1; + + const CUresult CU_err = cuda->cuStreamDestroy (*hStream); + + *hStream = NULL; if (CU_err != CUDA_SUCCESS) { @@ -1090,13 +1102,17 @@ int hc_cuEventCreate (void *hashcat_ctx, CUevent *phEvent, unsigned int Flags) return 0; } -int hc_cuEventDestroy (void *hashcat_ctx, CUevent hEvent) +int hc_cuEventDestroy (void *hashcat_ctx, CUevent *hEvent) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - const CUresult CU_err = cuda->cuEventDestroy (hEvent); + if (hEvent == NULL || *hEvent == NULL) return -1; + + const CUresult CU_err = cuda->cuEventDestroy (*hEvent); + + *hEvent = NULL; if (CU_err != CUDA_SUCCESS) { diff --git a/src/ext_hip.c b/src/ext_hip.c index 9e7eb00ab..3d402d235 100644 --- a/src/ext_hip.c +++ b/src/ext_hip.c @@ -565,13 +565,17 @@ int hc_hipEventCreateWithFlags (void *hashcat_ctx, hipEvent_t *phEvent, unsigned return 0; } -int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t hEvent) +int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t *hEvent) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const hipError_t HIP_err = hip->hipEventDestroy (hEvent); + if (hEvent == NULL || *hEvent == NULL) return -1; + + const hipError_t HIP_err = hip->hipEventDestroy (*hEvent); + + *hEvent = NULL; if (HIP_err != hipSuccess) { @@ -781,13 +785,17 @@ int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize) return 0; } -int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t dptr) +int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t *dptr) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const hipError_t HIP_err = hip->hipMemFree (dptr); + if (dptr == NULL || *dptr == NULL) return -1; + + const hipError_t HIP_err = hip->hipMemFree (*dptr); + + *dptr = NULL; if (HIP_err != hipSuccess) { @@ -1188,13 +1196,17 @@ int hc_hipModuleLoadDataEx (void *hashcat_ctx, hipModule_t *module, const void * return 0; } -int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t hmod) +int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t *hmod) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const hipError_t HIP_err = hip->hipModuleUnload (hmod); + if (hmod == NULL || *hmod == NULL) return -1; + + const hipError_t HIP_err = hip->hipModuleUnload (*hmod); + + *hmod = NULL; if (HIP_err != hipSuccess) { @@ -1350,13 +1362,17 @@ int hc_hipStreamCreateWithFlags (void *hashcat_ctx, hipStream_t *phStream, unsig return 0; } -int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t hStream) +int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t *hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const hipError_t HIP_err = hip->hipStreamDestroy (hStream); + if (hStream == NULL || *hStream == NULL) return -1; + + const hipError_t HIP_err = hip->hipStreamDestroy (*hStream); + + *hStream = NULL; if (HIP_err != hipSuccess) { diff --git a/src/ext_metal.m b/src/ext_metal.m index 091b13e3c..dd4e73d1f 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; diff --git a/src/memory.c b/src/memory.c index 15f4cbe19..5c26c3f3e 100644 --- a/src/memory.c +++ b/src/memory.c @@ -47,6 +47,8 @@ void *hcrealloc (void *ptr, const size_t oldsz, const size_t addsz) char *hcstrdup (const char *s) { + if (s == NULL) return (NULL); + const size_t len = strlen (s); char *b = (char *) hcmalloc (len + 1); From a1e811b28281223ff5dbbd012545f6c548347955 Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sat, 19 Jul 2025 02:17:36 +0200 Subject: [PATCH 2/4] manually set CUDA and HIP functions to NULL. Set '-g' to build_options only with HIP and OpenCL --- src/backend.c | 72 +++++++++++++++++++++++++++++++++++++++++++++++---- src/ext_hip.c | 2 +- 2 files changed, 68 insertions(+), 6 deletions(-) diff --git a/src/backend.c b/src/backend.c index 45c1df35a..8fb353b36 100644 --- a/src/backend.c +++ b/src/backend.c @@ -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)) @@ -17053,8 +17057,37 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx) hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module_amp); hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module_shared); - //if (device_param->cuda_context) hc_cuCtxDestroy (hashcat_ctx, device_param->cuda_context); - //device_param->cuda_context = NULL; + device_param->cuda_d_rules_c = 0; + device_param->cuda_d_bfs_c = 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_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; + + //if (device_param->cuda_context) hc_cuCtxDestroy (hashcat_ctx, device_param->cuda_context); + //device_param->cuda_context = NULL; } if (device_param->is_hip == true) @@ -17107,6 +17140,35 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx) hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_mp); hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_amp); hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_shared); + + 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_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; } #if defined (__APPLE__) diff --git a/src/ext_hip.c b/src/ext_hip.c index 3d402d235..00d938d57 100644 --- a/src/ext_hip.c +++ b/src/ext_hip.c @@ -795,7 +795,7 @@ int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t *dptr) const hipError_t HIP_err = hip->hipMemFree (*dptr); - *dptr = NULL; + *dptr = 0; if (HIP_err != hipSuccess) { From 5d3426adff14e79640910b103bd5c3fd67a15e13 Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sat, 19 Jul 2025 02:28:58 +0200 Subject: [PATCH 3/4] Update changes.txt --- docs/changes.txt | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/docs/changes.txt b/docs/changes.txt index 772069262..c22c0029e 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -145,11 +145,16 @@ - 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 the definitions of the following functions to use a pointer to the buffer instead of the value: + (OpenCL) hc_clReleaseMemObject, hc_clReleaseKernel, hc_clReleaseProgram + (CUDA) hc_cuModuleUnload, hc_cuMemFree, hc_cuStreamDestroy, hc_cuEventDestroy + (HIP) hc_hipEventDestroy, hc_hipMemFree, hc_hipModuleUnload, hc_hipStreamDestroy + (Metal) hc_mtlReleaseMemObject, hc_mtlReleaseFunction, hc_mtlReleaseLibrary - 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 From 6878234480a06cfcf668786d591c4f3594217e9e Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sat, 19 Jul 2025 11:35:11 +0200 Subject: [PATCH 4/4] restore API's original functions and move checks to new *Ptr functions. Restore original hcstrdup function. --- docs/changes.txt | 6 +- include/ext_OpenCL.h | 10 +- include/ext_cuda.h | 13 +- include/ext_hip.h | 13 +- src/backend.c | 320 +++++++++++++++++++++---------------------- src/ext_OpenCL.c | 63 ++++++--- src/ext_cuda.c | 84 ++++++++---- src/ext_hip.c | 84 ++++++++---- src/ext_metal.m | 6 +- src/memory.c | 2 - 10 files changed, 354 insertions(+), 247 deletions(-) diff --git a/docs/changes.txt b/docs/changes.txt index c22c0029e..62f0dd60a 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -145,11 +145,7 @@ - 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 the definitions of the following functions to use a pointer to the buffer instead of the value: - (OpenCL) hc_clReleaseMemObject, hc_clReleaseKernel, hc_clReleaseProgram - (CUDA) hc_cuModuleUnload, hc_cuMemFree, hc_cuStreamDestroy, hc_cuEventDestroy - (HIP) hc_hipEventDestroy, hc_hipMemFree, hc_hipModuleUnload, hc_hipStreamDestroy - (Metal) hc_mtlReleaseMemObject, hc_mtlReleaseFunction, hc_mtlReleaseLibrary +- 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 diff --git a/include/ext_OpenCL.h b/include/ext_OpenCL.h index 36e31af4e..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); @@ -148,9 +152,9 @@ int hc_clBuildProgram (void *hashcat_ctx, cl_program program, cl_uint int hc_clCompileProgram (void *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, cl_uint num_input_headers, const cl_program *input_headers, const char **header_include_names, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data); int hc_clLinkProgram (void *hashcat_ctx, cl_context context, cl_uint num_devices, const cl_device_id *device_list, const char *options, cl_uint num_input_programs, const cl_program *input_programs, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data, cl_program *program); int hc_clCreateKernel (void *hashcat_ctx, cl_program program, const char *kernel_name, cl_kernel *kernel); -int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem *mem); -int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel *kernel); -int hc_clReleaseProgram (void *hashcat_ctx, cl_program *program); +int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem mem); +int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel kernel); +int hc_clReleaseProgram (void *hashcat_ctx, cl_program program); int hc_clReleaseCommandQueue (void *hashcat_ctx, cl_command_queue command_queue); int hc_clReleaseContext (void *hashcat_ctx, cl_context context); int hc_clEnqueueMapBuffer (void *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_map, cl_map_flags map_flags, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event, void **buf); diff --git a/include/ext_cuda.h b/include/ext_cuda.h index 51cf3540f..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); @@ -1272,7 +1277,7 @@ int hc_cuDeviceGetName (void *hashcat_ctx, char *name, int len, CUdevice int hc_cuDeviceTotalMem (void *hashcat_ctx, size_t *bytes, CUdevice dev); int hc_cuDriverGetVersion (void *hashcat_ctx, int *driverVersion); int hc_cuEventCreate (void *hashcat_ctx, CUevent *phEvent, unsigned int Flags); -int hc_cuEventDestroy (void *hashcat_ctx, CUevent *hEvent); +int hc_cuEventDestroy (void *hashcat_ctx, CUevent hEvent); int hc_cuEventElapsedTime (void *hashcat_ctx, float *pMilliseconds, CUevent hStart, CUevent hEnd); int hc_cuEventQuery (void *hashcat_ctx, CUevent hEvent); int hc_cuEventRecord (void *hashcat_ctx, CUevent hEvent, CUstream hStream); @@ -1292,14 +1297,14 @@ int hc_cuMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, CUdeviceptr sr int hc_cuMemcpyHtoDAsync (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream); int hc_cuMemsetD32Async (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N, CUstream hStream); int hc_cuMemsetD8Async (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N, CUstream hStream); -int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr *dptr); +int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr); int hc_cuMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); int hc_cuModuleGetFunction (void *hashcat_ctx, CUfunction *hfunc, CUmodule hmod, const char *name); int hc_cuModuleGetGlobal (void *hashcat_ctx, CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name); int hc_cuModuleLoadDataEx (void *hashcat_ctx, CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues); -int hc_cuModuleUnload (void *hashcat_ctx, CUmodule *hmod); +int hc_cuModuleUnload (void *hashcat_ctx, CUmodule hmod); int hc_cuStreamCreate (void *hashcat_ctx, CUstream *phStream, unsigned int Flags); -int hc_cuStreamDestroy (void *hashcat_ctx, CUstream *hStream); +int hc_cuStreamDestroy (void *hashcat_ctx, CUstream hStream); int hc_cuStreamSynchronize (void *hashcat_ctx, CUstream hStream); int hc_cuCtxPushCurrent (void *hashcat_ctx, CUcontext ctx); int hc_cuCtxPopCurrent (void *hashcat_ctx, CUcontext *pctx); diff --git a/include/ext_hip.h b/include/ext_hip.h index 4b8bd678b..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 @@ -670,7 +675,7 @@ int hc_hipDeviceTotalMem (void *hashcat_ctx, size_t *bytes, hipDevice_t d int hc_hipDriverGetVersion (void *hashcat_ctx, int *driverVersion); int hc_hipEventCreate (void *hashcat_ctx, hipEvent_t *phEvent); int hc_hipEventCreateWithFlags (void *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags); -int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t *hEvent); +int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t hEvent); int hc_hipEventElapsedTime (void *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd); int hc_hipEventQuery (void *hashcat_ctx, hipEvent_t hEvent); int hc_hipEventRecord (void *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream); @@ -679,7 +684,7 @@ int hc_hipFuncGetAttribute (void *hashcat_ctx, int *pi, hipFunction_attribu int hc_hipInit (void *hashcat_ctx, unsigned int Flags); int hc_hipLaunchKernel (void *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra); int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize); -int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t *dptr); +int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t dptr); int hc_hipMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); int hc_hipMemcpyDtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount); int hc_hipMemcpyDtoH (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount); @@ -694,13 +699,13 @@ int hc_hipMemsetD8Async (void *hashcat_ctx, hipDeviceptr_t dstDevice, un int hc_hipModuleGetFunction (void *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name); int hc_hipModuleGetGlobal (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name); int hc_hipModuleLoadDataEx (void *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues); -int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t *hmod); +int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t hmod); int hc_hipRuntimeGetVersion (void *hashcat_ctx, int *runtimeVersion); int hc_hipSetDevice (void *hashcat_ctx, hipDevice_t dev); int hc_hipSetDeviceFlags (void *hashcat_ctx, unsigned int flags); int hc_hipStreamCreate (void *hashcat_ctx, hipStream_t *phStream); int hc_hipStreamCreateWithFlags (void *hashcat_ctx, hipStream_t *phStream, unsigned int flags); -int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t *hStream); +int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t hStream); int hc_hipStreamSynchronize (void *hashcat_ctx, hipStream_t hStream); int hc_hipGetDeviceProperties (void *hashcat_ctx, hipDeviceProp_t *prop, hipDevice_t dev); int hc_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor (void *hashcat_ctx, int *numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk); diff --git a/src/backend.c b/src/backend.c index 8fb353b36..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; } @@ -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) @@ -17008,54 +17008,54 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx) if (device_param->is_cuda == true) { - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_amp_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_comp_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_pws_idx); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_rules); - //hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_rules_c); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_combs); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_combs_c); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bfs); - //hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bfs_c); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_a); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_b); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_c); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s1_d); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_a); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_b); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_c); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_bitmap_s2_d); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_plain_bufs); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_digests_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_digests_shown); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_salt_bufs); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_esalt_bufs); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_tmps); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_hooks); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_result); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra0_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra1_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra2_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_extra3_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_root_css_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_markov_css_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_tm_c); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_st_digests_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_st_salts_buf); - hc_cuMemFree (hashcat_ctx, &device_param->cuda_d_st_esalts_buf); - 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); - hc_cuEventDestroy (hashcat_ctx, &device_param->cuda_event1); - hc_cuEventDestroy (hashcat_ctx, &device_param->cuda_event2); - 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); - hc_cuStreamDestroy (hashcat_ctx, &device_param->cuda_stream); + hc_cuStreamDestroyPtr (hashcat_ctx, &device_param->cuda_stream); - hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module); - hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module_mp); - hc_cuModuleUnload (hashcat_ctx, &device_param->cuda_module_amp); - 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); device_param->cuda_d_rules_c = 0; device_param->cuda_d_bfs_c = 0; @@ -17092,54 +17092,54 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx) if (device_param->is_hip == true) { - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_amp_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_comp_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_pws_idx); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_rules); - //hc_hipMemFree (hashcat_ctx, &device_param->hip_d_rules_c); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_combs); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_combs_c); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bfs); - //hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bfs_c); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_a); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_b); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_c); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s1_d); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_a); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_b); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_c); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_bitmap_s2_d); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_plain_bufs); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_digests_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_digests_shown); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_salt_bufs); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_esalt_bufs); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_tmps); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_hooks); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_result); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra0_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra1_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra2_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_extra3_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_root_css_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_markov_css_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_tm_c); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_st_digests_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_st_salts_buf); - hc_hipMemFree (hashcat_ctx, &device_param->hip_d_st_esalts_buf); - 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); - hc_hipEventDestroy (hashcat_ctx, &device_param->hip_event1); - hc_hipEventDestroy (hashcat_ctx, &device_param->hip_event2); - 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); - hc_hipStreamDestroy (hashcat_ctx, &device_param->hip_stream); + hc_hipStreamDestroyPtr (hashcat_ctx, &device_param->hip_stream); - hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module); - hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_mp); - hc_hipModuleUnload (hashcat_ctx, &device_param->hip_module_amp); - 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_rules_c = 0; device_param->hip_d_bfs_c = 0; @@ -17253,74 +17253,74 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx) if (device_param->is_opencl == true) { - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_amp_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_comp_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_pws_idx); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_rules); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_rules_c); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_combs); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_combs_c); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bfs); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bfs_c); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_a); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_b); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_c); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s1_d); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_a); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_b); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_c); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_bitmap_s2_d); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_plain_bufs); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_digests_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_digests_shown); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_salt_bufs); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_esalt_bufs); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_tmps); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_hooks); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_result); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra0_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra1_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra2_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_extra3_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_root_css_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_markov_css_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_tm_c); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_st_digests_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_st_salts_buf); - hc_clReleaseMemObject (hashcat_ctx, &device_param->opencl_d_st_esalts_buf); - 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); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel1); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel12); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel2p); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel2); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel2e); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel23); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel3); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel4); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_init2); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_loop2p); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_loop2); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_mp); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_mp_l); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_mp_r); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_tm); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_amp); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_memset); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_bzero); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_atinit); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_utf8toutf16le); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_decompress); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_aux1); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_aux2); - hc_clReleaseKernel (hashcat_ctx, &device_param->opencl_kernel_aux3); - 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); - hc_clReleaseProgram (hashcat_ctx, &device_param->opencl_program); - hc_clReleaseProgram (hashcat_ctx, &device_param->opencl_program_mp); - hc_clReleaseProgram (hashcat_ctx, &device_param->opencl_program_amp); - 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_context) hc_clReleaseContext (hashcat_ctx, device_param->opencl_context); diff --git a/src/ext_OpenCL.c b/src/ext_OpenCL.c index d86a5b1ba..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; @@ -663,17 +702,13 @@ int hc_clCreateKernel (void *hashcat_ctx, cl_program program, const char *kernel return 0; } -int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem *mem) +int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem mem) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; OCL_PTR *ocl = (OCL_PTR *) backend_ctx->ocl; - if (mem == NULL || *mem == NULL) return -1; - - const cl_int CL_err = ocl->clReleaseMemObject (*mem); - - *mem = NULL; + const cl_int CL_err = ocl->clReleaseMemObject (mem); if (CL_err != CL_SUCCESS) { @@ -685,17 +720,13 @@ int hc_clReleaseMemObject (void *hashcat_ctx, cl_mem *mem) return 0; } -int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel *kernel) +int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel kernel) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; OCL_PTR *ocl = (OCL_PTR *) backend_ctx->ocl; - if (kernel == NULL || *kernel == NULL) return -1; - - const cl_int CL_err = ocl->clReleaseKernel (*kernel); - - *kernel = NULL; + const cl_int CL_err = ocl->clReleaseKernel (kernel); if (CL_err != CL_SUCCESS) { @@ -707,17 +738,13 @@ int hc_clReleaseKernel (void *hashcat_ctx, cl_kernel *kernel) return 0; } -int hc_clReleaseProgram (void *hashcat_ctx, cl_program *program) +int hc_clReleaseProgram (void *hashcat_ctx, cl_program program) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; OCL_PTR *ocl = (OCL_PTR *) backend_ctx->ocl; - if (program == NULL || *program == NULL) return -1; - - const cl_int CL_err = ocl->clReleaseProgram (*program); - - *program = NULL; + const cl_int CL_err = ocl->clReleaseProgram (program); if (CL_err != CL_SUCCESS) { diff --git a/src/ext_cuda.c b/src/ext_cuda.c index 88ab2ee1d..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; @@ -414,17 +466,13 @@ int hc_cuModuleLoadDataEx (void *hashcat_ctx, CUmodule *module, const void *imag return 0; } -int hc_cuModuleUnload (void *hashcat_ctx, CUmodule *hmod) +int hc_cuModuleUnload (void *hashcat_ctx, CUmodule hmod) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - if (hmod == NULL || *hmod == NULL) return -1; - - const CUresult CU_err = cuda->cuModuleUnload (*hmod); - - *hmod = NULL; + const CUresult CU_err = cuda->cuModuleUnload (hmod); if (CU_err != CUDA_SUCCESS) { @@ -499,17 +547,13 @@ int hc_cuMemAlloc (void *hashcat_ctx, CUdeviceptr *dptr, size_t bytesize) return 0; } -int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr *dptr) +int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - if (dptr == NULL || *dptr == 0) return -1; - - const CUresult CU_err = cuda->cuMemFree (*dptr); - - *dptr = 0; + const CUresult CU_err = cuda->cuMemFree (dptr); if (CU_err != CUDA_SUCCESS) { @@ -963,17 +1007,13 @@ int hc_cuStreamCreate (void *hashcat_ctx, CUstream *phStream, unsigned int Flags return 0; } -int hc_cuStreamDestroy (void *hashcat_ctx, CUstream *hStream) +int hc_cuStreamDestroy (void *hashcat_ctx, CUstream hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - if (hStream == NULL || *hStream == NULL) return -1; - - const CUresult CU_err = cuda->cuStreamDestroy (*hStream); - - *hStream = NULL; + const CUresult CU_err = cuda->cuStreamDestroy (hStream); if (CU_err != CUDA_SUCCESS) { @@ -1102,17 +1142,13 @@ int hc_cuEventCreate (void *hashcat_ctx, CUevent *phEvent, unsigned int Flags) return 0; } -int hc_cuEventDestroy (void *hashcat_ctx, CUevent *hEvent) +int hc_cuEventDestroy (void *hashcat_ctx, CUevent hEvent) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; - if (hEvent == NULL || *hEvent == NULL) return -1; - - const CUresult CU_err = cuda->cuEventDestroy (*hEvent); - - *hEvent = NULL; + const CUresult CU_err = cuda->cuEventDestroy (hEvent); if (CU_err != CUDA_SUCCESS) { diff --git a/src/ext_hip.c b/src/ext_hip.c index 00d938d57..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; @@ -565,17 +617,13 @@ int hc_hipEventCreateWithFlags (void *hashcat_ctx, hipEvent_t *phEvent, unsigned return 0; } -int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t *hEvent) +int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t hEvent) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - if (hEvent == NULL || *hEvent == NULL) return -1; - - const hipError_t HIP_err = hip->hipEventDestroy (*hEvent); - - *hEvent = NULL; + const hipError_t HIP_err = hip->hipEventDestroy (hEvent); if (HIP_err != hipSuccess) { @@ -785,17 +833,13 @@ int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize) return 0; } -int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t *dptr) +int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t dptr) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - if (dptr == NULL || *dptr == NULL) return -1; - - const hipError_t HIP_err = hip->hipMemFree (*dptr); - - *dptr = 0; + const hipError_t HIP_err = hip->hipMemFree (dptr); if (HIP_err != hipSuccess) { @@ -1196,17 +1240,13 @@ int hc_hipModuleLoadDataEx (void *hashcat_ctx, hipModule_t *module, const void * return 0; } -int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t *hmod) +int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t hmod) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - if (hmod == NULL || *hmod == NULL) return -1; - - const hipError_t HIP_err = hip->hipModuleUnload (*hmod); - - *hmod = NULL; + const hipError_t HIP_err = hip->hipModuleUnload (hmod); if (HIP_err != hipSuccess) { @@ -1362,17 +1402,13 @@ int hc_hipStreamCreateWithFlags (void *hashcat_ctx, hipStream_t *phStream, unsig return 0; } -int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t *hStream) +int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - if (hStream == NULL || *hStream == NULL) return -1; - - const hipError_t HIP_err = hip->hipStreamDestroy (*hStream); - - *hStream = NULL; + const hipError_t HIP_err = hip->hipStreamDestroy (hStream); if (HIP_err != hipSuccess) { diff --git a/src/ext_metal.m b/src/ext_metal.m index dd4e73d1f..1bf348798 100644 --- a/src/ext_metal.m +++ b/src/ext_metal.m @@ -999,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; } @@ -1085,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; } diff --git a/src/memory.c b/src/memory.c index 5c26c3f3e..15f4cbe19 100644 --- a/src/memory.c +++ b/src/memory.c @@ -47,8 +47,6 @@ void *hcrealloc (void *ptr, const size_t oldsz, const size_t addsz) char *hcstrdup (const char *s) { - if (s == NULL) return (NULL); - const size_t len = strlen (s); char *b = (char *) hcmalloc (len + 1);