arno
/
hashcat
mirror of https://github.com/hashcat/hashcat.git
1
0
Fork 0
Code Issues Releases Wiki Activity

Add --gpu-architecture to NVRTC build options

Browse Source
pull/2022/head
Jens Steube 5 years ago
parent 6a32e8ef18
commit 06171958ee
2 changed files with 189 additions and 136 deletions
Show Stats Download Patch File Download Diff File

6
include/types.h
Unescape View File

@ -995,6 +995,8 @@ typedef struct link_speed
typedef struct hc_device_param
{
CUdevice device_cuda;
cl_device_id device;
cl_device_type device_type;
@ -1006,8 +1008,8 @@ typedef struct hc_device_param
st_status_t st_status;
u32 sm_major;
u32 sm_minor;
int sm_major;
int sm_minor;
u32 kernel_exec_timeout;
u8 pcie_bus;

319
src/backend.c
Unescape View File

@ -543,6 +543,8 @@ void generate_cached_kernel_amp_filename (const u32 attack_kern, char *profile_d
snprintf (cached_file, 255, "%s/kernels/amp_a%u.%s.kernel", profile_dir, attack_kern, device_name_chksum_amp_mp);
}
// NVRTC
int nvrtc_init (hashcat_ctx_t *hashcat_ctx)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
@ -611,6 +613,134 @@ void nvrtc_close (hashcat_ctx_t *hashcat_ctx)
}
}
int hc_nvrtcCreateProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcCreateProgram (prog, src, name, numHeaders, headers, includeNames);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcCreateProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcDestroyProgram (prog);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcDestroyProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcCompileProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, int numOptions, const char * const *options)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcCompileProgram (prog, numOptions, options);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcCompileProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *logSizeRet)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetProgramLogSize (prog, logSizeRet);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetProgramLogSize(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *log)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetProgramLog (prog, log);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetProgramLog(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetPTXSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *ptxSizeRet)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetPTXSize (prog, ptxSizeRet);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetPTXSize(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetPTX (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetPTX (prog, ptx);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetPTX(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
// CUDA
int cuda_init (hashcat_ctx_t *hashcat_ctx)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
@ -721,6 +851,35 @@ void cuda_close (hashcat_ctx_t *hashcat_ctx)
}
}
int hc_cuDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, CUdevice_attribute attrib, CUdevice dev)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
CUDA_PTR *cuda = backend_ctx->cuda;
const CUresult CU_err = cuda->cuDeviceGetAttribute (pi, attrib, dev);
if (CU_err != CUDA_SUCCESS)
{
const char *pStr = NULL;
if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS)
{
event_log_error (hashcat_ctx, "cuDeviceGetAttribute(): %s", pStr);
}
else
{
event_log_error (hashcat_ctx, "cuDeviceGetAttribute(): %d", CU_err);
}
return -1;
}
return 0;
}
// OpenCL
int ocl_init (hashcat_ctx_t *hashcat_ctx)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
@ -832,132 +991,6 @@ void ocl_close (hashcat_ctx_t *hashcat_ctx)
}
}
int hc_nvrtcCreateProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcCreateProgram (prog, src, name, numHeaders, headers, includeNames);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcCreateProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcDestroyProgram (prog);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcDestroyProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcCompileProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, int numOptions, const char * const *options)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcCompileProgram (prog, numOptions, options);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcCompileProgram(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *logSizeRet)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetProgramLogSize (prog, logSizeRet);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetProgramLogSize(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *log)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetProgramLog (prog, log);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetProgramLog(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetPTXSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *ptxSizeRet)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetPTXSize (prog, ptxSizeRet);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetPTXSize(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_nvrtcGetPTX (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
NVRTC_PTR *nvrtc = backend_ctx->nvrtc;
const nvrtcResult NVRTC_err = nvrtc->nvrtcGetPTX (prog, ptx);
if (NVRTC_err != NVRTC_SUCCESS)
{
event_log_error (hashcat_ctx, "nvrtcGetPTX(): %s", nvrtc->nvrtcGetErrorString (NVRTC_err));
return -1;
}
return 0;
}
int hc_clEnqueueNDRangeKernel (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
{
backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx;
@ -3571,6 +3604,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
// platform vendor
int CL_rc;
int CU_rc;
CL_rc = hc_clGetPlatformInfo (hashcat_ctx, platform, CL_PLATFORM_VENDOR, 0, NULL, &param_value_size);
@ -4150,16 +4184,30 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
device_param->pcie_device = (u8) (pci_slot_id_nv >> 3);
device_param->pcie_function = (u8) (pci_slot_id_nv & 7);
cl_uint sm_minor = 0;
cl_uint sm_major = 0;
int sm_minor = 0;
int sm_major = 0;
CL_rc = hc_clGetDeviceInfo (hashcat_ctx, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL);
//if (backend_ctx->cuda)
if (0)
{
CU_rc = hc_cuDeviceGetAttribute (hashcat_ctx, &sm_minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device_param->device_cuda);
if (CL_rc == -1) return -1;
if (CU_rc == -1) return -1;
CL_rc = hc_clGetDeviceInfo (hashcat_ctx, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL);
CU_rc = hc_cuDeviceGetAttribute (hashcat_ctx, &sm_major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device_param->device_cuda);
if (CL_rc == -1) return -1;
if (CU_rc == -1) return -1;
}
else
{
CL_rc = hc_clGetDeviceInfo (hashcat_ctx, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL);
if (CL_rc == -1) return -1;
CL_rc = hc_clGetDeviceInfo (hashcat_ctx, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL);
if (CL_rc == -1) return -1;
}
device_param->sm_minor = sm_minor;
device_param->sm_major = sm_major;
@ -5490,13 +5538,16 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
if (rc_nvrtcCreateProgram == -1) return -1;
char **nvrtc_options = (char **) hccalloc (1 + strlen (build_options_module_buf) + 1, sizeof (char *)); // ...
char **nvrtc_options = (char **) hccalloc (3 + strlen (build_options_module_buf) + 1, sizeof (char *)); // ...
nvrtc_options[0] = "--device-as-default-execution-space";
nvrtc_options[1] = "--gpu-architecture";
hc_asprintf (&nvrtc_options[2], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
char *nvrtc_options_string = hcstrdup (build_options_module_buf);
const int num_options = 1 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 1);
const int num_options = 3 + nvrtc_make_options_array_from_string (nvrtc_options_string, nvrtc_options + 3);
const int rc_nvrtcCompileProgram = hc_nvrtcCompileProgram (hashcat_ctx, program, num_options, (const char * const *) nvrtc_options);

Write Preview
Loading…
Cancel
Save