arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-23 00:28:11 +00:00
Code Issues Releases Wiki Activity

Disable inline assembly instruction tests for CUDA and refer to documented requirements

Browse Source
This commit is contained in:
Jens Steube 2020-01-03 12:27:27 +01:00
parent d0fb171da9
commit df5e2361d3
1 changed files with 28 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

28
src/backend.c
View File

@ -280,6 +280,7 @@ static bool setup_opencl_device_types_filter (hashcat_ctx_t *hashcat_ctx, const
return true; return true;
} }
/*
static bool cuda_test_instruction (hashcat_ctx_t *hashcat_ctx, const int sm_major, const int sm_minor, const char *kernel_buf) static bool cuda_test_instruction (hashcat_ctx_t *hashcat_ctx, const int sm_major, const int sm_minor, const char *kernel_buf)
{ {
nvrtcProgram program; nvrtcProgram program;
@ -356,6 +357,7 @@ static bool cuda_test_instruction (hashcat_ctx_t *hashcat_ctx, const int sm_majo
return true; return true;
} }
*/
static bool opencl_test_instruction (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_device_id device, const char *kernel_buf) static bool opencl_test_instruction (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_device_id device, const char *kernel_buf)
{ {
@ -5444,6 +5446,18 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
// //
//if (rc_cuCtxSetCacheConfig == -1) return -1; //if (rc_cuCtxSetCacheConfig == -1) return -1;
const int sm = (device_param->sm_major * 10) + device_param->sm_minor;
device_param->has_add = (sm >= 12) ? true : false;
device_param->has_addc = (sm >= 12) ? true : false;
device_param->has_sub = (sm >= 12) ? true : false;
device_param->has_subc = (sm >= 12) ? true : false;
device_param->has_bfe = (sm >= 20) ? true : false;
device_param->has_lop3 = (sm >= 50) ? true : false;
device_param->has_mov64 = (sm >= 10) ? true : false;
device_param->has_prmt = (sm >= 20) ? true : false;
/*
#define RUN_INSTRUCTION_CHECKS() \ #define RUN_INSTRUCTION_CHECKS() \
device_param->has_add = cuda_test_instruction (hashcat_ctx, sm_major, sm_minor, "__global__ void test () { unsigned int r; asm volatile (\"add.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \ device_param->has_add = cuda_test_instruction (hashcat_ctx, sm_major, sm_minor, "__global__ void test () { unsigned int r; asm volatile (\"add.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \
device_param->has_addc = cuda_test_instruction (hashcat_ctx, sm_major, sm_minor, "__global__ void test () { unsigned int r; asm volatile (\"addc.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \ device_param->has_addc = cuda_test_instruction (hashcat_ctx, sm_major, sm_minor, "__global__ void test () { unsigned int r; asm volatile (\"addc.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \
@ -5480,6 +5494,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
} }
#undef RUN_INSTRUCTION_CHECKS #undef RUN_INSTRUCTION_CHECKS
*/
// device_available_mem // device_available_mem
@ -6214,6 +6229,18 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
if ((device_param->opencl_device_type & CL_DEVICE_TYPE_GPU) && (device_param->opencl_platform_vendor_id == VENDOR_ID_NV)) if ((device_param->opencl_device_type & CL_DEVICE_TYPE_GPU) && (device_param->opencl_platform_vendor_id == VENDOR_ID_NV))
{ {
const int sm = (device_param->sm_major * 10) + device_param->sm_minor;
device_param->has_add = (sm >= 12) ? true : false;
device_param->has_addc = (sm >= 12) ? true : false;
device_param->has_sub = (sm >= 12) ? true : false;
device_param->has_subc = (sm >= 12) ? true : false;
device_param->has_bfe = (sm >= 20) ? true : false;
device_param->has_lop3 = (sm >= 50) ? true : false;
device_param->has_mov64 = (sm >= 10) ? true : false;
device_param->has_prmt = (sm >= 20) ? true : false;
/*
#define RUN_INSTRUCTION_CHECKS() \ #define RUN_INSTRUCTION_CHECKS() \
device_param->has_add = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r; asm volatile (\"add.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \ device_param->has_add = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r; asm volatile (\"add.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \
device_param->has_addc = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r; asm volatile (\"addc.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \ device_param->has_addc = opencl_test_instruction (hashcat_ctx, context, device_param->opencl_device, "__kernel void test () { uint r; asm volatile (\"addc.cc.u32 %0, 0, 0;\" : \"=r\"(r)); }"); \
@ -6250,6 +6277,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
} }
#undef RUN_INSTRUCTION_CHECKS #undef RUN_INSTRUCTION_CHECKS
*/
} }
// device_available_mem // device_available_mem