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

Added workaround to get rid of internal runtimes memory leaks

Browse Source 
As of now, especially in the benchmark mode, hashcat will not go to create and destroy context and command-queue for each enabled device each time it switches from one hash-mode to the next.
Specifically using OpenCL with an NVIDIA device, it was not possible to complete the benchmark because clCreateContext has memory leaks that slowly consume all available GPU memory until hashcat can activate a new context and disable the device.

Avoid deprecated HIP functions

All hipCtx* features have been declared deprecated, so we have replaced them with the new ones, also fixing a critical bug on handling multiple AMD devices in the same system.
This commit is contained in:
Gabriele Gristina 2025-07-06 21:28:37 +02:00
parent 0576c41491
commit f663abee44
No known key found for this signature in database
GPG Key ID: 9F68B59298F311F0
7 changed files with 793 additions and 790 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
docs/changes.txt
View File

@ -129,6 +129,7 @@
- Alias Devices: Prevents hashcat, when started with x86_64 emulation on Apple Silicon, from showing the Apple M1 OpenCL CPU as an alias for the Apple M1 Metal GPU
- Apple Driver: Automatically enable GPU support on Apple OpenCL instead of CPU support
- 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: Splitting backend_ctx_devices_init into smaller runtime-specific functions
- Backend Checks: Describe workaround in error message when detecting more than 64 backend devices
@ -141,6 +142,7 @@
- Building: Support building windows binaries on macOS using MinGW
- Dependencies: Updated OpenCL-Headers to v2024.10.24 (commit 265df85)
- Documents: Updated BUILD.md and added BUILD_macOS.md (containing instructions for building windows binaries on macOS)
- HIP Backend: Avoid deprecated functions
- Modules: Added support for non-zero IVs for -m 6800 (Lastpass). Also added `tools/lastpass2hashcat.py`
- Modules: Updated module_unstable_warning
- Open Document Format: Added support for small documents with content length < 1024

976
include/ext_hip.h
View File

File diff suppressed because it is too large Load Diff

7
src/autotune.c
View File

@ -679,7 +679,7 @@ HC_API_CALL void *thread_autotune (void *p)
if (device_param->is_hip == true)
{
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL;
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1) return NULL;
}
// check for autotune failure
@ -695,11 +695,6 @@ HC_API_CALL void *thread_autotune (void *p)
if (hc_cuCtxPopCurrent (hashcat_ctx, &device_param->cuda_context) == -1) return NULL;
}
if (device_param->is_hip == true)
{
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL;
}
return NULL;
}

370
src/backend.c
View File

