From 45b8672270ba9d0abf751e74da567fe42c861afb Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sat, 28 Jun 2025 10:42:47 +0200 Subject: [PATCH 1/7] add --machine-readable format to --backend-info --- src/terminal.c | 705 +++++++++++++++++++++++++++++++++++++++++-------- 1 file changed, 593 insertions(+), 112 deletions(-) diff --git a/src/terminal.c b/src/terminal.c index 948b4795d..2976bc7c2 100644 --- a/src/terminal.c +++ b/src/terminal.c @@ -1075,18 +1075,43 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) const user_options_t *user_options = hashcat_ctx->user_options; const folder_config_t *folder_config = hashcat_ctx->folder_config; + if (user_options->machine_readable == true) + { + printf ("{ "); + } + if (user_options->backend_info > 1) { - event_log_info (hashcat_ctx, "System Info:"); - event_log_info (hashcat_ctx, "============"); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "System Info:"); + event_log_info (hashcat_ctx, "============"); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"SystemInfo\": { "); + } #if defined (_WIN) || defined (__CYGWIN__) || defined (__MSYS__) // TODO - event_log_info (hashcat_ctx, "OS.Name......: Windows"); - event_log_info (hashcat_ctx, "OS.Release...: N/A"); - event_log_info (hashcat_ctx, "HW.Platform..: N/A"); - event_log_info (hashcat_ctx, "HW.Model.....: N/A"); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OS.Name......: Windows"); + event_log_info (hashcat_ctx, "OS.Release...: N/A"); + event_log_info (hashcat_ctx, "HW.Platform..: N/A"); + event_log_info (hashcat_ctx, "HW.Model.....: N/A"); + } + else + { + printf ("\"OS\": { "); + printf ("\"Name\": \"%s\", ", "Windows"); + printf ("\"Release\": \"%s\" }, ", "N/A"); + printf ("\"Hardware\": { "); + printf ("\"Platform\": \"%s\", ", "N/A"); + printf ("\"Model\": \"%s\" } ", "N/A"); + printf ("}, "); + } #else struct utsname utsbuf; @@ -1123,10 +1148,23 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) rc_uname = true; } - event_log_info (hashcat_ctx, "OS.Name......: %s", (rc_uname == true) ? utsbuf.sysname : "N/A"); - event_log_info (hashcat_ctx, "OS.Release...: %s", (rc_uname == true) ? utsbuf.release : "N/A"); - event_log_info (hashcat_ctx, "HW.Model.....: %s", (rc_sysctl == true) ? hw_model_buf : "N/A"); - event_log_info (hashcat_ctx, "HW.Platform..: %s", (rc_uname == true) ? utsbuf.machine : "N/A"); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OS.Name......: %s", (rc_uname == true) ? utsbuf.sysname : "N/A"); + event_log_info (hashcat_ctx, "OS.Release...: %s", (rc_uname == true) ? utsbuf.release : "N/A"); + event_log_info (hashcat_ctx, "HW.Platform..: %s", (rc_uname == true) ? utsbuf.machine : "N/A"); + event_log_info (hashcat_ctx, "HW.Model.....: %s", (rc_sysctl == true) ? hw_model_buf : "N/A"); + } + else + { + printf ("\"OS\": { "); + printf ("\"Name\": \"%s\", ", (rc_uname == true) ? utsbuf.sysname : "N/A"); + printf ("\"Release\": \"%s\" }, ", (rc_uname == true) ? utsbuf.release : "N/A"); + printf ("\"Hardware\": { "); + printf ("\"Platform\": \"%s\", ", (rc_uname == true) ? utsbuf.machine : "N/A"); + printf ("\"Model\": \"%s\" } ", (rc_sysctl == true) ? hw_model_buf : "N/A"); + printf ("}, "); + } if (rc_sysctl == true) { @@ -1134,38 +1172,72 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) } #endif // _WIN || __CYGWIN__ || __MSYS__ - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, NULL); - event_log_info (hashcat_ctx, "Environment Info:"); - event_log_info (hashcat_ctx, "================="); - event_log_info (hashcat_ctx, NULL); + event_log_info (hashcat_ctx, "Environment Info:"); + event_log_info (hashcat_ctx, "================="); + event_log_info (hashcat_ctx, NULL); - event_log_info (hashcat_ctx, "Cur.Work.Dir.: %s", folder_config->cwd); - event_log_info (hashcat_ctx, "Install.Dir..: %s", folder_config->install_dir); - event_log_info (hashcat_ctx, "Profile.Dir..: %s", folder_config->profile_dir); - event_log_info (hashcat_ctx, "Cache.Dir....: %s", folder_config->cache_dir); - // uninitialized at this point, for instance if the user uses --session - //event_log_info (hashcat_ctx, "Session.Dir..: %s", folder_config->session_dir); - event_log_info (hashcat_ctx, "Shared.Dir...: %s", folder_config->shared_dir); - event_log_info (hashcat_ctx, "CL.Inc.Path..: %s", folder_config->cpath_real); + event_log_info (hashcat_ctx, "Cur.Work.Dir.: %s", folder_config->cwd); + event_log_info (hashcat_ctx, "Install.Dir..: %s", folder_config->install_dir); + event_log_info (hashcat_ctx, "Profile.Dir..: %s", folder_config->profile_dir); + event_log_info (hashcat_ctx, "Cache.Dir....: %s", folder_config->cache_dir); + // uninitialized at this point, for instance if the user uses --session + //event_log_info (hashcat_ctx, "Session.Dir..: %s", folder_config->session_dir); + event_log_info (hashcat_ctx, "Shared.Dir...: %s", folder_config->shared_dir); + event_log_info (hashcat_ctx, "CL.Inc.Path..: %s", folder_config->cpath_real); - event_log_info (hashcat_ctx, NULL); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"EnvironmentInfo\": { "); + printf ("\"CurrentWorkingDirectory\": \"%s\", ", folder_config->cwd); + printf ("\"InstallDirectory\": \"%s\", ", folder_config->install_dir); + printf ("\"ProfileDirectory\": \"%s\", ", folder_config->profile_dir); + printf ("\"CacheDirectory\": \"%s\", ", folder_config->cache_dir); + printf ("\"SharedDirectory\": \"%s\", ", folder_config->shared_dir); + printf ("\"CLIncludePath\": \"%s\" ", folder_config->cpath_real); + printf ("}, "); + } } if (backend_ctx->cuda) { - event_log_info (hashcat_ctx, "CUDA Info:"); - event_log_info (hashcat_ctx, "=========="); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "CUDA Info:"); + event_log_info (hashcat_ctx, "=========="); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"CUDAInfo\": { "); + } int cuda_devices_cnt = backend_ctx->cuda_devices_cnt; int cuda_driver_version = backend_ctx->cuda_driver_version; - event_log_info (hashcat_ctx, "CUDA.Version.: %u.%u", cuda_driver_version / 1000, (cuda_driver_version % 100) / 10); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "CUDA.Version.: %u.%u", cuda_driver_version / 1000, (cuda_driver_version % 100) / 10); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%u.%u\", ", cuda_driver_version / 1000, (cuda_driver_version % 100) / 10); + printf ("\"BackendDevices\": [ "); + } for (int cuda_devices_idx = 0; cuda_devices_idx < cuda_devices_cnt; cuda_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_cuda[cuda_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1184,29 +1256,88 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); - event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%u\", ", device_maxclock_frequency); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + printf ("\"PCI.Addr.BDFe\": \"%04x:%02x:%02x.%u\" ", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + } + + if (user_options->machine_readable == true) + { + if ((cuda_devices_idx + 1) < cuda_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if (backend_ctx->hip || backend_ctx->mtl || backend_ctx->ocl) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } if (backend_ctx->hip) { - event_log_info (hashcat_ctx, "HIP Info:"); - event_log_info (hashcat_ctx, "========="); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "HIP Info:"); + event_log_info (hashcat_ctx, "========="); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"HIPInfo\": { "); + } int hip_devices_cnt = backend_ctx->hip_devices_cnt; int hip_runtimeVersion = backend_ctx->hip_runtimeVersion; @@ -1217,17 +1348,41 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) int hip_version_minor = (hip_runtimeVersion - (hip_version_major * 10000000)) / 100000; int hip_version_patch = (hip_runtimeVersion - (hip_version_major * 10000000) - (hip_version_minor * 100000)); - event_log_info (hashcat_ctx, "HIP.Version.: %u.%u.%u", hip_version_major, hip_version_minor, hip_version_patch); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "HIP.Version.: %u.%u.%u", hip_version_major, hip_version_minor, hip_version_patch); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%u.%u.%u\", ", hip_version_major, hip_version_minor, hip_version_patch); + } } else { - event_log_info (hashcat_ctx, "HIP.Version.: %u.%u", hip_runtimeVersion / 100, hip_runtimeVersion % 10); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "HIP.Version.: %u.%u", hip_runtimeVersion / 100, hip_runtimeVersion % 10); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%u.%u\", ", hip_runtimeVersion / 100, hip_runtimeVersion % 10); + } + } + + if (user_options->machine_readable == true) + { + printf ("\"BackendDevices\": [ "); } for (int hip_devices_idx = 0; hip_devices_idx < hip_devices_cnt; hip_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_hip[hip_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1246,40 +1401,116 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); - event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%u\", ", device_maxclock_frequency); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + printf ("\"PCI.Addr.BDFe\": \"%04x:%02x:%02x.%u\" ", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + } + + if (user_options->machine_readable == true) + { + if ((hip_devices_idx + 1) < hip_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if (backend_ctx->mtl || backend_ctx->ocl) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } #if defined (__APPLE__) if (backend_ctx->mtl) { - event_log_info (hashcat_ctx, "Metal Info:"); - event_log_info (hashcat_ctx, "==========="); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Metal Info:"); + event_log_info (hashcat_ctx, "==========="); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"MetalInfo\": { "); + } int metal_devices_cnt = backend_ctx->metal_devices_cnt; char *metal_runtimeVersionStr = backend_ctx->metal_runtimeVersionStr; - event_log_info (hashcat_ctx, "Metal.Version.: %s", metal_runtimeVersionStr); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Metal.Version.: %s", metal_runtimeVersionStr); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%s\", ", metal_runtimeVersionStr); + } + + if (user_options->machine_readable == true) + { + printf ("\"BackendDevices\": [ "); + } for (int metal_devices_idx = 0; metal_devices_idx < metal_devices_cnt; metal_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_metal[metal_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1310,30 +1541,111 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); - event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); - event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: N/A"); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); + event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: N/A"); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + } + else + { + printf ("\"Type\": \"%s\", ", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + printf ("\"VendorID\": \"%u\", ", opencl_device_vendor_id); + printf ("\"Vendor\": \"%s\", ", opencl_device_vendor); + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%s\", ", "N/A"); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryAllocPerBlock\": \"%" PRIu64 " MB\", ", device_maxmem_alloc / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + } switch (device_physical_location) { - case MTL_DEVICE_LOCATION_BUILTIN: event_log_info (hashcat_ctx, " Phys.Location..: built-in"); break; - case MTL_DEVICE_LOCATION_SLOT: event_log_info (hashcat_ctx, " Phys.Location..: connected to slot %u", device_location_number); break; - case MTL_DEVICE_LOCATION_EXTERNAL: event_log_info (hashcat_ctx, " Phys.Location..: connected via an external interface (port %u)", device_location_number); break; - case MTL_DEVICE_LOCATION_UNSPECIFIED: event_log_info (hashcat_ctx, " Phys.Location..: unspecified"); break; - default: event_log_info (hashcat_ctx, " Phys.Location..: N/A"); break; + case MTL_DEVICE_LOCATION_BUILTIN: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: built-in"); + } + else + { + printf ("\"PhysicalLocation\": \"built-in\", "); + } + + break; + case MTL_DEVICE_LOCATION_SLOT: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: connected to slot %u", device_location_number); + } + else + { + printf ("\"PhysicalLocation\": \"connected to slot %u\", ", device_location_number); + } + + break; + case MTL_DEVICE_LOCATION_EXTERNAL: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: connected via an external interface (port %u)", device_location_number); + } + else + { + printf ("\"PhysicalLocation\": \"connected via an external interface (port %u)\", ", device_location_number); + } + + break; + case MTL_DEVICE_LOCATION_UNSPECIFIED: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: unspecified"); + } + else + { + printf ("\"PhysicalLocation\": \"unspecified\", "); + } + + break; + default: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: N/A"); + } + else + { + printf ("\"PhysicalLocation\": \"%s\", ", "N/A"); + } + + break; } /* @@ -1347,28 +1659,92 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) } */ - event_log_info (hashcat_ctx, " Registry.ID....: %u", device_registryID); - - if (device_physical_location != MTL_DEVICE_LOCATION_BUILTIN) + if (user_options->machine_readable == false) { - event_log_info (hashcat_ctx, " Max.TX.Rate....: %u MB/sec", device_max_transfer_rate); + event_log_info (hashcat_ctx, " Registry.ID....: %u", device_registryID); } else { - event_log_info (hashcat_ctx, " Max.TX.Rate....: N/A"); + printf ("\"RegistryID\": \"%u\", ", device_registryID); } - event_log_info (hashcat_ctx, " GPU.Properties.: headless %u, low-power %u, removable %u", device_is_headless, device_is_low_power, device_is_removable); - event_log_info (hashcat_ctx, NULL); + if (device_physical_location != MTL_DEVICE_LOCATION_BUILTIN) + { + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Max.TX.Rate....: %u MB/sec", device_max_transfer_rate); + } + else + { + printf ("\"MaxTXRate\": \"%u MB/sec\", ", device_max_transfer_rate); + } + } + else + { + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Max.TX.Rate....: N/A"); + } + else + { + printf ("\"MaxTXRate\": \"%s\", ", "N/A"); + } + } + + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " GPU.Properties.: headless %u, low-power %u, removable %u", device_is_headless, device_is_low_power, device_is_removable); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"GPUProperties\": { "); + printf ("\"headless\": \"%u\", ", device_is_headless); + printf ("\"low_power\": \"%u\", ", device_is_low_power); + printf ("\"removable\": \"%u\" ", device_is_removable); + printf ("} "); + } + + if (user_options->machine_readable == true) + { + if ((metal_devices_idx + 1) < metal_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if (backend_ctx->ocl) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } #endif if (backend_ctx->ocl) { - event_log_info (hashcat_ctx, "OpenCL Info:"); - event_log_info (hashcat_ctx, "============"); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OpenCL Info:"); + event_log_info (hashcat_ctx, "============"); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"OpenCLInfo\": { "); + printf ("\"Platforms\": [ "); + } cl_uint opencl_platforms_cnt = backend_ctx->opencl_platforms_cnt; cl_uint *opencl_platforms_devices_cnt = backend_ctx->opencl_platforms_devices_cnt; @@ -1378,19 +1754,44 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) for (cl_uint opencl_platforms_idx = 0; opencl_platforms_idx < opencl_platforms_cnt; opencl_platforms_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + char *opencl_platform_vendor = opencl_platforms_vendor[opencl_platforms_idx]; char *opencl_platform_name = opencl_platforms_name[opencl_platforms_idx]; char *opencl_platform_version = opencl_platforms_version[opencl_platforms_idx]; cl_uint opencl_platform_devices_cnt = opencl_platforms_devices_cnt[opencl_platforms_idx]; - event_log_info (hashcat_ctx, "OpenCL Platform ID #%u", opencl_platforms_idx + 1); - event_log_info (hashcat_ctx, " Vendor..: %s", opencl_platform_vendor); - event_log_info (hashcat_ctx, " Name....: %s", opencl_platform_name); - event_log_info (hashcat_ctx, " Version.: %s", opencl_platform_version); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OpenCL Platform ID #%u", opencl_platforms_idx + 1); + event_log_info (hashcat_ctx, " Vendor..: %s", opencl_platform_vendor); + event_log_info (hashcat_ctx, " Name....: %s", opencl_platform_name); + event_log_info (hashcat_ctx, " Version.: %s", opencl_platform_version); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"PlatformID\": \"%u\", ", opencl_platforms_idx + 1); + printf ("\"Vendor\": \"%s\", ", opencl_platform_vendor); + printf ("\"Name\": \"%s\", ", opencl_platform_name); + printf ("\"Version\": \"%s\", ", opencl_platform_version); + } + + if (user_options->machine_readable == true) + { + printf ("\"BackendDevices\": [ "); + } for (cl_uint opencl_platform_devices_idx = 0; opencl_platform_devices_idx < opencl_platform_devices_cnt; opencl_platform_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_opencl_platform[opencl_platforms_idx][opencl_platform_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1412,25 +1813,58 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, " Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, " Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); - event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); - event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Version........: %s", opencl_device_version); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); - event_log_info (hashcat_ctx, " OpenCL.Version.: %s", opencl_device_c_version); - event_log_info (hashcat_ctx, " Driver.Version.: %s", opencl_driver_version); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); + event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Version........: %s", opencl_device_version); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + event_log_info (hashcat_ctx, " OpenCL.Version.: %s", opencl_device_c_version); + event_log_info (hashcat_ctx, " Driver.Version.: %s", opencl_driver_version); + } + else + { + printf ("\"Type\": \"%s\", ", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + printf ("\"VendorID\": \"%u\", ", opencl_device_vendor_id); + printf ("\"Vendor\": \"%s\", ", opencl_device_vendor); + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%u\", ", device_maxclock_frequency); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryAllocPerBlock\": \"%" PRIu64 " MB\", ", device_maxmem_alloc / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + printf ("\"OpenCLVersion\": \"%s\", ", opencl_device_c_version); + printf ("\"DriverVersion\": \"%s\" ", opencl_device_version); + } if (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU) { @@ -1440,18 +1874,68 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if ((device_param->opencl_platform_vendor_id == VENDOR_ID_AMD) && (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)) { - event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + } + else + { + printf (", \"PCI.Addr.BDF\": \"%02x:%02x.%u\" ", pcie_bus, pcie_device, pcie_function); + } } if ((device_param->opencl_platform_vendor_id == VENDOR_ID_NV) && (device_param->opencl_device_vendor_id == VENDOR_ID_NV)) { - event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + } + else + { + printf (", \"PCI.Addr.BDF\": \"%02x:%02x.%u\" ", pcie_bus, pcie_device, pcie_function); + } } } - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, NULL); + } + else + { + if ((opencl_platform_devices_idx + 1) < opencl_platform_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if ((opencl_platforms_idx + 1) < opencl_platforms_cnt) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } + + if (user_options->machine_readable == true) + { + printf ("] } "); + } + } + + if (user_options->machine_readable == true) + { + printf ("}"); } } @@ -2676,10 +3160,7 @@ void status_display (hashcat_ctx_t *hashcat_ctx) digests_remain, digests_remain_percent); } - } - if (hashcat_status->digests_cnt > 1000) - { event_log_info (hashcat_ctx, "Recovered/Time...: %s", hashcat_status->cpt); From 92b2f996b29fd72f79b82fae18d774b4f683d379 Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sat, 28 Jun 2025 10:45:24 +0200 Subject: [PATCH 2/7] update docs/changes.txt --- docs/changes.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/docs/changes.txt b/docs/changes.txt index d9ff56f2b..5ff9f0f63 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -125,6 +125,7 @@ - Apple Driver: Updated requirements to use Apple OpenCL API to macOS 13.0 - use - Backend: Updated filename chksum format to prevent invalid cache on Apple Silicon when switching arch - 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 - Dependencies: Added sse2neon v1.8.0 (commit 658eeac) - Dependencies: Updated LZMA SDK to 24.09 From c275c35cedd9817e237652c06af48cdab46a9a8f Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sat, 28 Jun 2025 22:54:36 +0200 Subject: [PATCH 3/7] workaround for HIP bug and avoiding a potential same bug on CUDA --- src/backend.c | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/src/backend.c b/src/backend.c index 15e4badbb..fc0051dd4 100644 --- a/src/backend.c +++ b/src/backend.c @@ -15948,6 +15948,12 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { u32 threads_per_block_with_regs = (floor) ((float) device_param->regsPerBlock / num_regs); + if (threads_per_block_with_regs == 0) + { + // prevent threads_per_block from resulting in 0 due to a bug on the runtime + threads_per_block_with_regs = threads_per_block; + } + if (threads_per_block_with_regs > device_param->kernel_preferred_wgs_multiple) threads_per_block_with_regs -= threads_per_block_with_regs % device_param->kernel_preferred_wgs_multiple; threads_per_block = MIN (threads_per_block, threads_per_block_with_regs); @@ -15967,6 +15973,14 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { u32 threads_per_block_with_regs = (floor) ((float) device_param->regsPerBlock / num_regs); + if (threads_per_block_with_regs == 0) + { + // https://rocm.docs.amd.com/projects/HIP/en/docs-develop/doxygen/html/bug.html + // HIP-Clang always returns 0 for regsPerBlock due to a known bug + // prevent threads_per_block from resulting in 0, otherwise hashcat crashes + threads_per_block_with_regs = threads_per_block; + } + if (threads_per_block_with_regs > device_param->kernel_preferred_wgs_multiple) threads_per_block_with_regs -= threads_per_block_with_regs % device_param->kernel_preferred_wgs_multiple; threads_per_block = MIN (threads_per_block, threads_per_block_with_regs); From 0c2ed0d1991e97c1ae76006adfee656e637a93f2 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Sun, 29 Jun 2025 14:39:14 +0200 Subject: [PATCH 4/7] Update plugins that benefit from an artificially limited register count (NVIDIA). Update default hash settings to 64MiB:3:4 for Argon2 in -m 70000, following RFC 9106 recommendations. Add option OPTS_TYPE_THREAD_MULTI_DISABLE: allows plugin developers to disable scaling the password candidate batch size based on device thread count. This can be useful for super slow hash algorithms that utilize threads differently, e.g., when the algorithm allows parallelization. Note: thread count for the device can still be set normally. Add options OPTI_TYPE_SLOW_HASH_DIMY_INIT/LOOP/COMP: enable 2D launches for slow hash init/loop/comp kernel with dimensions X and Y. The Y value must be set via salt->salt_dimy attribute. Change autotune kernel-loops start value to the lowest multiple of the target hash iteration count, if kernel_loops_min permits. Fixed a bug in autotune where kernel_threads_max was not respected during initial init and loop-prepare kernel runs. --- OpenCL/inc_types.h | 1 + include/types.h | 18 ++++--- src/autotune.c | 22 ++++---- src/backend.c | 101 +++++++++++++++++++++++++++++------ src/modules/module_01800.c | 1 + src/modules/module_06400.c | 1 + src/modules/module_06800.c | 1 + src/modules/module_07700.c | 1 + src/modules/module_07701.c | 1 + src/modules/module_08300.c | 3 +- src/modules/module_13751.c | 3 +- src/modules/module_13752.c | 3 +- src/modules/module_13753.c | 3 +- src/modules/module_13761.c | 3 +- src/modules/module_13762.c | 3 +- src/modules/module_13763.c | 3 +- src/modules/module_14800.c | 1 + src/modules/module_14900.c | 3 +- src/modules/module_15600.c | 1 + src/modules/module_16300.c | 1 + src/modules/module_16700.c | 1 + src/modules/module_16900.c | 1 + src/modules/module_18100.c | 1 + src/modules/module_20800.c | 1 + src/modules/module_21300.c | 1 + src/modules/module_21500.c | 1 + src/modules/module_21700.c | 1 + src/modules/module_22100.c | 3 +- src/modules/module_22300.c | 1 + src/modules/module_22921.c | 3 +- src/modules/module_22941.c | 3 +- src/modules/module_23400.c | 1 + src/modules/module_23600.c | 1 + src/modules/module_23800.c | 3 +- src/modules/module_24200.c | 1 + src/modules/module_24420.c | 1 + src/modules/module_25500.c | 1 + src/modules/module_25900.c | 3 +- src/modules/module_26000.c | 1 + src/modules/module_26100.c | 1 + src/modules/module_26600.c | 1 + src/modules/module_26700.c | 3 +- src/modules/module_26800.c | 3 +- src/modules/module_26900.c | 1 + src/modules/module_27300.c | 1 + src/modules/module_27400.c | 1 + src/modules/module_27500.c | 1 + src/modules/module_27600.c | 1 + src/modules/module_29451.c | 3 +- src/modules/module_29452.c | 3 +- src/modules/module_29453.c | 3 +- src/modules/module_29461.c | 3 +- src/modules/module_29462.c | 3 +- src/modules/module_29463.c | 3 +- src/modules/module_70000.c | 2 +- src/shared.c | 6 +++ tools/test_modules/m70000.pm | 2 +- 57 files changed, 190 insertions(+), 54 deletions(-) diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h index 233e28958..a13c89b8e 100644 --- a/OpenCL/inc_types.h +++ b/OpenCL/inc_types.h @@ -2008,6 +2008,7 @@ typedef struct salt u32 salt_len_pc; u32 salt_iter; u32 salt_iter2; + u32 salt_dimy; u32 salt_sign[2]; u32 salt_repeats; diff --git a/include/types.h b/include/types.h index 22c57d85a..600e7bd93 100644 --- a/include/types.h +++ b/include/types.h @@ -412,6 +412,9 @@ typedef enum opti_type OPTI_TYPE_REGISTER_LIMIT = (1 << 20), // We'll limit the register count to 128 OPTI_TYPE_SLOW_HASH_SIMD_INIT2 = (1 << 21), OPTI_TYPE_SLOW_HASH_SIMD_LOOP2 = (1 << 22), + OPTI_TYPE_SLOW_HASH_DIMY_INIT = (1 << 23), + OPTI_TYPE_SLOW_HASH_DIMY_LOOP = (1 << 24), + OPTI_TYPE_SLOW_HASH_DIMY_COMP = (1 << 25), } opti_type_t; @@ -476,14 +479,17 @@ typedef enum opts_type OPTS_TYPE_DYNAMIC_SHARED = (1ULL << 53), // use dynamic shared memory (note: needs special kernel changes) OPTS_TYPE_SELF_TEST_DISABLE = (1ULL << 54), // some algos use JiT in combinations with a salt or create too much startup time OPTS_TYPE_MP_MULTI_DISABLE = (1ULL << 55), // do not multiply the kernel-accel with the multiprocessor count per device to allow more fine-tuned workload settings - OPTS_TYPE_NATIVE_THREADS = (1ULL << 56), // forces "native" thread count: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefront), GPU-NV=32 (warps) - OPTS_TYPE_MAXIMUM_THREADS = (1ULL << 57), // disable else branch in pre-compilation thread count optimization setting - OPTS_TYPE_POST_AMP_UTF16LE = (1ULL << 58), // run the utf8 to utf16le conversion kernel after they have been processed from amplifiers + OPTS_TYPE_THREAD_MULTI_DISABLE // do not multiply the kernel-power with the thread count per device for super slow algos + = (1ULL << 56), + OPTS_TYPE_NATIVE_THREADS = (1ULL << 57), // forces "native" thread count: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefront), GPU-NV=32 (warps) + OPTS_TYPE_MAXIMUM_THREADS = (1ULL << 58), // disable else branch in pre-compilation thread count optimization setting + OPTS_TYPE_POST_AMP_UTF16LE = (1ULL << 59), // run the utf8 to utf16le conversion kernel after they have been processed from amplifiers OPTS_TYPE_AUTODETECT_DISABLE - = (1ULL << 59), // skip autodetect engine - OPTS_TYPE_STOCK_MODULE = (1ULL << 60), // module included with hashcat default distribution + = (1ULL << 60), // skip autodetect engine + OPTS_TYPE_STOCK_MODULE = (1ULL << 61), // module included with hashcat default distribution OPTS_TYPE_MULTIHASH_DESPITE_ESALT - = (1ULL << 61), // overrule multihash cracking check same salt but not same esalt + = (1ULL << 62), // overrule multihash cracking check same salt but not same esalt + OPTS_TYPE_MAXIMUM_ACCEL = (1ULL << 63) // try to maximize kernel-accel during autotune } opts_type_t; diff --git a/src/autotune.c b/src/autotune.c index 87637b29d..065c0a217 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -43,7 +43,8 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par device_param->kernel_param.loop_cnt = kernel_loops; // not a bug, both need to be set device_param->kernel_param.il_cnt = kernel_loops; // because there's two variables for inner iters for slow and fast hashes - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads); u32 kernel_power_try = hardware_power * kernel_accel; @@ -133,7 +134,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_accel = kernel_accel_min; device_param->kernel_loops = kernel_loops_min; device_param->kernel_threads = kernel_threads_min; - device_param->hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads_min; + device_param->hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads_min); device_param->kernel_power = device_param->hardware_power * kernel_accel_min; } @@ -212,7 +214,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // from here it's clear we are allowed to autotune // so let's init some fake words - const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads_max; + const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads_max); u32 kernel_power_max = hardware_power_max * kernel_accel_max; @@ -298,13 +301,13 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { const u32 kernel_threads_sav = device_param->kernel_threads; - device_param->kernel_threads = device_param->kernel_wgs1; + device_param->kernel_threads = MIN (device_param->kernel_wgs1, kernel_threads_max); run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_max, false, 0, true); if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE) { - device_param->kernel_threads = device_param->kernel_wgs2p; + device_param->kernel_threads = MIN (device_param->kernel_wgs2p, kernel_threads_max); run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_max, false, 0, true); } @@ -330,8 +333,6 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // v7 autotuner is a lot more straight forward - u32 kernel_loops_min_start = kernel_loops_min; - if (hashes && hashes->st_salts_buf) { u32 start = kernel_loops_max; @@ -348,12 +349,12 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if ((start >= kernel_loops_min) && (start <= kernel_loops_max)) { - kernel_loops_min_start = start; + kernel_loops = start; } } } - for (u32 kernel_loops_test = kernel_loops_min_start; kernel_loops_test <= kernel_loops_max; kernel_loops_test <<= 1) + for (u32 kernel_loops_test = kernel_loops; kernel_loops_test <= kernel_loops_max; kernel_loops_test <<= 1) { double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_test, kernel_threads_min, 2); @@ -564,7 +565,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_loops = kernel_loops; device_param->kernel_threads = kernel_threads; - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * device_param->kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param->kernel_threads); device_param->hardware_power = hardware_power; diff --git a/src/backend.c b/src/backend.c index fc0051dd4..00cac2245 100644 --- a/src/backend.c +++ b/src/backend.c @@ -2598,7 +2598,10 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con if (kernel_threads == 0) kernel_threads = 1; - num_elements = CEILDIV (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = CEILDIV (num_elements, kernel_threads); + } if (kern_run == KERN_RUN_1) { @@ -2636,14 +2639,29 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + u32 gridDimX = num_elements; + u32 gridDimY = 1; + u32 gridDimZ = 1; + + u32 blockDimX = kernel_threads; + u32 blockDimY = 1; + u32 blockDimZ = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if (is_autotune == true) { - if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; } if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event1, device_param->cuda_stream) == -1) return -1; - if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event2, device_param->cuda_stream) == -1) return -1; @@ -2699,7 +2717,10 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con if (kernel_threads == 0) kernel_threads = 1; - num_elements = CEILDIV (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = CEILDIV (num_elements, kernel_threads); + } if (kern_run == KERN_RUN_1) { @@ -2737,14 +2758,31 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + u32 gridDimX = num_elements; + u32 gridDimY = 1; + u32 gridDimZ = 1; + + u32 blockDimX = kernel_threads; + u32 blockDimY = 1; + u32 blockDimZ = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + + //printf ("%d %d %d %d %d %d %d\n", kern_run, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ); + if (is_autotune == true) { - if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; } if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event1, device_param->hip_stream) == -1) return -1; - if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event2, device_param->hip_stream) == -1) return -1; @@ -2984,17 +3022,44 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } - num_elements = round_up_multiple_64 (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = round_up_multiple_64 (num_elements, kernel_threads); + } + else + { + num_elements = num_elements * kernel_threads; + } - const size_t global_work_size[3] = { num_elements, 1, 1 }; - const size_t local_work_size[3] = { kernel_threads, 1, 1 }; + size_t global_work_size[3] = { num_elements, 1, 1 }; + size_t local_work_size[3] = { kernel_threads, 1, 1 }; + + cl_uint work_dim = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } if (is_autotune == true) { - if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; } - if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; // spin damper section @@ -8952,7 +9017,8 @@ void backend_ctx_devices_sync_tuning (hashcat_ctx_t *hashcat_ctx) device_param_dst->kernel_loops = device_param_src->kernel_loops; device_param_dst->kernel_threads = device_param_src->kernel_threads; - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param_dst->device_processors) * device_param_dst->kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param_dst->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param_dst->kernel_threads); device_param_dst->hardware_power = hardware_power; @@ -9522,7 +9588,11 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "-D MAX_THREADS_PER_BLOCK=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-architecture=%s", device_param->gcnArchName); - hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-max-threads-per-block=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); + + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-max-threads-per-block=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); + } // untested but it should work #if defined (_WIN) || defined (__CYGWIN__) || defined (__MSYS__) @@ -10436,7 +10506,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) * device properties */ - const u32 device_processors = device_param->device_processors; + //const u32 device_processors = device_param->device_processors; /** * device threads @@ -16059,7 +16129,8 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) // device_param->kernel_threads = kernel_threads; device_param->kernel_threads = 0; - u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_processors) * device_param->kernel_threads_max; + const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param->kernel_threads_max); u32 kernel_accel_min = device_param->kernel_accel_min; u32 kernel_accel_max = device_param->kernel_accel_max; diff --git a/src/modules/module_01800.c b/src/modules/module_01800.c index fea4851c6..f0d2e9623 100644 --- a/src/modules/module_01800.c +++ b/src/modules/module_01800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OS; static const char *HASH_NAME = "sha512crypt $6$, SHA512 (Unix)"; static const u64 KERN_TYPE = 1800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_06400.c b/src/modules/module_06400.c index b7ac9d87f..a553846c7 100644 --- a/src/modules/module_06400.c +++ b/src/modules/module_06400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OS; static const char *HASH_NAME = "AIX {ssha256}"; static const u64 KERN_TYPE = 6400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_06800.c b/src/modules/module_06800.c index 215f26992..e6cf96fd6 100644 --- a/src/modules/module_06800.c +++ b/src/modules/module_06800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "LastPass + LastPass sniffed"; static const u64 KERN_TYPE = 6800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_07700.c b/src/modules/module_07700.c index 2642a0bb7..1d783d351 100644 --- a/src/modules/module_07700.c +++ b/src/modules/module_07700.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SAP CODVN B (BCODE)"; static const u64 KERN_TYPE = 7700; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_PRECOMPUTE_INIT + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_07701.c b/src/modules/module_07701.c index bc852639b..232d7ca4d 100644 --- a/src/modules/module_07701.c +++ b/src/modules/module_07701.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SAP CODVN B (BCODE) from RFC_READ_TABLE"; static const u64 KERN_TYPE = 7701; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_PRECOMPUTE_INIT + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_08300.c b/src/modules/module_08300.c index c499667ae..80dd3475d 100644 --- a/src/modules/module_08300.c +++ b/src/modules/module_08300.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_5; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_SERVER; static const char *HASH_NAME = "DNSSEC (NSEC3)"; static const u64 KERN_TYPE = 8300; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE | OPTS_TYPE_ST_HEX diff --git a/src/modules/module_13751.c b/src/modules/module_13751.c index dac31a633..ee27acfd6 100644 --- a/src/modules/module_13751.c +++ b/src/modules/module_13751.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit (legacy)"; static const u64 KERN_TYPE = 13751; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13752.c b/src/modules/module_13752.c index e959a80af..095758dd2 100644 --- a/src/modules/module_13752.c +++ b/src/modules/module_13752.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit (legacy)"; static const u64 KERN_TYPE = 13752; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13753.c b/src/modules/module_13753.c index 277afaf02..4eee98625 100644 --- a/src/modules/module_13753.c +++ b/src/modules/module_13753.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit (legacy)"; static const u64 KERN_TYPE = 13753; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13761.c b/src/modules/module_13761.c index e33a693e2..5b1b82b27 100644 --- a/src/modules/module_13761.c +++ b/src/modules/module_13761.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13751; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13762.c b/src/modules/module_13762.c index 0f5f18545..6f1a27929 100644 --- a/src/modules/module_13762.c +++ b/src/modules/module_13762.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13752; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13763.c b/src/modules/module_13763.c index c9e28d4c7..fb50002a2 100644 --- a/src/modules/module_13763.c +++ b/src/modules/module_13763.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13753; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_14800.c b/src/modules/module_14800.c index 9f3c1ca0f..1e4b91a15 100644 --- a/src/modules/module_14800.c +++ b/src/modules/module_14800.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "iTunes backup >= 10.0"; static const u64 KERN_TYPE = 14800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_14900.c b/src/modules/module_14900.c index cbe009024..fcc639688 100644 --- a/src/modules/module_14900.c +++ b/src/modules/module_14900.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_CIPHER_KPA; static const char *HASH_NAME = "Skip32 (PT = $salt, key = $pass)"; static const u64 KERN_TYPE = 14900; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_SUGGEST_KG; diff --git a/src/modules/module_15600.c b/src/modules/module_15600.c index 1f2b7b9a3..66469d9bf 100644 --- a/src/modules/module_15600.c +++ b/src/modules/module_15600.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Ethereum Wallet, PBKDF2-HMAC-SHA256"; static const u64 KERN_TYPE = 15600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_16300.c b/src/modules/module_16300.c index e7507f1fa..6a12a1b39 100644 --- a/src/modules/module_16300.c +++ b/src/modules/module_16300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Ethereum Pre-Sale Wallet, PBKDF2-HMAC-SHA256"; static const u64 KERN_TYPE = 16300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_16700.c b/src/modules/module_16700.c index 9dce98e62..2fec2f557 100644 --- a/src/modules/module_16700.c +++ b/src/modules/module_16700.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "FileVault 2"; static const u64 KERN_TYPE = 16200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_16900.c b/src/modules/module_16900.c index 443da6007..70eca3f16 100644 --- a/src/modules/module_16900.c +++ b/src/modules/module_16900.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Ansible Vault"; static const u64 KERN_TYPE = 16900; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_18100.c b/src/modules/module_18100.c index 46dadfe60..678cb8265 100644 --- a/src/modules/module_18100.c +++ b/src/modules/module_18100.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OTP; static const char *HASH_NAME = "TOTP (HMAC-SHA1)"; static const u64 KERN_TYPE = 18100; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE diff --git a/src/modules/module_20800.c b/src/modules/module_20800.c index 6cece615d..1330b99bc 100644 --- a/src/modules/module_20800.c +++ b/src/modules/module_20800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "sha256(md5($pass))"; static const u64 KERN_TYPE = 20800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_21300.c b/src/modules/module_21300.c index 75e57c1f4..34ebff7f5 100644 --- a/src/modules/module_21300.c +++ b/src/modules/module_21300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "md5($salt.sha1($salt.$pass))"; static const u64 KERN_TYPE = 21300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_21500.c b/src/modules/module_21500.c index dba580a7c..8936be11c 100644 --- a/src/modules/module_21500.c +++ b/src/modules/module_21500.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SolarWinds Orion"; static const u64 KERN_TYPE = 21500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64 + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_21700.c b/src/modules/module_21700.c index 65536027b..e00cf232f 100644 --- a/src/modules/module_21700.c +++ b/src/modules/module_21700.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Electrum Wallet (Salt-Type 4)"; static const u64 KERN_TYPE = 21700; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_22100.c b/src/modules/module_22100.c index 344218243..4d79d30bf 100644 --- a/src/modules/module_22100.c +++ b/src/modules/module_22100.c @@ -20,7 +20,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "BitLocker"; static const u64 KERN_TYPE = 22100; -static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP; +static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_MP_MULTI_DISABLE; diff --git a/src/modules/module_22300.c b/src/modules/module_22300.c index a85d1d753..04a4bc08c 100644 --- a/src/modules/module_22300.c +++ b/src/modules/module_22300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "sha256($salt.$pass.$salt)"; static const u64 KERN_TYPE = 22300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_22921.c b/src/modules/module_22921.c index 163aacee9..c6471a232 100644 --- a/src/modules/module_22921.c +++ b/src/modules/module_22921.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "RSA/DSA/EC/OpenSSH Private Keys ($6$)"; static const u64 KERN_TYPE = 22921; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; diff --git a/src/modules/module_22941.c b/src/modules/module_22941.c index 4b4bf09ac..65e7d7c2b 100644 --- a/src/modules/module_22941.c +++ b/src/modules/module_22941.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "RSA/DSA/EC/OpenSSH Private Keys ($4$)"; static const u64 KERN_TYPE = 22941; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; diff --git a/src/modules/module_23400.c b/src/modules/module_23400.c index 13743815f..c2cc2a2e2 100644 --- a/src/modules/module_23400.c +++ b/src/modules/module_23400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Bitwarden"; static const u64 KERN_TYPE = 23400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_23600.c b/src/modules/module_23600.c index fdb3467d1..0b8994ff7 100644 --- a/src/modules/module_23600.c +++ b/src/modules/module_23600.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "AxCrypt 2 AES-256"; static const u64 KERN_TYPE = 23600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_23800.c b/src/modules/module_23800.c index c6c8fec61..c75a9aaaf 100644 --- a/src/modules/module_23800.c +++ b/src/modules/module_23800.c @@ -20,7 +20,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; // actually only DGST_SIZE_4_ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "RAR3-p (Compressed)"; static const u64 KERN_TYPE = 23800; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_HOOK23 diff --git a/src/modules/module_24200.c b/src/modules/module_24200.c index 8b6a8e5c8..20ea6ccec 100644 --- a/src/modules/module_24200.c +++ b/src/modules/module_24200.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_DATABASE_SERVER; static const char *HASH_NAME = "MongoDB ServerKey SCRAM-SHA-256"; static const u64 KERN_TYPE = 24200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_24420.c b/src/modules/module_24420.c index 39820834a..544e07ee1 100644 --- a/src/modules/module_24420.c +++ b/src/modules/module_24420.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "PKCS#8 Private Keys (PBKDF2-HMAC-SHA256 + 3DES/AES)"; static const u64 KERN_TYPE = 24420; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_25500.c b/src/modules/module_25500.c index afc451fb9..06a1e795d 100644 --- a/src/modules/module_25500.c +++ b/src/modules/module_25500.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Stargazer Stellar Wallet XLM"; static const u64 KERN_TYPE = 25500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_25900.c b/src/modules/module_25900.c index 7ad951f6d..a460a3bd8 100644 --- a/src/modules/module_25900.c +++ b/src/modules/module_25900.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_SERVER; static const char *HASH_NAME = "KNX IP Secure - Device Authentication Code"; static const u64 KERN_TYPE = 25900; -static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP; +static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_DEEP_COMP_KERNEL; diff --git a/src/modules/module_26000.c b/src/modules/module_26000.c index 35a6dd937..8ed90d5ee 100644 --- a/src/modules/module_26000.c +++ b/src/modules/module_26000.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Mozilla key3.db"; static const u64 KERN_TYPE = 26000; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE; diff --git a/src/modules/module_26100.c b/src/modules/module_26100.c index c42958289..3414fcbdb 100644 --- a/src/modules/module_26100.c +++ b/src/modules/module_26100.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Mozilla key4.db"; static const u64 KERN_TYPE = 26100; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26600.c b/src/modules/module_26600.c index 926555e94..92b5c8efc 100644 --- a/src/modules/module_26600.c +++ b/src/modules/module_26600.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "MetaMask Wallet (needs all data, checks AES-GCM tag)"; static const u64 KERN_TYPE = 26600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26700.c b/src/modules/module_26700.c index 63854c0d2..47427ba32 100644 --- a/src/modules/module_26700.c +++ b/src/modules/module_26700.c @@ -21,7 +21,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA224-128"; static const u64 KERN_TYPE = 26700; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26800.c b/src/modules/module_26800.c index 024bfa636..44d72f7df 100644 --- a/src/modules/module_26800.c +++ b/src/modules/module_26800.c @@ -21,7 +21,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_6; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA256-192"; static const u64 KERN_TYPE = 26800; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26900.c b/src/modules/module_26900.c index 8d916ebf2..29e3cdc92 100644 --- a/src/modules/module_26900.c +++ b/src/modules/module_26900.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA384-256"; static const u64 KERN_TYPE = 26900; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS diff --git a/src/modules/module_27300.c b/src/modules/module_27300.c index ff20acc9d..4d607051e 100644 --- a/src/modules/module_27300.c +++ b/src/modules/module_27300.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA512-384"; static const u64 KERN_TYPE = 27300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS diff --git a/src/modules/module_27400.c b/src/modules/module_27400.c index 7d9f0fbd4..059b50cb5 100644 --- a/src/modules/module_27400.c +++ b/src/modules/module_27400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VMware VMX (PBKDF2-HMAC-SHA1 + AES-256-CBC)"; static const u64 KERN_TYPE = 27400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_27500.c b/src/modules/module_27500.c index ebabe24e5..6b79eba3a 100644 --- a/src/modules/module_27500.c +++ b/src/modules/module_27500.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VirtualBox (PBKDF2-HMAC-SHA256 & AES-128-XTS)"; static const u64 KERN_TYPE = 27500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_27600.c b/src/modules/module_27600.c index b4cdade54..3259ce05f 100644 --- a/src/modules/module_27600.c +++ b/src/modules/module_27600.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VirtualBox (PBKDF2-HMAC-SHA256 & AES-256-XTS)"; static const u64 KERN_TYPE = 27600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_29451.c b/src/modules/module_29451.c index c8e10bdcf..fe3c0737f 100644 --- a/src/modules/module_29451.c +++ b/src/modules/module_29451.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit"; static const u64 KERN_TYPE = 13751; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29452.c b/src/modules/module_29452.c index 395839f64..842f36f47 100644 --- a/src/modules/module_29452.c +++ b/src/modules/module_29452.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit"; static const u64 KERN_TYPE = 13752; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29453.c b/src/modules/module_29453.c index 97bf9a333..6162337de 100644 --- a/src/modules/module_29453.c +++ b/src/modules/module_29453.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit"; static const u64 KERN_TYPE = 13753; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29461.c b/src/modules/module_29461.c index be7dcfc2e..ff2705da7 100644 --- a/src/modules/module_29461.c +++ b/src/modules/module_29461.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit + boot-mode"; static const u64 KERN_TYPE = 13751; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29462.c b/src/modules/module_29462.c index 7e9a6535c..14b53d6ad 100644 --- a/src/modules/module_29462.c +++ b/src/modules/module_29462.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit + boot-mode"; static const u64 KERN_TYPE = 13752; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29463.c b/src/modules/module_29463.c index e98c71289..7a5c8ac92 100644 --- a/src/modules/module_29463.c +++ b/src/modules/module_29463.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit + boot-mode"; static const u64 KERN_TYPE = 13753; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_70000.c b/src/modules/module_70000.c index 01504ea09..57bd77d12 100644 --- a/src/modules/module_70000.c +++ b/src/modules/module_70000.c @@ -29,7 +29,7 @@ static const u64 BRIDGE_TYPE = BRIDGE_TYPE_MATCH_TUNINGS // optional - impr | BRIDGE_TYPE_REPLACE_LOOP; static const char *BRIDGE_NAME = "argon2id_reference"; static const char *ST_PASS = "hashcat"; -static const char *ST_HASH = "$argon2id$v=19$m=4096,t=3,p=1$FoIjFnZlM2JSJWYXUgMFAw$eYKMzhbW8uyT1LLtKRdRcJj2CQeRrdr2pKv/Y71YbAQ"; +static const char *ST_HASH = "$argon2id$v=19$m=65536,t=3,p=1$FBMjI4RJBhIykCgol1KEJA$2ky5GAdhT1kH4kIgPN/oERE3Taiy43vNN70a3HpiKQU"; u32 module_attack_exec (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC; } u32 module_dgst_pos0 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0; } diff --git a/src/shared.c b/src/shared.c index a00025352..f7dbcd78d 100644 --- a/src/shared.c +++ b/src/shared.c @@ -104,6 +104,9 @@ static const char *const OPTI_STR_USES_BITS_8 = "Uses-8-Bit"; static const char *const OPTI_STR_USES_BITS_16 = "Uses-16-Bit"; static const char *const OPTI_STR_USES_BITS_32 = "Uses-32-Bit"; static const char *const OPTI_STR_USES_BITS_64 = "Uses-64-Bit"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_INIT = "Slow-Hash-DimensionY-INIT"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_COMP = "Slow-Hash-DimensionY-LOOP"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_LOOP = "Slow-Hash-DimensionY-COMP"; static const char *const HASH_CATEGORY_UNDEFINED_STR = "Undefined"; static const char *const HASH_CATEGORY_RAW_HASH_STR = "Raw Hash"; @@ -1072,6 +1075,9 @@ const char *stroptitype (const u32 opti_type) case OPTI_TYPE_SLOW_HASH_SIMD_LOOP: return OPTI_STR_SLOW_HASH_SIMD_LOOP; case OPTI_TYPE_SLOW_HASH_SIMD_LOOP2: return OPTI_STR_SLOW_HASH_SIMD_LOOP2; case OPTI_TYPE_SLOW_HASH_SIMD_COMP: return OPTI_STR_SLOW_HASH_SIMD_COMP; + case OPTI_TYPE_SLOW_HASH_DIMY_INIT: return OPTI_STR_SLOW_HASH_DIMY_INIT; + case OPTI_TYPE_SLOW_HASH_DIMY_LOOP: return OPTI_STR_SLOW_HASH_DIMY_LOOP; + case OPTI_TYPE_SLOW_HASH_DIMY_COMP: return OPTI_STR_SLOW_HASH_DIMY_COMP; case OPTI_TYPE_USES_BITS_8: return OPTI_STR_USES_BITS_8; case OPTI_TYPE_USES_BITS_16: return OPTI_STR_USES_BITS_16; case OPTI_TYPE_USES_BITS_32: return OPTI_STR_USES_BITS_32; diff --git a/tools/test_modules/m70000.pm b/tools/test_modules/m70000.pm index 4b44334ac..72861ae49 100644 --- a/tools/test_modules/m70000.pm +++ b/tools/test_modules/m70000.pm @@ -17,7 +17,7 @@ sub module_generate_hash { my $word = shift; my $salt = shift; - my $m = shift // 4096; + my $m = shift // 65536; my $t = shift // 3; my $p = shift // 1; my $len = shift // random_number (1, 2) * 16; From ca7111996863cbc0552c3a9e016cecfa20480f8f Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sun, 29 Jun 2025 16:59:39 +0200 Subject: [PATCH 5/7] Selftest: rename selftest function to process_selftest and splitting into 3 smaller functions --- docs/changes.txt | 1 + src/selftest.c | 72 ++++++++++++++++++++++++++++++++++-------------- 2 files changed, 53 insertions(+), 20 deletions(-) diff --git a/docs/changes.txt b/docs/changes.txt index 7e15685c3..f4dc61531 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -141,6 +141,7 @@ - Modules: Updated module_unstable_warning - Open Document Format: Added support for small documents with content length < 1024 - OpenCL Backend: added workaround to set device_available_memory from CUDA/HIP alias device +- Selftest: rename selftest function to process_selftest and splitting into 3 smaller functions - Status Code: Add specific return code for self-test fail (-11) - Scrypt: Increase buffer sizes in module for hash mode 8900 to allow longer scrypt digests - Unicode: Update UTF-8 to UTF-16 conversion to match RFC 3629 diff --git a/src/selftest.c b/src/selftest.c index f1b2b2dd5..aea6f3128 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -12,18 +12,14 @@ #include "thread.h" #include "selftest.h" -static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) +static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u32 *highest_pw_len) { - bridge_ctx_t *bridge_ctx = hashcat_ctx->bridge_ctx; - hashconfig_t *hashconfig = hashcat_ctx->hashconfig; hashes_t *hashes = hashcat_ctx->hashes; module_ctx_t *module_ctx = hashcat_ctx->module_ctx; - status_ctx_t *status_ctx = hashcat_ctx->status_ctx; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; user_options_t *user_options = hashcat_ctx->user_options; user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; - if (hashconfig->st_hash == NULL) return 0; - // init : replace hashes with selftest hash if (device_param->is_cuda == true) @@ -86,8 +82,6 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param pw_t comb; bf_t bf; - u32 highest_pw_len = 0; - if (user_options->slow_candidates == true) { if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) @@ -460,7 +454,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; } - highest_pw_len = pw.pw_len; + *highest_pw_len = pw.pw_len; } } } @@ -500,6 +494,16 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param } } + return 0; +} + +static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u32 highest_pw_len) +{ + bridge_ctx_t *bridge_ctx = hashcat_ctx->bridge_ctx; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; + hashes_t *hashes = hashcat_ctx->hashes; + module_ctx_t *module_ctx = hashcat_ctx->module_ctx; + // main : run the kernel const u32 kernel_threads_sav = device_param->kernel_threads; @@ -933,22 +937,28 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_threads = kernel_threads_sav; - // check : check if cracked + return 0; +} - u32 num_cracked = 0; +static int selftest_cleanup (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u32 *num_cracked) +{ + user_options_t *user_options = hashcat_ctx->user_options; + user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; + + // check : check if cracked cl_event opencl_event; if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoHAsync (hashcat_ctx, num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1; if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event3, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoHAsync (hashcat_ctx, num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1; if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event3, device_param->hip_stream) == -1) return -1; } @@ -956,13 +966,13 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param #if defined (__APPLE__) if (device_param->is_metal == true) { - if (hc_mtlMemcpyDtoH (hashcat_ctx, device_param->metal_command_queue, &num_cracked, device_param->metal_d_result, 0, sizeof (u32)) == -1) return -1; + if (hc_mtlMemcpyDtoH (hashcat_ctx, device_param->metal_command_queue, num_cracked, device_param->metal_d_result, 0, sizeof (u32)) == -1) return -1; } #endif if (device_param->is_opencl == true) { - if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_FALSE, 0, sizeof (u32), &num_cracked, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_FALSE, 0, sizeof (u32), num_cracked, 0, NULL, &opencl_event) == -1) return -1; if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; } @@ -1154,7 +1164,27 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_clReleaseEvent (hashcat_ctx, opencl_event) == -1) return -1; } + return 0; +} + +static int process_selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) +{ + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; + status_ctx_t *status_ctx = hashcat_ctx->status_ctx; + + if (hashconfig->st_hash == NULL) return 0; + + u32 highest_pw_len = 0; + u32 num_cracked = 0; + + if (selftest_init (hashcat_ctx, device_param, &highest_pw_len) == -1) return -1; + + if (selftest_run_kernel (hashcat_ctx, device_param, highest_pw_len) == -1) return -1; + + if (selftest_cleanup (hashcat_ctx, device_param, &num_cracked) == -1) return -1; + // check return + if (num_cracked == 0) { hc_thread_mutex_lock (status_ctx->mux_display); @@ -1169,20 +1199,22 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! HIP kernel self-test failed.", device_param->device_id + 1); } - #if defined (__APPLE__) if (device_param->is_metal == true) { event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! Metal kernel self-test failed.", device_param->device_id + 1); } - #endif if (device_param->is_opencl == true) { event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! OpenCL kernel self-test failed.", device_param->device_id + 1); } - event_log_warning (hashcat_ctx, "Your device driver installation is probably broken."); - event_log_warning (hashcat_ctx, "See also: https://hashcat.net/faq/wrongdriver"); + if (device_param->is_metal == false) + { + event_log_warning (hashcat_ctx, "Your device driver installation is probably broken."); + event_log_warning (hashcat_ctx, "See also: https://hashcat.net/faq/wrongdriver"); + } + event_log_warning (hashcat_ctx, NULL); hc_thread_mutex_unlock (status_ctx->mux_display); @@ -1232,7 +1264,7 @@ HC_API_CALL void *thread_selftest (void *p) if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL; } - const int rc_selftest = selftest (hashcat_ctx, device_param); + const int rc_selftest = process_selftest (hashcat_ctx, device_param); if (user_options->benchmark == true) { From f848163b0ac594f43695a716167d75a798d7eddc Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sun, 29 Jun 2025 23:08:43 +0200 Subject: [PATCH 6/7] Fixed race condition in selftest_init on OpenCL with non-blocking write --- docs/changes.txt | 1 + src/selftest.c | 27 +++++++++++++++++++-------- 2 files changed, 20 insertions(+), 8 deletions(-) diff --git a/docs/changes.txt b/docs/changes.txt index f4dc61531..2d8b5007a 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -92,6 +92,7 @@ - Fixed bug in grep out-of-memory workaround on Unit Test - Fixed bug in input_tokenizer when TOKEN_ATTR_FIXED_LENGTH is used and refactor modules - Fixed bug in --stdout that caused certain rules to malfunction +- Fixed race condition in selftest_init on OpenCL with non-blocking write - Fixed build failed for 10700 optimized with Apple Metal - Fixed build failed for 13772 and 13773 with Apple Metal - Fixed build failed for 18400 with Apple Metal diff --git a/src/selftest.c b/src/selftest.c index aea6f3128..d8b105d2e 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -78,6 +78,8 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ tmp.pw_len = (u32) tmp_len; } + cl_event opencl_event; + pw_t pw; pw_t comb; bf_t bf; @@ -118,7 +120,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else @@ -163,7 +165,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) @@ -239,9 +241,9 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, 1 * sizeof (pw_t), &comb, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, 1 * sizeof (pw_t), &comb, 0, NULL, &opencl_event) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else if (user_options_extra->attack_kern == ATTACK_KERN_BF) @@ -284,7 +286,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else @@ -345,7 +347,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs_c, CL_FALSE, 0, 1 * sizeof (bf_t), &bf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs_c, CL_FALSE, 0, 1 * sizeof (bf_t), &bf, 0, NULL, &opencl_event) == -1) return -1; } memset (&pw, 0, sizeof (pw)); @@ -451,7 +453,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } *highest_pw_len = pw.pw_len; @@ -489,11 +491,20 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } } + // prevent race condition on OpenCL with non-blocking write + + if (device_param->is_opencl == true) + { + if (hc_clWaitForEvents (hashcat_ctx, 1, &opencl_event) == -1) return -1; + + if (hc_clReleaseEvent (hashcat_ctx, opencl_event) == -1) return -1; + } + return 0; } From f8df94f4571d557e50ca3a25e5e62111df18dcf4 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Mon, 30 Jun 2025 11:26:05 +0200 Subject: [PATCH 7/7] Switched all async and non-blocking calls to synchronous and blocking ones. Kept the original async bindings intact. This avoids race conditions like the one fixed in the previous commit, with no performance impact. Fixed a typedef issue for clEnqueueReadBuffer(). Updated Python/hcshared.py with missing entry for new salt_dimy attribute in salt_t struct. Fixed a bug in the autotuner when determining the starting value for kernel loops, in cases where the iteration count is N-1 and not a multiple of 1024. Updated additional plugins to use OPTI_TYPE_REGISTER_LIMIT. --- Python/hcshared.py | 3 +- include/ext_OpenCL.h | 2 +- include/ext_cuda.h | 27 +++- include/ext_hip.h | 15 ++ src/autotune.c | 8 +- src/backend.c | 280 ++++++++++++++++++------------------- src/ext_cuda.c | 145 ++++++++++++++++++- src/ext_hip.c | 143 ++++++++++++++++++- src/hashes.c | 14 +- src/modules/module_01300.c | 1 + src/modules/module_07100.c | 11 +- src/modules/module_08200.c | 1 + src/selftest.c | 90 ++++++------ 13 files changed, 523 insertions(+), 217 deletions(-) diff --git a/Python/hcshared.py b/Python/hcshared.py index 2fcf56c70..ad4390083 100644 --- a/Python/hcshared.py +++ b/Python/hcshared.py @@ -6,13 +6,14 @@ import sys def extract_salts(salts_buf) -> list: salts=[] - for salt_buf, salt_buf_pc, salt_len, salt_len_pc, salt_iter, salt_iter2, salt_sign, salt_repeats, orig_pos, digests_cnt, digests_done, digests_offset, scrypt_N, scrypt_r, scrypt_p in struct.iter_unpack("256s 256s I I I I 8s I I I I I I I I", salts_buf): + for salt_buf, salt_buf_pc, salt_len, salt_len_pc, salt_iter, salt_iter2, salt_dimy, salt_sign, salt_repeats, orig_pos, digests_cnt, digests_done, digests_offset, scrypt_N, scrypt_r, scrypt_p in struct.iter_unpack("256s 256s I I I I I 8s I I I I I I I I", salts_buf): salt_buf = salt_buf[0:salt_len] salt_buf_pc = salt_buf_pc[0:salt_len_pc] salts.append({ "salt_buf": salt_buf, \ "salt_buf_pc": salt_buf_pc, \ "salt_iter": salt_iter, \ "salt_iter2": salt_iter2, \ + "salt_dimy": salt_dimy, \ "salt_sign": salt_sign, \ "salt_repeats": salt_repeats, \ "orig_pos": orig_pos, \ diff --git a/include/ext_OpenCL.h b/include/ext_OpenCL.h index d49f82c79..7044ac667 100644 --- a/include/ext_OpenCL.h +++ b/include/ext_OpenCL.h @@ -51,7 +51,7 @@ typedef cl_int (CL_API_CALL *OCL_CLENQUEUEFILLBUFFER) (cl_comman typedef cl_int (CL_API_CALL *OCL_CLENQUEUECOPYBUFFER) (cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t, cl_uint, const cl_event *, cl_event *); typedef void * (CL_API_CALL *OCL_CLENQUEUEMAPBUFFER) (cl_command_queue, cl_mem, cl_bool, cl_map_flags, size_t, size_t, cl_uint, const cl_event *, cl_event *, cl_int *); typedef cl_int (CL_API_CALL *OCL_CLENQUEUENDRANGEKERNEL) (cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *); -typedef cl_int (CL_API_CALL *OCL_CLENQUEUEREADBUFFER) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, cl_event *); +typedef cl_int (CL_API_CALL *OCL_CLENQUEUEREADBUFFER) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *OCL_CLENQUEUEUNMAPMEMOBJECT) (cl_command_queue, cl_mem, void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *OCL_CLENQUEUEWRITEBUFFER) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *OCL_CLFINISH) (cl_command_queue); diff --git a/include/ext_cuda.h b/include/ext_cuda.h index 44ae01cd6..a80af5367 100644 --- a/include/ext_cuda.h +++ b/include/ext_cuda.h @@ -1154,14 +1154,19 @@ typedef CUresult (CUDA_API_CALL *CUDA_CUINIT) (unsigned int); typedef CUresult (CUDA_API_CALL *CUDA_CULAUNCHKERNEL) (CUfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, CUstream, void **, void **); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOC) (CUdeviceptr *, size_t); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOCHOST) (void **, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOD) (CUdeviceptr, CUdeviceptr, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOH) (void *, CUdeviceptr, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTOD) (CUdeviceptr, const void *, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32) (CUdeviceptr, unsigned int, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8) (CUdeviceptr, unsigned char, size_t); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTODASYNC) (CUdeviceptr, CUdeviceptr, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOHASYNC) (void *, CUdeviceptr, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTODASYNC) (CUdeviceptr, const void *, size_t, CUstream); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32ASYNC) (CUdeviceptr, unsigned int, size_t, CUstream); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8ASYNC) (CUdeviceptr, unsigned char, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREE) (CUdeviceptr); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREEHOST) (void *); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMGETINFO) (size_t *, size_t *); -typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32ASYNC) (CUdeviceptr, unsigned int, size_t, CUstream); -typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8ASYNC) (CUdeviceptr, unsigned char, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETFUNCTION) (CUfunction *, CUmodule, const char *); typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETGLOBAL) (CUdeviceptr *, size_t *, CUmodule, const char *); typedef CUresult (CUDA_API_CALL *CUDA_CUMODULELOAD) (CUmodule *, const char *); @@ -1217,14 +1222,19 @@ typedef struct hc_cuda_lib CUDA_CULAUNCHKERNEL cuLaunchKernel; CUDA_CUMEMALLOC cuMemAlloc; CUDA_CUMEMALLOCHOST cuMemAllocHost; + CUDA_CUMEMCPYDTOD cuMemcpyDtoD; + CUDA_CUMEMCPYDTOH cuMemcpyDtoH; + CUDA_CUMEMCPYHTOD cuMemcpyHtoD; + CUDA_CUMEMSETD32 cuMemsetD32; + CUDA_CUMEMSETD8 cuMemsetD8; CUDA_CUMEMCPYDTODASYNC cuMemcpyDtoDAsync; CUDA_CUMEMCPYDTOHASYNC cuMemcpyDtoHAsync; CUDA_CUMEMCPYHTODASYNC cuMemcpyHtoDAsync; + CUDA_CUMEMSETD32ASYNC cuMemsetD32Async; + CUDA_CUMEMSETD8ASYNC cuMemsetD8Async; CUDA_CUMEMFREE cuMemFree; CUDA_CUMEMFREEHOST cuMemFreeHost; CUDA_CUMEMGETINFO cuMemGetInfo; - CUDA_CUMEMSETD32ASYNC cuMemsetD32Async; - CUDA_CUMEMSETD8ASYNC cuMemsetD8Async; CUDA_CUMODULEGETFUNCTION cuModuleGetFunction; CUDA_CUMODULEGETGLOBAL cuModuleGetGlobal; CUDA_CUMODULELOAD cuModuleLoad; @@ -1272,13 +1282,18 @@ int hc_cuFuncSetAttribute (void *hashcat_ctx, CUfunction hfunc, CUfunction_ int hc_cuInit (void *hashcat_ctx, unsigned int Flags); int hc_cuLaunchKernel (void *hashcat_ctx, CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra); int hc_cuMemAlloc (void *hashcat_ctx, CUdeviceptr *dptr, size_t bytesize); +int hc_cuMemcpyDtoD (void *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount); +int hc_cuMemcpyDtoH (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount); +int hc_cuMemcpyHtoD (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount); +int hc_cuMemsetD32 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N); +int hc_cuMemsetD8 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N); int hc_cuMemcpyDtoDAsync (void *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream); int hc_cuMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream); int hc_cuMemcpyHtoDAsync (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream); -int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr); -int hc_cuMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); 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_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); diff --git a/include/ext_hip.h b/include/ext_hip.h index 4cc6fb797..d0f53d173 100644 --- a/include/ext_hip.h +++ b/include/ext_hip.h @@ -588,6 +588,11 @@ typedef hipError_t (HIP_API_CALL *HIP_HIPLAUNCHKERNEL) (hipFunction_t, typedef hipError_t (HIP_API_CALL *HIP_HIPMEMALLOC) (hipDeviceptr_t *, size_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMFREE) (hipDeviceptr_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMGETINFO) (size_t *, size_t *); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOD) (hipDeviceptr_t, hipDeviceptr_t, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOH) (void *, hipDeviceptr_t, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTOD) (hipDeviceptr_t, const void *, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD32) (hipDeviceptr_t, unsigned int, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD8) (hipDeviceptr_t, unsigned char, size_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTODASYNC) (hipDeviceptr_t, hipDeviceptr_t, size_t, hipStream_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOHASYNC) (void *, hipDeviceptr_t, size_t, hipStream_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTODASYNC) (hipDeviceptr_t, const void *, size_t, hipStream_t); @@ -633,6 +638,11 @@ typedef struct hc_hip_lib HIP_HIPMEMALLOC hipMemAlloc; HIP_HIPMEMFREE hipMemFree; HIP_HIPMEMGETINFO hipMemGetInfo; + HIP_HIPMEMCPYDTOD hipMemcpyDtoD; + HIP_HIPMEMCPYDTOH hipMemcpyDtoH; + HIP_HIPMEMCPYHTOD hipMemcpyHtoD; + HIP_HIPMEMSETD32 hipMemsetD32; + HIP_HIPMEMSETD8 hipMemsetD8; HIP_HIPMEMCPYDTODASYNC hipMemcpyDtoDAsync; HIP_HIPMEMCPYDTOHASYNC hipMemcpyDtoHAsync; HIP_HIPMEMCPYHTODASYNC hipMemcpyHtoDAsync; @@ -680,6 +690,11 @@ int hc_hipLaunchKernel (void *hashcat_ctx, hipFunction_t f, unsigned int int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize); 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); +int hc_hipMemcpyHtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount); +int hc_hipMemsetD32 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N); +int hc_hipMemsetD8 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N); int hc_hipMemcpyDtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); int hc_hipMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); int hc_hipMemcpyHtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream); diff --git a/src/autotune.c b/src/autotune.c index 065c0a217..a599e65be 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -268,12 +268,12 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t)) == -1) return -1; } #if defined (__APPLE__) @@ -344,8 +344,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param start = MIN (start, smallest_repeat_double (hashes->st_salts_buf->salt_iter)); start = MIN (start, smallest_repeat_double (hashes->st_salts_buf->salt_iter + 1)); - if ((hashes->st_salts_buf->salt_iter % 125) == 0) start = MIN (start, 125); - if ((hashes->st_salts_buf->salt_iter + 1 % 125) == 0) start = MIN (start, 125); + if (((hashes->st_salts_buf->salt_iter + 0) % 125) == 0) start = MIN (start, 125); + if (((hashes->st_salts_buf->salt_iter + 1) % 125) == 0) start = MIN (start, 125); if ((start >= kernel_loops_min) && (start <= kernel_loops_max)) { diff --git a/src/backend.c b/src/backend.c index 00cac2245..2d34caee1 100644 --- a/src/backend.c +++ b/src/backend.c @@ -985,7 +985,7 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c { if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1) return -1; - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &pw_idx, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, &pw_idx, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } @@ -994,7 +994,7 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c { if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &pw_idx, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t), device_param->hip_stream) == -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; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1020,14 +1020,14 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, pw->i, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, pw->i, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, pw->i, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, pw->i, 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; } @@ -1072,7 +1072,7 @@ int copy_pws_idx (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u { if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1) return -1; - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, dest, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t)), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, dest, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t))) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; @@ -1083,7 +1083,7 @@ int copy_pws_idx (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u { if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, dest, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t)), device_param->hip_stream) == -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; @@ -1112,7 +1112,7 @@ int copy_pws_comp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1) return -1; - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, dest, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, dest, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; @@ -1123,7 +1123,7 @@ int copy_pws_comp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, dest, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->hip_stream) == -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; @@ -1199,12 +1199,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_tm_c, size_tm, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_tm_c, size_tm) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_tm_c, size_tm, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_tm_c, size_tm) == -1) return -1; } #if defined (__APPLE__) @@ -1309,12 +1309,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, device_param->cuda_d_pws_amp_buf, pws_cnt * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_pws_buf, device_param->cuda_d_pws_amp_buf, pws_cnt * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, device_param->hip_d_pws_amp_buf, pws_cnt * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_pws_buf, device_param->hip_d_pws_amp_buf, pws_cnt * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -1373,14 +1373,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1432,12 +1432,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } #if defined (__APPLE__) @@ -1534,14 +1534,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1563,14 +1563,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1625,14 +1625,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1684,12 +1684,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } #if defined (__APPLE__) @@ -1774,14 +1774,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1803,14 +1803,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -2059,12 +2059,12 @@ int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t return 0; } -int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u8 value, const u64 size) +int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u8 value, const u64 size) { - return hc_cuMemsetD8Async (hashcat_ctx, buf + offset, value, size, device_param->cuda_stream); + return hc_cuMemsetD8 (hashcat_ctx, buf + offset, value, size); } -int run_cuda_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u32 value, const u64 size) +int run_cuda_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u32 value, const u64 size) { /* check that the size is multiple of element size */ if (size % 4 != 0) @@ -2072,7 +2072,7 @@ int run_cuda_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *dev return CUDA_ERROR_INVALID_VALUE; } - return hc_cuMemsetD32Async (hashcat_ctx, buf + offset, value, size / 4, device_param->cuda_stream); + return hc_cuMemsetD32 (hashcat_ctx, buf + offset, value, size / 4); } int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 size) @@ -2096,7 +2096,7 @@ int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device if (num16m) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, buf + (num16d * 16), bzeros, num16m, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, buf + (num16d * 16), bzeros, num16m) == -1) return -1; } return 0; @@ -2138,12 +2138,12 @@ int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t return 0; } -int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u8 value, const u64 size) +int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u8 value, const u64 size) { - return hc_hipMemsetD8Async (hashcat_ctx, buf + offset, value, size, device_param->hip_stream); + return hc_hipMemsetD8 (hashcat_ctx, buf + offset, value, size); } -int run_hip_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u32 value, const u64 size) +int run_hip_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u32 value, const u64 size) { /* check that the size is multiple of element size */ if (size % 4 != 0) @@ -2151,7 +2151,7 @@ int run_hip_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *devi return hipErrorInvalidValue; } - return hc_hipMemsetD32Async (hashcat_ctx, buf + offset, value, size / 4, device_param->hip_stream); + return hc_hipMemsetD32 (hashcat_ctx, buf + offset, value, size / 4); } int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 size) @@ -2175,7 +2175,7 @@ int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (num16m) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, buf + (num16d * 16), bzeros, num16m, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, buf + (num16d * 16), bzeros, num16m) == -1) return -1; } return 0; @@ -2466,7 +2466,7 @@ int run_opencl_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *devi } else { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, buf, CL_FALSE, num16d * 16, num16m, bzeros, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, buf, CL_TRUE, num16d * 16, num16m, bzeros, 0, NULL, NULL) == -1) return -1; } } @@ -2592,7 +2592,7 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con case KERN_RUN_AUX4: cuda_function = device_param->cuda_function_aux4; break; } - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params) == -1) return -1; if (hc_cuFuncSetAttribute (hashcat_ctx, cuda_function, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shared_mem) == -1) return -1; @@ -2711,7 +2711,7 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con case KERN_RUN_AUX4: hip_function = device_param->hip_function_aux4; break; } - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params) == -1) return -1; //if (hc_hipFuncSetAttribute (hashcat_ctx, hip_function, HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shared_mem) == -1) return -1; @@ -2982,7 +2982,7 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con if (hc_clSetKernelArg (hashcat_ctx, opencl_kernel, i, sizeof (cl_mem), device_param->kernel_params[i]) == -1) return -1; } - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_kernel_param, CL_FALSE, 0, device_param->size_kernel_params, &device_param->kernel_param, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_kernel_param, CL_TRUE, 0, device_param->size_kernel_params, &device_param->kernel_param, 0, NULL, NULL) == -1) return -1; /* for (u32 i = 24; i <= 34; i++) @@ -3616,7 +3616,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3624,13 +3624,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3638,7 +3638,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3660,7 +3660,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3668,7 +3668,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3680,7 +3680,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3688,13 +3688,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3702,7 +3702,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3724,7 +3724,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3732,7 +3732,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3778,7 +3778,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3786,13 +3786,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3800,7 +3800,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3822,7 +3822,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3830,7 +3830,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3842,7 +3842,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3850,13 +3850,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3864,7 +3864,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3886,7 +3886,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3894,7 +3894,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3904,7 +3904,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3912,13 +3912,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3926,7 +3926,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3948,7 +3948,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3956,7 +3956,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -4190,12 +4190,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4322,12 +4322,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4339,7 +4339,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; } } else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) @@ -4352,12 +4352,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4382,12 +4382,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4515,12 +4515,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4532,7 +4532,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; } } else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) @@ -4545,12 +4545,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4577,12 +4577,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_bfs, innerloop_left * sizeof (bf_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_bfs, innerloop_left * sizeof (bf_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_bfs, innerloop_left * sizeof (bf_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_bfs, innerloop_left * sizeof (bf_t)) == -1) return -1; } #if defined (__APPLE__) @@ -11863,16 +11863,16 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_st_salts_buf, size_st_salts) == -1) return -1; if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_kernel_param, size_kernel_params) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_digests_buf, hashes->digests_buf, size_digests, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_salt_bufs, hashes->salts_buf, size_salts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_digests_buf, hashes->digests_buf, size_digests) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_salt_bufs, hashes->salts_buf, size_salts) == -1) return -1; /** * special buffers @@ -11899,7 +11899,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_rules_c, size_rules_c) == -1) return -1; } - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_rules, straight_ctx->kernel_rules_buf, size_rules_src, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_rules, straight_ctx->kernel_rules_buf, size_rules_src) == -1) return -1; } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { @@ -11934,19 +11934,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_esalt_bufs, size_esalts) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_esalt_bufs, hashes->esalts_buf, size_esalts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_esalt_bufs, hashes->esalts_buf, size_esalts) == -1) return -1; } if (hashconfig->st_hash != NULL) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_st_digests_buf, hashes->st_digests_buf, size_st_digests, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_st_salts_buf, hashes->st_salts_buf, size_st_salts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_st_digests_buf, hashes->st_digests_buf, size_st_digests) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_st_salts_buf, hashes->st_salts_buf, size_st_salts) == -1) return -1; if (size_esalts) { if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_st_esalts_buf, size_st_esalts) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts) == -1) return -1; } } } @@ -11974,16 +11974,16 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_st_salts_buf, size_st_salts) == -1) return -1; if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_kernel_param, size_kernel_params) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_digests_buf, hashes->digests_buf, size_digests, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_salt_bufs, hashes->salts_buf, size_salts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_digests_buf, hashes->digests_buf, size_digests) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_salt_bufs, hashes->salts_buf, size_salts) == -1) return -1; /** * special buffers @@ -12010,7 +12010,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_rules_c, size_rules_c) == -1) return -1; } - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_rules, straight_ctx->kernel_rules_buf, size_rules_src, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_rules, straight_ctx->kernel_rules_buf, size_rules_src) == -1) return -1; } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { @@ -12045,19 +12045,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_esalt_bufs, size_esalts) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_esalt_bufs, hashes->esalts_buf, size_esalts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_esalt_bufs, hashes->esalts_buf, size_esalts) == -1) return -1; } if (hashconfig->st_hash != NULL) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_st_digests_buf, hashes->st_digests_buf, size_st_digests, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_st_salts_buf, hashes->st_salts_buf, size_st_salts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_digests_buf, hashes->st_digests_buf, size_st_digests) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_salts_buf, hashes->st_salts_buf, size_st_salts) == -1) return -1; if (size_esalts) { if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_st_esalts_buf, size_st_esalts) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts) == -1) return -1; } } } @@ -12195,15 +12195,15 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_st_salts, NULL, &device_param->opencl_d_st_salts_buf) == -1) return -1; if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_kernel_params, NULL, &device_param->opencl_d_kernel_param) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_a, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_a, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_b, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_b, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_c, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_c, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_d, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_d, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_a, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_a, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_b, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_b, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_c, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_c, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_d, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_d, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_buf, CL_FALSE, 0, size_digests, hashes->digests_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_a, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_a, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_b, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_b, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_c, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_c, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_d, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_d, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_a, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_a, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_b, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_b, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_c, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_c, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_d, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_d, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_buf, CL_TRUE, 0, size_digests, hashes->digests_buf, 0, NULL, NULL) == -1) return -1; if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_salt_bufs, CL_FALSE, 0, size_salts, hashes->salts_buf, 0, NULL, NULL) == -1) return -1; /** @@ -12221,7 +12221,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_rules, NULL, &device_param->opencl_d_rules) == -1) return -1; if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_rules_c, NULL, &device_param->opencl_d_rules_c) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, CL_FALSE, 0, size_rules_src, straight_ctx->kernel_rules_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, CL_TRUE, 0, size_rules_src, straight_ctx->kernel_rules_buf, 0, NULL, NULL) == -1) return -1; } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { @@ -12244,7 +12244,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_esalts, NULL, &device_param->opencl_d_esalt_bufs) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_esalt_bufs, CL_FALSE, 0, size_esalts, hashes->esalts_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_esalt_bufs, CL_TRUE, 0, size_esalts, hashes->esalts_buf, 0, NULL, NULL) == -1) return -1; } if (hashconfig->st_hash != NULL) @@ -12256,7 +12256,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_st_esalts, NULL, &device_param->opencl_d_st_esalts_buf) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_st_esalts_buf, CL_FALSE, 0, size_st_esalts, hashes->st_esalts_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_st_esalts_buf, CL_TRUE, 0, size_st_esalts, hashes->st_esalts_buf, 0, NULL, NULL) == -1) return -1; } } @@ -17537,14 +17537,14 @@ int backend_session_update_mp (hashcat_ctx_t *hashcat_ctx) if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } #if defined (__APPLE__) @@ -17558,7 +17558,7 @@ int backend_session_update_mp (hashcat_ctx_t *hashcat_ctx) if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_root_css_buf, CL_FALSE, 0, device_param->size_root_css, mask_ctx->root_css_buf, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_FALSE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; } @@ -17593,14 +17593,14 @@ int backend_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } #if defined (__APPLE__) @@ -17614,7 +17614,7 @@ int backend_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_ if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_root_css_buf, CL_FALSE, 0, device_param->size_root_css, mask_ctx->root_css_buf, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_FALSE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; } diff --git a/src/ext_cuda.c b/src/ext_cuda.c index 4c0db9ffb..878977ade 100644 --- a/src/ext_cuda.c +++ b/src/ext_cuda.c @@ -87,14 +87,19 @@ int cuda_init (void *hashcat_ctx) HC_LOAD_FUNC_CUDA (cuda, cuLaunchKernel, cuLaunchKernel, CUDA_CULAUNCHKERNEL, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemAlloc, cuMemAlloc_v2, CUDA_CUMEMALLOC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemAllocHost, cuMemAllocHost_v2, CUDA_CUMEMALLOCHOST, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoD, cuMemcpyDtoD_v2, CUDA_CUMEMCPYDTOD, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoH, cuMemcpyDtoH_v2, CUDA_CUMEMCPYDTOH, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemcpyHtoD, cuMemcpyHtoD_v2, CUDA_CUMEMCPYHTOD, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD32, cuMemsetD32, CUDA_CUMEMSETD32, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD8, cuMemsetD8, CUDA_CUMEMSETD8, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoDAsync, cuMemcpyDtoDAsync_v2, CUDA_CUMEMCPYDTODASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoHAsync, cuMemcpyDtoHAsync_v2, CUDA_CUMEMCPYDTOHASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemcpyHtoDAsync, cuMemcpyHtoDAsync_v2, CUDA_CUMEMCPYHTODASYNC, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD32Async, cuMemsetD32Async, CUDA_CUMEMSETD32ASYNC, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD8Async, cuMemsetD8Async, CUDA_CUMEMSETD8ASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemFree, cuMemFree_v2, CUDA_CUMEMFREE, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemFreeHost, cuMemFreeHost, CUDA_CUMEMFREEHOST, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemGetInfo, cuMemGetInfo_v2, CUDA_CUMEMGETINFO, CUDA, 1); - HC_LOAD_FUNC_CUDA (cuda, cuMemsetD32Async, cuMemsetD32Async, CUDA_CUMEMSETD32ASYNC, CUDA, 1); - HC_LOAD_FUNC_CUDA (cuda, cuMemsetD8Async, cuMemsetD8Async, CUDA_CUMEMSETD8ASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuModuleGetFunction, cuModuleGetFunction, CUDA_CUMODULEGETFUNCTION, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuModuleGetGlobal, cuModuleGetGlobal_v2, CUDA_CUMODULEGETGLOBAL, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuModuleLoad, cuModuleLoad, CUDA_CUMODULELOAD, CUDA, 1); @@ -517,6 +522,142 @@ int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr) return 0; } + +int hc_cuMemcpyDtoH (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount) +{ + 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->cuMemcpyDtoH (dstHost, srcDevice, ByteCount); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemcpyDtoH(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemcpyDtoH(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemcpyDtoD (void *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount) +{ + 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->cuMemcpyDtoD (dstDevice, srcDevice, ByteCount); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemcpyDtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemcpyDtoD(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemcpyHtoD (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount) +{ + 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->cuMemcpyHtoD (dstDevice, srcHost, ByteCount); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemcpyHtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemcpyHtoD(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemsetD32 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N) +{ + 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->cuMemsetD32 (dstDevice, ui, N); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemsetD32(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemsetD32(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemsetD8 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N) +{ + 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->cuMemsetD8 (dstDevice, uc, N); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemsetD8(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemsetD8(): %d", CU_err); + } + + return -1; + } + + return 0; +} + int hc_cuMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; diff --git a/src/ext_hip.c b/src/ext_hip.c index decfa987c..4a26a7e9f 100644 --- a/src/ext_hip.c +++ b/src/ext_hip.c @@ -140,12 +140,16 @@ int hip_init (void *hashcat_ctx) 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, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, 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); @@ -800,6 +804,143 @@ int hc_hipMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total) return 0; } + + +int hc_hipMemcpyDtoH (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount) +{ + 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->hipMemcpyDtoH (dstHost, srcDevice, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyDtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount) +{ + 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->hipMemcpyDtoD (dstDevice, srcDevice, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyHtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount) +{ + 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->hipMemcpyHtoD (dstDevice, srcHost, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemsetD32 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N) +{ + 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->hipMemsetD32 (dstDevice, ui, N); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemsetD32(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemsetD32(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemsetD8 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N) +{ + 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->hipMemsetD8 (dstDevice, uc, N); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemsetD8(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemsetD8(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + int hc_hipMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; diff --git a/src/hashes.c b/src/hashes.c index 78a6d5b90..72ab14433 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -334,7 +334,7 @@ int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pla if (device_param->is_cuda == true) { - rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->cuda_stream); + rc = hc_cuMemcpyDtoH (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size); if (rc == 0) { @@ -351,7 +351,7 @@ int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pla if (device_param->is_hip == true) { - rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->hip_stream); + rc = hc_hipMemcpyDtoH (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size); if (rc == 0) { @@ -382,7 +382,7 @@ int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pla if (device_param->is_opencl == true) { - rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_FALSE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, &opencl_event); + rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_TRUE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, &opencl_event); if (rc == 0) { @@ -587,14 +587,14 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32)) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -624,7 +624,7 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (device_param->is_cuda == true) { - rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->cuda_stream); + rc = hc_cuMemcpyDtoH (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t)); if (rc == 0) { @@ -641,7 +641,7 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (device_param->is_hip == true) { - rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->hip_stream); + rc = hc_hipMemcpyDtoH (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t)); if (rc == 0) { diff --git a/src/modules/module_01300.c b/src/modules/module_01300.c index b50a0e99c..c79a3493c 100644 --- a/src/modules/module_01300.c +++ b/src/modules/module_01300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH; static const char *HASH_NAME = "SHA2-224"; static const u64 KERN_TYPE = 1300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_07100.c b/src/modules/module_07100.c index ec3924cec..78c8f8081 100644 --- a/src/modules/module_07100.c +++ b/src/modules/module_07100.c @@ -85,15 +85,6 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con return pw_max; } -char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param) -{ - char *jit_build_options = NULL; - - hc_asprintf (&jit_build_options, "-D NO_UNROLL"); - - return jit_build_options; -} - int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len) { u64 *digest = (u64 *) digest_buf; @@ -398,7 +389,7 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_hook23 = MODULE_DEFAULT; module_ctx->module_hook_salt_size = MODULE_DEFAULT; module_ctx->module_hook_size = MODULE_DEFAULT; - module_ctx->module_jit_build_options = module_jit_build_options; + module_ctx->module_jit_build_options = MODULE_DEFAULT; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_08200.c b/src/modules/module_08200.c index 54d5c62b5..90049b1ed 100644 --- a/src/modules/module_08200.c +++ b/src/modules/module_08200.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "1Password, cloudkeychain"; static const u64 KERN_TYPE = 8200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/selftest.c b/src/selftest.c index d8b105d2e..5e744f88b 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -103,12 +103,12 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -120,7 +120,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else @@ -148,12 +148,12 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -165,7 +165,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) @@ -218,16 +218,16 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, &comb, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, &comb, 1 * sizeof (pw_t)) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, &comb, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, &comb, 1 * sizeof (pw_t)) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -241,9 +241,9 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, 1 * sizeof (pw_t), &comb, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, 1 * sizeof (pw_t), &comb, 0, NULL, &opencl_event) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else if (user_options_extra->attack_kern == ATTACK_KERN_BF) @@ -269,12 +269,12 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -286,7 +286,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else @@ -330,12 +330,12 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bfs_c, &bf, 1 * sizeof (bf_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bfs_c, &bf, 1 * sizeof (bf_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bfs_c, &bf, 1 * sizeof (bf_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bfs_c, &bf, 1 * sizeof (bf_t)) == -1) return -1; } #if defined (__APPLE__) @@ -347,7 +347,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs_c, CL_FALSE, 0, 1 * sizeof (bf_t), &bf, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs_c, CL_TRUE, 0, 1 * sizeof (bf_t), &bf, 0, NULL, &opencl_event) == -1) return -1; } memset (&pw, 0, sizeof (pw)); @@ -436,12 +436,12 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -453,7 +453,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } *highest_pw_len = pw.pw_len; @@ -474,12 +474,12 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -491,7 +491,7 @@ static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } } @@ -587,14 +587,14 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -616,12 +616,12 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } #if defined (__APPLE__) @@ -633,7 +633,7 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_FALSE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; } } @@ -683,14 +683,14 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -724,14 +724,14 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -758,14 +758,14 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -787,12 +787,12 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } #if defined (__APPLE__) @@ -804,7 +804,7 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_FALSE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; } } } @@ -846,14 +846,14 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -887,12 +887,12 @@ static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *d { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; } #if defined (__APPLE__) @@ -962,14 +962,14 @@ static int selftest_cleanup (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *devi if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, num_cracked, device_param->cuda_d_result, sizeof (u32)) == -1) return -1; if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event3, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, num_cracked, device_param->hip_d_result, sizeof (u32)) == -1) return -1; if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event3, device_param->hip_stream) == -1) return -1; } @@ -983,7 +983,7 @@ static int selftest_cleanup (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *devi if (device_param->is_opencl == true) { - if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_FALSE, 0, sizeof (u32), num_cracked, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), num_cracked, 0, NULL, &opencl_event) == -1) return -1; if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; }