arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-02-02 02:41:35 +00:00
Code Issues Releases Wiki Activity

OpenCL Runtime: Improve ROCM detection and make sure to not confuse with recent AMDGPU drivers

Browse Source
This commit is contained in:
jsteube 2019-02-23 22:50:19 +01:00
parent 9b916918de
commit 761475b224
2 changed files with 33 additions and 23 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
docs/changes.txt
View File

@ -27,6 +27,7 @@
- Tuning Database: Updated hashcat.hctune with new models and refreshed vector width values
- Bitcoin Wallet: Be more user friendly by allowing a larger data range for ckey and public_key
- OpenCL Runtime: Improve ROCM detection and make sure to not confuse with recent AMDGPU drivers
##
## Technical

55
src/opencl.c
View File

@ -3701,29 +3701,6 @@ int opencl_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
device_param->pcie_bus = amdtopo.pcie.bus;
device_param->pcie_device = amdtopo.pcie.device;
device_param->pcie_function = amdtopo.pcie.function;
#if defined (__linux__)
// check for AMD ROCm driver (only available on linux)
const char *t1 = strstr (device_param->driver_version, "HSA");
const char *t2 = strstr (device_param->driver_version, "LC");
const char *t3 = strstr (device_param->driver_version, "PAL");
if ((t1 == NULL) && (t2 == NULL) && (t3 == NULL))
{
device_param->is_rocm = false;
}
else
{
device_param->is_rocm = true;
}
#else
device_param->is_rocm = false;
#endif
}
if ((device_param->platform_vendor_id == VENDOR_ID_NV) && (device_param->device_vendor_id == VENDOR_ID_NV))
@ -4613,6 +4590,38 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
if (CL_rc == -1) return -1;
if ((device_param->device_type & CL_DEVICE_TYPE_GPU) && (device_param->platform_vendor_id == VENDOR_ID_AMD))
{
char *kernel_buf = "__kernel void test (__global int *p) { __asm__ (\"DS_NOP\"); }";
const size_t kernel_len = strlen (kernel_buf);
cl_program program;
CL_rc = hc_clCreateProgramWithSource (hashcat_ctx, device_param->context, 1, &kernel_buf, &kernel_len, &program);
if (CL_rc == -1) return -1;
opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx;
OCL_PTR *ocl = opencl_ctx->ocl;
CL_rc = ocl->clBuildProgram (program, 1, &device_param->device, NULL, NULL, NULL);
if (CL_rc == CL_SUCCESS)
{
device_param->is_rocm = true;
}
else
{
device_param->is_rocm = false;
}
CL_rc = hc_clReleaseProgram (hashcat_ctx, program);
if (CL_rc == -1) return -1;
}
// device_available_mem
#define MAX_ALLOC_CHECKS_CNT 8192