@ -993,7 +993,7 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c
if (device_param->is_hip == true)
{
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1;
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1) return -1;
if (hc_hipMemcpyDtoH (hashcat_ctx, &pw_idx, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t)) == -1) return -1;
@ -1059,11 +1059,6 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c
if (hc_cuCtxPopCurrent (hashcat_ctx, &device_param->cuda_context) == -1) return -1;
}
if (device_param->is_hip == true)
{
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return -1;
}
return 0;
}
@ -1082,13 +1077,11 @@ int copy_pws_idx (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u
if (device_param->is_hip == true)
{
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1;
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1) return -1;
if (hc_hipMemcpyDtoH (hashcat_ctx, dest, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t))) == -1) return -1;
if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1;
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return -1;
}
#if defined (__APPLE__)
@ -1122,13 +1115,11 @@ int copy_pws_comp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
if (device_param->is_hip == true)
{
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1;
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1) return -1;
if (hc_hipMemcpyDtoH (hashcat_ctx, dest, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1;
if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1;
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return -1;
}
#if defined (__APPLE__)
@ -5937,24 +5928,24 @@ static void backend_ctx_devices_init_cuda (hashcat_ctx_t *hashcat_ctx, int *virt
device_param->has_prmt = (sm >= 20) ? true : false;
device_param->has_shfw = (sm >= 70) ? true : false;
// one-time init cuda context
if (hc_cuCtxCreate (hashcat_ctx, &device_param->cuda_context, CU_CTX_SCHED_BLOCKING_SYNC, device_param->cuda_device) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1)
{
device_param->skipped = true;
continue;
}
// device_available_mem
CUcontext cuda_context;
if (hc_cuCtxCreate (hashcat_ctx, &cuda_context, CU_CTX_SCHED_BLOCKING_SYNC, device_param->cuda_device) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_cuCtxPushCurrent (hashcat_ctx, cuda_context) == -1)
{
device_param->skipped = true;
continue;
}
size_t free = 0;
size_t total = 0;
@ -5967,14 +5958,7 @@ static void backend_ctx_devices_init_cuda (hashcat_ctx_t *hashcat_ctx, int *virt
device_param->device_available_mem = ((u64) free * (100 - user_options->backend_devices_keepfree)) / 100;
if (hc_cuCtxPopCurrent (hashcat_ctx, &cuda_context) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_cuCtxDestroy (hashcat_ctx, cuda_context) == -1)
if (hc_cuCtxPopCurrent (hashcat_ctx, &device_param->cuda_context) == -1)
{
device_param->skipped = true;
@ -6440,24 +6424,24 @@ static void backend_ctx_devices_init_hip (hashcat_ctx_t *hashcat_ctx, int *virth
device_param->has_prmt = false;
device_param->has_shfw = true; // always reports false : prop.arch.hasFunnelShift;
// one-time init hip context
if (hc_hipSetDeviceFlags (hashcat_ctx, hipDeviceScheduleBlockingSync) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1)
{
device_param->skipped = true;
continue;
}
// device_available_mem
hipCtx_t hip_context;
if (hc_hipCtxCreate (hashcat_ctx, &hip_context, hipDeviceScheduleBlockingSync, device_param->hip_device) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_hipCtxPushCurrent (hashcat_ctx, hip_context) == -1)
{
device_param->skipped = true;
continue;
}
size_t free = 0;
size_t total = 0;
@ -6470,20 +6454,6 @@ static void backend_ctx_devices_init_hip (hashcat_ctx_t *hashcat_ctx, int *virth
device_param->device_available_mem = ((u64) free * (100 - user_options->backend_devices_keepfree)) / 100;
if (hc_hipCtxPopCurrent (hashcat_ctx, &hip_context) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_hipCtxDestroy (hashcat_ctx, hip_context) == -1)
{
device_param->skipped = true;
continue;
}
#if defined (__linux__)
if (strchr (folder_config->cpath_real, ' ') != NULL)
{
@ -8623,13 +8593,9 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
if (device_param->skipped == true) continue;
}
/**
* create command-queue
*/
// one-time init metal command-queue
mtl_command_queue command_queue;
if (hc_mtlCreateCommandQueue (hashcat_ctx, device_param->metal_device, &command_queue) == -1)
if (hc_mtlCreateCommandQueue (hashcat_ctx, device_param->metal_device, &device_param->metal_command_queue) == -1)
{
device_param->skipped = true;
@ -8678,11 +8644,11 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
u8 tmp_host[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
if (hc_mtlMemcpyHtoD (hashcat_ctx, command_queue, tmp_device[c], 0, tmp_host, sizeof (tmp_host)) == -1) break;
if (hc_mtlMemcpyDtoH (hashcat_ctx, command_queue, tmp_host, tmp_device[c], 0, sizeof (tmp_host)) == -1) break;
if (hc_mtlMemcpyHtoD (hashcat_ctx, device_param->metal_command_queue, tmp_device[c], 0, tmp_host, sizeof (tmp_host)) == -1) break;
if (hc_mtlMemcpyDtoH (hashcat_ctx, device_param->metal_command_queue, tmp_host, tmp_device[c], 0, sizeof (tmp_host)) == -1) break;
if (hc_mtlMemcpyHtoD (hashcat_ctx, command_queue, tmp_device[c], MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), tmp_host, sizeof (tmp_host)) == -1) break;
if (hc_mtlMemcpyDtoH (hashcat_ctx, command_queue, tmp_host, tmp_device[c], MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), sizeof (tmp_host)) == -1) break;
if (hc_mtlMemcpyHtoD (hashcat_ctx, device_param->metal_command_queue, tmp_device[c], MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), tmp_host, sizeof (tmp_host)) == -1) break;
if (hc_mtlMemcpyDtoH (hashcat_ctx, device_param->metal_command_queue, tmp_host, tmp_device[c], MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), sizeof (tmp_host)) == -1) break;
}
device_param->device_available_mem = MAX_ALLOC_CHECKS_SIZE;
@ -8707,8 +8673,6 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
hcfree (tmp_device);
}
hc_mtlReleaseCommandQueue (hashcat_ctx, command_queue);
if (device_param->device_host_unified_memory == 1)
{
// so, we actually have only half the memory because we need the same buffers on host side
@ -8734,11 +8698,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
if (device_param->skipped == true) continue;
}
/**
* create context for each device
*/
cl_context context;
// one-time init opencl context
/*
cl_context_properties properties[3];
@ -8747,10 +8707,10 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
properties[1] = (cl_context_properties) device_param->opencl_platform;
properties[2] = 0;
CL_rc = hc_clCreateContext (hashcat_ctx, properties, 1, &device_param->opencl_device, NULL, NULL, &context);
CL_rc = hc_clCreateContext (hashcat_ctx, properties, 1, &device_param->opencl_device, NULL, NULL, &device_param->opencl_context);
*/
if (hc_clCreateContext (hashcat_ctx, NULL, 1, &device_param->opencl_device, NULL, NULL, &context) == -1)
if (hc_clCreateContext (hashcat_ctx, NULL, 1, &device_param->opencl_device, NULL, NULL, &device_param->opencl_context) == -1)
{
device_param->skipped = true;
@ -8760,13 +8720,9 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
continue;
}
/**
* create command-queue
*/
// one-time init open command-queue
cl_command_queue command_queue;
if (hc_clCreateCommandQueue (hashcat_ctx, context, device_param->opencl_device, 0, &command_queue) == -1)
if (hc_clCreateCommandQueue (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, CL_QUEUE_PROFILING_ENABLE, &device_param->opencl_command_queue) == -1)
{
device_param->skipped = true;
@ -8781,17 +8737,17 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
if ((device_param->opencl_device_type & CL_DEVICE_TYPE_GPU) && (device_param->opencl_platform_vendor_id == VENDOR_ID_AMD))
{
#define RUN_INSTRUCTION_CHECKS() \
device_param->has_vadd = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADD_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vaddc = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADDC_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vadd_co = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADD_CO_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vaddc_co = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADDC_CO_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vsub = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUB_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vsubb = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUBB_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vsub_co = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUB_CO_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vsubb_co = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUBB_CO_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vadd3 = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADD3_U32 %0, 0, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vbfe = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_BFE_U32 %0, 0, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vperm = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_PERM_B32 %0, 0, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vadd = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADD_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vaddc = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADDC_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vadd_co = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADD_CO_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vaddc_co = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADDC_CO_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vsub = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUB_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vsubb = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUBB_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vsub_co = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUB_CO_U32 %0, vcc, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vsubb_co = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_SUBB_CO_U32 %0, vcc, 0, 0, vcc;\" : \"=v\"(r1)); }"); \
device_param->has_vadd3 = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_ADD3_U32 %0, 0, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vbfe = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_BFE_U32 %0, 0, 0, 0;\" : \"=v\"(r1)); }"); \
device_param->has_vperm = opencl_test_instruction (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, "__kernel void test () { uint r1; __asm__ __volatile__ (\"V_PERM_B32 %0, 0, 0, 0;\" : \"=v\"(r1)); }"); \
if (backend_devices_idx > 0)
{
@ -8979,7 +8935,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
OCL_PTR *ocl = (OCL_PTR *) backend_ctx->ocl;
tmp_device[c] = ocl->clCreateBuffer (context, CL_MEM_READ_WRITE, MAX_ALLOC_CHECKS_SIZE, NULL, &CL_err);
tmp_device[c] = ocl->clCreateBuffer (device_param->opencl_context, CL_MEM_READ_WRITE, MAX_ALLOC_CHECKS_SIZE, NULL, &CL_err);
if (CL_err != CL_SUCCESS)
{
@ -8992,11 +8948,11 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
u8 tmp_host[8];
if (ocl->clEnqueueReadBuffer (command_queue, tmp_device[c], CL_TRUE, 0, sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
if (ocl->clEnqueueWriteBuffer (command_queue, tmp_device[c], CL_TRUE, 0, sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
if (ocl->clEnqueueReadBuffer (device_param->opencl_command_queue, tmp_device[c], CL_TRUE, 0, sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
if (ocl->clEnqueueWriteBuffer (device_param->opencl_command_queue, tmp_device[c], CL_TRUE, 0, sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
if (ocl->clEnqueueReadBuffer (command_queue, tmp_device[c], CL_TRUE, MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
if (ocl->clEnqueueWriteBuffer (command_queue, tmp_device[c], CL_TRUE, MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
if (ocl->clEnqueueReadBuffer (device_param->opencl_command_queue, tmp_device[c], CL_TRUE, MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
if (ocl->clEnqueueWriteBuffer (device_param->opencl_command_queue, tmp_device[c], CL_TRUE, MAX_ALLOC_CHECKS_SIZE - sizeof (tmp_host), sizeof (tmp_host), tmp_host, 0, NULL, NULL) != CL_SUCCESS) break;
}
device_param->device_available_mem = MAX_ALLOC_CHECKS_SIZE;
@ -9008,24 +8964,28 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
// clean up
int r = 0;
for (c = 0; c < MAX_ALLOC_CHECKS_CNT; c++)
{
if (((c + 1 + 1) * MAX_ALLOC_CHECKS_SIZE) >= device_param->device_global_mem) break;
if (tmp_device[c] != NULL)
{
if (hc_clReleaseMemObject (hashcat_ctx, tmp_device[c]) == -1) return -1;
if (hc_clReleaseMemObject (hashcat_ctx, tmp_device[c]) == -1) r = -1;
}
}
hcfree (tmp_device);
if (r == -1)
{
// return -1 here is blocking, to be better evaluated
//return -1;
}
}
}
hc_clReleaseCommandQueue (hashcat_ctx, command_queue);
hc_clReleaseContext (hashcat_ctx, context);
if (device_param->device_host_unified_memory == 1)
{
// so, we actually have only half the memory because we need the same buffers on host side
@ -9065,23 +9025,61 @@ void backend_ctx_devices_destroy (hashcat_ctx_t *hashcat_ctx)
hcfree (backend_ctx->opencl_platforms_version[opencl_platforms_idx]);
}
// one-time release context/command-queue from all runtimes
for (int backend_devices_idx = 0; backend_devices_idx < backend_ctx->backend_devices_cnt; backend_devices_idx++)
{
hc_device_param_t *device_param = &backend_ctx->devices_param[backend_devices_idx];
hcfree (device_param->device_name);
if (device_param->is_cuda == true)
{
if (device_param->cuda_context)
{
hc_cuCtxDestroy (hashcat_ctx, device_param->cuda_context);
device_param->cuda_context = NULL;
}
}
if (device_param->is_hip == true)
{
hcfree (device_param->gcnArchName);
}
#if defined (__APPLE__)
if (device_param->is_metal == true)
{
if (device_param->metal_command_queue)
{
hc_mtlReleaseCommandQueue (hashcat_ctx, device_param->metal_command_queue);
device_param->metal_command_queue = NULL;
}
}
#endif
if (device_param->is_opencl == true)
{
hcfree (device_param->opencl_driver_version);
hcfree (device_param->opencl_device_version);
hcfree (device_param->opencl_device_c_version);
hcfree (device_param->opencl_device_vendor);
}
if (device_param->is_hip == true)
{
hcfree (device_param->gcnArchName);
if (device_param->opencl_command_queue)
{
hc_clReleaseCommandQueue (hashcat_ctx, device_param->opencl_command_queue);
device_param->opencl_command_queue = NULL;
}
if (device_param->opencl_context)
{
hc_clReleaseContext (hashcat_ctx, device_param->opencl_context);
device_param->opencl_context = NULL;
}
}
}
@ -10730,93 +10728,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
}
#endif
/**
* create context for each device
*/
if (device_param->is_cuda == true)
{
if (hc_cuCtxCreate (hashcat_ctx, &device_param->cuda_context, CU_CTX_SCHED_BLOCKING_SYNC, device_param->cuda_device) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1)
{
device_param->skipped = true;
continue;
}
}
if (device_param->is_hip == true)
{
if (hc_hipCtxCreate (hashcat_ctx, &device_param->hip_context, hipDeviceScheduleBlockingSync, device_param->hip_device) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1)
{
device_param->skipped = true;
continue;
}
}
#if defined (__APPLE__)
if (device_param->is_metal == true)
{
/**
* create command-queue
*/
if (hc_mtlCreateCommandQueue (hashcat_ctx, device_param->metal_device, &device_param->metal_command_queue) == -1)
{
device_param->skipped = true;
continue;
}
}
#endif
if (device_param->is_opencl == true)
{
/*
cl_context_properties properties[3];
properties[0] = CL_CONTEXT_PLATFORM;
properties[1] = (cl_context_properties) device_param->opencl_platform;
properties[2] = 0;
CL_rc = hc_clCreateContext (hashcat_ctx, properties, 1, &device_param->opencl_device, NULL, NULL, &device_param->opencl_context);
*/
if (hc_clCreateContext (hashcat_ctx, NULL, 1, &device_param->opencl_device, NULL, NULL, &device_param->opencl_context) == -1)
{
device_param->skipped = true;
continue;
}
/**
* create command-queue
*/
// not supported with NV
// device_param->opencl_command_queue = hc_clCreateCommandQueueWithProperties (hashcat_ctx, device_param->opencl_device, NULL);
if (hc_clCreateCommandQueue (hashcat_ctx, device_param->opencl_context, device_param->opencl_device, CL_QUEUE_PROFILING_ENABLE, &device_param->opencl_command_queue) == -1)
{
device_param->skipped = true;
continue;
}
}
// re-using context/command-queue, there is no need to re-initialize them
/**
* create stream for CUDA devices
@ -10824,6 +10736,13 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
if (device_param->is_cuda == true)
{
if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_cuStreamCreate (hashcat_ctx, &device_param->cuda_stream, CU_STREAM_DEFAULT) == -1)
{
device_param->skipped = true;
@ -10838,7 +10757,14 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
if (device_param->is_hip == true)
{
if (hc_hipStreamCreate (hashcat_ctx, &device_param->hip_stream, hipStreamDefault) == -1)
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_hipStreamCreateWithFlags (hashcat_ctx, &device_param->hip_stream, hipStreamDefault) == -1)
{
device_param->skipped = true;
@ -10880,21 +10806,21 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
if (device_param->is_hip == true)
{
if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event1, hipEventBlockingSync) == -1)
if (hc_hipEventCreateWithFlags (hashcat_ctx, &device_param->hip_event1, hipEventBlockingSync) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event2, hipEventBlockingSync) == -1)
if (hc_hipEventCreateWithFlags (hashcat_ctx, &device_param->hip_event2, hipEventBlockingSync) == -1)
{
device_param->skipped = true;
continue;
}
if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event3, hipEventDisableTiming) == -1)
if (hc_hipEventCreateWithFlags (hashcat_ctx, &device_param->hip_event3, hipEventDisableTiming) == -1)
{
device_param->skipped = true;
@ -16927,16 +16853,6 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
}
}
if (device_param->is_hip == true)
{
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1)
{
device_param->skipped = true;
continue;
}
}
hardware_power_all += hardware_power_max;
EVENT_DATA (EVENT_BACKEND_DEVICE_INIT_POST, &backend_devices_idx, sizeof (int));
@ -17058,7 +16974,7 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
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);
if (device_param->cuda_context) hc_cuCtxDestroy (hashcat_ctx, device_param->cuda_context);
//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;
@ -17135,7 +17051,7 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
device_param->cuda_module_amp = NULL;
device_param->cuda_module_shared = NULL;
device_param->cuda_context = NULL;
//device_param->cuda_context = NULL;
}
if (device_param->is_hip == true)
@ -17189,8 +17105,6 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
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);
if (device_param->hip_context) hc_hipCtxDestroy (hashcat_ctx, device_param->hip_context);
device_param->hip_d_pws_buf = 0;
device_param->hip_d_pws_amp_buf = 0;
device_param->hip_d_pws_comp_buf = 0;
@ -17265,8 +17179,6 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
device_param->hip_module_mp = NULL;
device_param->hip_module_amp = NULL;
device_param->hip_module_shared = NULL;
device_param->hip_context = NULL;
}
#if defined (__APPLE__)
@ -17341,7 +17253,7 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
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);
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);
@ -17411,7 +17323,7 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
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_command_queue = NULL;
//device_param->metal_device = NULL;
}
#endif // __APPLE__
@ -17487,9 +17399,9 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
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);
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;
@ -17557,8 +17469,8 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx)
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;

14
src/dispatch.c
View File

@ -381,7 +381,7 @@ HC_API_CALL void *thread_calc_stdin (void *p)
if (device_param->is_hip == true)
{
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL;
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1) return NULL;
}
if (calc_stdin (hashcat_ctx, device_param) == -1)
@ -396,11 +396,6 @@ HC_API_CALL void *thread_calc_stdin (void *p)
if (hc_cuCtxPopCurrent (hashcat_ctx, &device_param->cuda_context) == -1) return NULL;
}
if (device_param->is_hip == true)
{
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL;
}
if (bridge_ctx->enabled == true)
{
if (bridge_ctx->thread_term != BRIDGE_DEFAULT)
@ -1685,7 +1680,7 @@ HC_API_CALL void *thread_calc (void *p)
if (device_param->is_hip == true)
{
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL;
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1) return NULL;
}
if (calc (hashcat_ctx, device_param) == -1)
@ -1700,11 +1695,6 @@ HC_API_CALL void *thread_calc (void *p)
if (hc_cuCtxPopCurrent (hashcat_ctx, &device_param->cuda_context) == -1) return NULL;
}
if (device_param->is_hip == true)
{
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL;
}
if (bridge_ctx->enabled == true)
{
if (bridge_ctx->thread_term != BRIDGE_DEFAULT)

210
src/ext_hip.c
View File

@ -115,51 +115,55 @@ int hip_init (void *hashcat_ctx)
// a good reference is cuda.h itself
// this needs to be verified for each new cuda release
HC_LOAD_FUNC_HIP (hip, hipCtxCreate, hipCtxCreate, HIP_HIPCTXCREATE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxDestroy, hipCtxDestroy, HIP_HIPCTXDESTROY, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxPopCurrent, hipCtxPopCurrent, HIP_HIPCTXPOPCURRENT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxPushCurrent, hipCtxPushCurrent, HIP_HIPCTXPUSHCURRENT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxSetCurrent, hipCtxSetCurrent, HIP_HIPCTXSETCURRENT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxSynchronize, hipCtxSynchronize, HIP_HIPCTXSYNCHRONIZE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGet, hipDeviceGet, HIP_HIPDEVICEGET, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGetAttribute, hipDeviceGetAttribute, HIP_HIPDEVICEGETATTRIBUTE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGetCount, hipGetDeviceCount, HIP_HIPDEVICEGETCOUNT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGetName, hipDeviceGetName, HIP_HIPDEVICEGETNAME, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceTotalMem, hipDeviceTotalMem, HIP_HIPDEVICETOTALMEM, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDriverGetVersion, hipDriverGetVersion, HIP_HIPDRIVERGETVERSION, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventCreate, hipEventCreateWithFlags, HIP_HIPEVENTCREATE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventDestroy, hipEventDestroy, HIP_HIPEVENTDESTROY, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventElapsedTime, hipEventElapsedTime, HIP_HIPEVENTELAPSEDTIME, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventRecord, hipEventRecord, HIP_HIPEVENTRECORD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventSynchronize, hipEventSynchronize, HIP_HIPEVENTSYNCHRONIZE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipFuncGetAttribute, hipFuncGetAttribute, HIP_HIPFUNCGETATTRIBUTE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipGetErrorName, hipDrvGetErrorName, HIP_HIPGETERRORNAME, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipGetErrorString, hipDrvGetErrorString, HIP_HIPGETERRORSTRING, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipInit, hipInit, HIP_HIPINIT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipLaunchKernel, hipModuleLaunchKernel, HIP_HIPLAUNCHKERNEL, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemAlloc, hipMalloc, HIP_HIPMEMALLOC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemFree, hipFree, HIP_HIPMEMFREE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemGetInfo, hipMemGetInfo, HIP_HIPMEMGETINFO, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoD, hipMemcpyDtoD, HIP_HIPMEMCPYDTOD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoH, hipMemcpyDtoH, HIP_HIPMEMCPYDTOH, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoD, hipMemcpyHtoD, HIP_HIPMEMCPYHTOD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD32, hipMemsetD32, HIP_HIPMEMSETD32, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD8, hipMemsetD8, HIP_HIPMEMSETD8, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoDAsync, hipMemcpyDtoDAsync, HIP_HIPMEMCPYDTODASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoHAsync, hipMemcpyDtoHAsync, HIP_HIPMEMCPYDTOHASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD32Async, hipMemsetD32Async, HIP_HIPMEMSETD32ASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD8Async, hipMemsetD8Async, HIP_HIPMEMSETD8ASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleGetFunction, hipModuleGetFunction, HIP_HIPMODULEGETFUNCTION, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleGetGlobal, hipModuleGetGlobal, HIP_HIPMODULEGETGLOBAL, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleLoadDataEx, hipModuleLoadDataEx, HIP_HIPMODULELOADDATAEX, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleUnload, hipModuleUnload, HIP_HIPMODULEUNLOAD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipRuntimeGetVersion, hipRuntimeGetVersion, HIP_HIPRUNTIMEGETVERSION, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipStreamCreate, hipStreamCreate, HIP_HIPSTREAMCREATE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipStreamDestroy, hipStreamDestroy, HIP_HIPSTREAMDESTROY, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipStreamSynchronize, hipStreamSynchronize, HIP_HIPSTREAMSYNCHRONIZE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipGetDeviceProperties, hipGetDevicePropertiesR0600, HIP_HIPGETDEVICEPROPERTIES, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleOccupancyMaxActiveBlocksPerMultiprocessor, hipModuleOccupancyMaxActiveBlocksPerMultiprocessor, HIP_HIPMODULEOCCUPANCYMAXACTIVEBLOCKSPERMULTIPROCESSOR, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxCreate, hipCtxCreate, HIP_HIPCTXCREATE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxDestroy, hipCtxDestroy, HIP_HIPCTXDESTROY, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxPopCurrent, hipCtxPopCurrent, HIP_HIPCTXPOPCURRENT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxPushCurrent, hipCtxPushCurrent, HIP_HIPCTXPUSHCURRENT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxSetCurrent, hipCtxSetCurrent, HIP_HIPCTXSETCURRENT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipCtxSynchronize, hipCtxSynchronize, HIP_HIPCTXSYNCHRONIZE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGet, hipDeviceGet, HIP_HIPDEVICEGET, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGetAttribute, hipDeviceGetAttribute, HIP_HIPDEVICEGETATTRIBUTE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGetCount, hipGetDeviceCount, HIP_HIPDEVICEGETCOUNT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceGetName, hipDeviceGetName, HIP_HIPDEVICEGETNAME, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDeviceTotalMem, hipDeviceTotalMem, HIP_HIPDEVICETOTALMEM, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipDriverGetVersion, hipDriverGetVersion, HIP_HIPDRIVERGETVERSION, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventCreate, hipEventCreate, HIP_HIPEVENTCREATE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventCreateWithFlags, hipEventCreateWithFlags, HIP_HIPEVENTCREATEWITHFLAGS, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventDestroy, hipEventDestroy, HIP_HIPEVENTDESTROY, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventElapsedTime, hipEventElapsedTime, HIP_HIPEVENTELAPSEDTIME, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventRecord, hipEventRecord, HIP_HIPEVENTRECORD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipEventSynchronize, hipEventSynchronize, HIP_HIPEVENTSYNCHRONIZE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipFuncGetAttribute, hipFuncGetAttribute, HIP_HIPFUNCGETATTRIBUTE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipGetErrorName, hipDrvGetErrorName, HIP_HIPGETERRORNAME, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipGetErrorString, hipDrvGetErrorString, HIP_HIPGETERRORSTRING, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipInit, hipInit, HIP_HIPINIT, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipLaunchKernel, hipModuleLaunchKernel, HIP_HIPLAUNCHKERNEL, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemAlloc, hipMalloc, HIP_HIPMEMALLOC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemFree, hipFree, HIP_HIPMEMFREE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemGetInfo, hipMemGetInfo, HIP_HIPMEMGETINFO, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoD, hipMemcpyDtoD, HIP_HIPMEMCPYDTOD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoH, hipMemcpyDtoH, HIP_HIPMEMCPYDTOH, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoD, hipMemcpyHtoD, HIP_HIPMEMCPYHTOD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD32, hipMemsetD32, HIP_HIPMEMSETD32, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD8, hipMemsetD8, HIP_HIPMEMSETD8, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoDAsync, hipMemcpyDtoDAsync, HIP_HIPMEMCPYDTODASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoHAsync, hipMemcpyDtoHAsync, HIP_HIPMEMCPYDTOHASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD32Async, hipMemsetD32Async, HIP_HIPMEMSETD32ASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipMemsetD8Async, hipMemsetD8Async, HIP_HIPMEMSETD8ASYNC, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleGetFunction, hipModuleGetFunction, HIP_HIPMODULEGETFUNCTION, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleGetGlobal, hipModuleGetGlobal, HIP_HIPMODULEGETGLOBAL, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleLoadDataEx, hipModuleLoadDataEx, HIP_HIPMODULELOADDATAEX, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleUnload, hipModuleUnload, HIP_HIPMODULEUNLOAD, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipRuntimeGetVersion, hipRuntimeGetVersion, HIP_HIPRUNTIMEGETVERSION, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipSetDevice, hipSetDevice, HIP_HIPSETDEVICE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipSetDeviceFlags, hipSetDeviceFlags, HIP_HIPSETDEVICEFLAGS, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipStreamCreate, hipStreamCreate, HIP_HIPSTREAMCREATE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipStreamCreateWithFlags, hipStreamCreateWithFlags, HIP_HIPSTREAMCREATEWITHFLAGS, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipStreamDestroy, hipStreamDestroy, HIP_HIPSTREAMDESTROY, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipStreamSynchronize, hipStreamSynchronize, HIP_HIPSTREAMSYNCHRONIZE, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipGetDeviceProperties, hipGetDevicePropertiesR0600, HIP_HIPGETDEVICEPROPERTIES, HIP, 1);
HC_LOAD_FUNC_HIP (hip, hipModuleOccupancyMaxActiveBlocksPerMultiprocessor, hipModuleOccupancyMaxActiveBlocksPerMultiprocessor, HIP_HIPMODULEOCCUPANCYMAXACTIVEBLOCKSPERMULTIPROCESSOR, HIP, 1);
return 0;
}
@ -507,13 +511,13 @@ int hc_hipDriverGetVersion (void *hashcat_ctx, int *driverVersion)
return 0;
}
int hc_hipEventCreate (void *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags)
int hc_hipEventCreate (void *hashcat_ctx, hipEvent_t *phEvent)
{
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->hipEventCreate (phEvent, Flags);
const hipError_t HIP_err = hip->hipEventCreate (phEvent);
if (HIP_err != hipSuccess)
{
@ -534,6 +538,33 @@ int hc_hipEventCreate (void *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flag
return 0;
}
int hc_hipEventCreateWithFlags (void *hashcat_ctx, hipEvent_t *phEvent, unsigned int flags)
{
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->hipEventCreateWithFlags (phEvent, flags);
if (HIP_err != hipSuccess)
{
const char *pStr = NULL;
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
{
event_log_error (hashcat_ctx, "hipEventCreateWithFlags(): %s", pStr);
}
else
{
event_log_error (hashcat_ctx, "hipEventCreateWithFlags(): %d", HIP_err);
}
return -1;
}
return 0;
}
int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t hEvent)
{
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
@ -1211,13 +1242,67 @@ int hc_hipRuntimeGetVersion (void *hashcat_ctx, int *runtimeVersion)
return 0;
}
int hc_hipStreamCreate (void *hashcat_ctx, hipStream_t *phStream, unsigned int Flags)
int hc_hipSetDevice (void *hashcat_ctx, hipDevice_t dev)
{
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->hipStreamCreate (phStream, Flags);
const hipError_t HIP_err = hip->hipSetDevice (dev);
if (HIP_err != hipSuccess)
{
const char *pStr = NULL;
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
{
event_log_error (hashcat_ctx, "hipSetDevice(): %s", pStr);
}
else
{
event_log_error (hashcat_ctx, "hipSetDevice(): %d", HIP_err);
}
return -1;
}
return 0;
}
int hc_hipSetDeviceFlags (void *hashcat_ctx, unsigned int flags)
{
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->hipSetDeviceFlags (flags);
if (HIP_err != hipSuccess)
{
const char *pStr = NULL;
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
{
event_log_error (hashcat_ctx, "hipSetDeviceFlags(): %s", pStr);
}
else
{
event_log_error (hashcat_ctx, "hipSetDeviceFlags(): %d", HIP_err);
}
return -1;
}
return 0;
}
int hc_hipStreamCreate (void *hashcat_ctx, hipStream_t *phStream)
{
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->hipStreamCreate (phStream);
if (HIP_err != hipSuccess)
{
@ -1238,6 +1323,33 @@ int hc_hipStreamCreate (void *hashcat_ctx, hipStream_t *phStream, unsigned int F
return 0;
}
int hc_hipStreamCreateWithFlags (void *hashcat_ctx, hipStream_t *phStream, unsigned int Flags)
{
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->hipStreamCreateWithFlags (phStream, Flags);
if (HIP_err != hipSuccess)
{
const char *pStr = NULL;
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
{
event_log_error (hashcat_ctx, "hipStreamCreateWithFlags(): %s", pStr);
}
else
{
event_log_error (hashcat_ctx, "hipStreamCreateWithFlags(): %d", HIP_err);
}
return -1;
}
return 0;
}
int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t hStream)
{
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;

4
src/selftest.c
View File

@ -1272,7 +1272,7 @@ HC_API_CALL void *thread_selftest (void *p)
if (device_param->is_hip == true)
{
if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL;
if (hc_hipSetDevice (hashcat_ctx, device_param->hip_device) == -1) return NULL;
}
const int rc_selftest = process_selftest (hashcat_ctx, device_param);
@ -1303,8 +1303,6 @@ HC_API_CALL void *thread_selftest (void *p)
if (device_param->is_hip == true)
{
if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return NULL;
if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL;
}
if (bridge_ctx->enabled == true)