From ca3beacd93abe1c30a1bb6b33c8ae391f792156d Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Sun, 11 Jul 2021 14:30:49 +0200 Subject: [PATCH] Disable dynamic shared memory on HIP, because hipFuncSetAttribute() maps to cudaFuncSetAttribute() and not to cuFuncSetAttribute() --- include/backend.h | 2 +- src/backend.c | 77 ++++++++++++++++++++++++++--------------------- 2 files changed, 44 insertions(+), 35 deletions(-) diff --git a/include/backend.h b/include/backend.h index a024aa37d..a3d997ca8 100644 --- a/include/backend.h +++ b/include/backend.h @@ -112,7 +112,7 @@ int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream); int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc); -int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value); +//int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value); int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags); int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, HIPstream hStream, void **kernelParams, void **extra); int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize); diff --git a/src/backend.c b/src/backend.c index 5b12cd2c5..eafaec62e 100644 --- a/src/backend.c +++ b/src/backend.c @@ -2454,7 +2454,7 @@ int hip_init (hashcat_ctx_t *hashcat_ctx) HC_LOAD_FUNC_HIP (hip, hipEventRecord, hipEventRecord, HIP_HIPEVENTRECORD, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipEventSynchronize, hipEventSynchronize, HIP_HIPEVENTSYNCHRONIZE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipFuncGetAttribute, hipFuncGetAttribute, HIP_HIPFUNCGETATTRIBUTE, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipFuncSetAttribute, hipFuncSetAttribute, HIP_HIPFUNCSETATTRIBUTE, HIP, 1); + //HC_LOAD_FUNC_HIP (hip, hipFuncSetAttribute, hipFuncSetAttribute, HIP_HIPFUNCSETATTRIBUTE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipFuncSetCacheConfig, hipFuncSetCacheConfig, HIP_HIPFUNCSETCACHECONFIG, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipFuncSetSharedMemConfig, hipFuncSetSharedMemConfig, HIP_HIPFUNCSETSHAREDMEMCONFIG, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipGetErrorName, hipGetErrorName, HIP_HIPGETERRORNAME, HIP, 1); @@ -3079,6 +3079,10 @@ int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_att return 0; } +/* + +// ATTENTION, this one maps to cudaFuncSetAttribute not cuFuncSetAttribute !!! + int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; @@ -3105,6 +3109,7 @@ int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfu return 0; } +*/ int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags) { @@ -5474,7 +5479,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_hipFuncSetAttribute (hashcat_ctx, hip_function, HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shared_mem) == -1) return -1; + //if (hc_hipFuncSetAttribute (hashcat_ctx, hip_function, HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shared_mem) == -1) return -1; } if (kernel_threads == 0) kernel_threads = 1; @@ -10030,6 +10035,9 @@ static int get_hip_kernel_local_mem_size (hashcat_ctx_t *hashcat_ctx, HIPfunctio return 0; } +/* +not supported because there's no cuFuncSetAttribute equivalent + static int get_hip_kernel_dynamic_local_mem_size (hashcat_ctx_t *hashcat_ctx, HIPfunction function, u64 *result) { // AFAIK there's no way to query the maximum value for dynamic shared memory available (because it depends on kernel code). @@ -10063,6 +10071,7 @@ static int get_hip_kernel_dynamic_local_mem_size (hashcat_ctx_t *hashcat_ctx, HI return 0; } +*/ static int get_opencl_kernel_wgs (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_kernel kernel, u32 *result) { @@ -11706,7 +11715,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_memset, &device_param->kernel_local_mem_size_memset) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_memset, &device_param->kernel_dynamic_local_mem_size_memset) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_memset, &device_param->kernel_dynamic_local_mem_size_memset) == -1) return -1; device_param->kernel_preferred_wgs_multiple_memset = device_param->hip_warp_size; @@ -11722,7 +11731,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_atinit, &device_param->kernel_local_mem_size_atinit) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_atinit, &device_param->kernel_dynamic_local_mem_size_atinit) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_atinit, &device_param->kernel_dynamic_local_mem_size_atinit) == -1) return -1; device_param->kernel_preferred_wgs_multiple_atinit = device_param->hip_warp_size; @@ -11737,7 +11746,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_decompress, &device_param->kernel_local_mem_size_decompress) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_decompress, &device_param->kernel_dynamic_local_mem_size_decompress) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_decompress, &device_param->kernel_dynamic_local_mem_size_decompress) == -1) return -1; device_param->kernel_preferred_wgs_multiple_decompress = device_param->hip_warp_size; @@ -11749,7 +11758,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_utf8toutf16le, &device_param->kernel_local_mem_size_utf8toutf16le) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_utf8toutf16le, &device_param->kernel_dynamic_local_mem_size_utf8toutf16le) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_utf8toutf16le, &device_param->kernel_dynamic_local_mem_size_utf8toutf16le) == -1) return -1; device_param->kernel_preferred_wgs_multiple_utf8toutf16le = device_param->hip_warp_size; } @@ -13375,7 +13384,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_local_mem_size1) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; device_param->kernel_preferred_wgs_multiple1 = device_param->hip_warp_size; @@ -13389,7 +13398,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_local_mem_size2) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; device_param->kernel_preferred_wgs_multiple2 = device_param->hip_warp_size; @@ -13403,7 +13412,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_local_mem_size3) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; device_param->kernel_preferred_wgs_multiple3 = device_param->hip_warp_size; } @@ -13417,7 +13426,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_local_mem_size4) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_dynamic_local_mem_size4) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_dynamic_local_mem_size4) == -1) return -1; device_param->kernel_preferred_wgs_multiple4 = device_param->hip_warp_size; } @@ -13436,7 +13445,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_local_mem_size1) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; device_param->kernel_preferred_wgs_multiple1 = device_param->hip_warp_size; @@ -13450,7 +13459,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_local_mem_size2) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; device_param->kernel_preferred_wgs_multiple2 = device_param->hip_warp_size; @@ -13464,7 +13473,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_local_mem_size3) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; device_param->kernel_preferred_wgs_multiple3 = device_param->hip_warp_size; } @@ -13478,7 +13487,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_local_mem_size4) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_dynamic_local_mem_size4) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_dynamic_local_mem_size4) == -1) return -1; device_param->kernel_preferred_wgs_multiple4 = device_param->hip_warp_size; } @@ -13501,7 +13510,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_tm, &device_param->kernel_local_mem_size_tm) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_tm, &device_param->kernel_dynamic_local_mem_size_tm) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_tm, &device_param->kernel_dynamic_local_mem_size_tm) == -1) return -1; device_param->kernel_preferred_wgs_multiple_tm = device_param->hip_warp_size; } @@ -13520,7 +13529,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_local_mem_size1) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; device_param->kernel_preferred_wgs_multiple1 = device_param->hip_warp_size; @@ -13534,7 +13543,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_local_mem_size2) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; device_param->kernel_preferred_wgs_multiple2 = device_param->hip_warp_size; @@ -13548,7 +13557,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_local_mem_size3) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; device_param->kernel_preferred_wgs_multiple3 = device_param->hip_warp_size; @@ -13564,7 +13573,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2p, &device_param->kernel_local_mem_size2p) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2p, &device_param->kernel_dynamic_local_mem_size2p) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2p, &device_param->kernel_dynamic_local_mem_size2p) == -1) return -1; device_param->kernel_preferred_wgs_multiple2p = device_param->hip_warp_size; } @@ -13581,7 +13590,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2e, &device_param->kernel_local_mem_size2e) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2e, &device_param->kernel_dynamic_local_mem_size2e) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2e, &device_param->kernel_dynamic_local_mem_size2e) == -1) return -1; device_param->kernel_preferred_wgs_multiple2e = device_param->hip_warp_size; } @@ -13598,7 +13607,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function12, &device_param->kernel_local_mem_size12) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function12, &device_param->kernel_dynamic_local_mem_size12) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function12, &device_param->kernel_dynamic_local_mem_size12) == -1) return -1; device_param->kernel_preferred_wgs_multiple12 = device_param->hip_warp_size; } @@ -13615,7 +13624,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function23, &device_param->kernel_local_mem_size23) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function23, &device_param->kernel_dynamic_local_mem_size23) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function23, &device_param->kernel_dynamic_local_mem_size23) == -1) return -1; device_param->kernel_preferred_wgs_multiple23 = device_param->hip_warp_size; } @@ -13632,7 +13641,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_init2, &device_param->kernel_local_mem_size_init2) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_init2, &device_param->kernel_dynamic_local_mem_size_init2) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_init2, &device_param->kernel_dynamic_local_mem_size_init2) == -1) return -1; device_param->kernel_preferred_wgs_multiple_init2 = device_param->hip_warp_size; } @@ -13649,7 +13658,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_loop2p, &device_param->kernel_local_mem_size_loop2p) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_loop2p, &device_param->kernel_dynamic_local_mem_size_loop2p) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_loop2p, &device_param->kernel_dynamic_local_mem_size_loop2p) == -1) return -1; device_param->kernel_preferred_wgs_multiple_loop2p = device_param->hip_warp_size; } @@ -13666,7 +13675,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_loop2, &device_param->kernel_local_mem_size_loop2) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_loop2, &device_param->kernel_dynamic_local_mem_size_loop2) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_loop2, &device_param->kernel_dynamic_local_mem_size_loop2) == -1) return -1; device_param->kernel_preferred_wgs_multiple_loop2 = device_param->hip_warp_size; } @@ -13683,7 +13692,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux1, &device_param->kernel_local_mem_size_aux1) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux1, &device_param->kernel_dynamic_local_mem_size_aux1) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux1, &device_param->kernel_dynamic_local_mem_size_aux1) == -1) return -1; device_param->kernel_preferred_wgs_multiple_aux1 = device_param->hip_warp_size; } @@ -13700,7 +13709,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux2, &device_param->kernel_local_mem_size_aux2) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux2, &device_param->kernel_dynamic_local_mem_size_aux2) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux2, &device_param->kernel_dynamic_local_mem_size_aux2) == -1) return -1; device_param->kernel_preferred_wgs_multiple_aux2 = device_param->hip_warp_size; } @@ -13717,7 +13726,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux3, &device_param->kernel_local_mem_size_aux3) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux3, &device_param->kernel_dynamic_local_mem_size_aux3) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux3, &device_param->kernel_dynamic_local_mem_size_aux3) == -1) return -1; device_param->kernel_preferred_wgs_multiple_aux3 = device_param->hip_warp_size; } @@ -13734,7 +13743,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux4, &device_param->kernel_local_mem_size_aux4) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux4, &device_param->kernel_dynamic_local_mem_size_aux4) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux4, &device_param->kernel_dynamic_local_mem_size_aux4) == -1) return -1; device_param->kernel_preferred_wgs_multiple_aux4 = device_param->hip_warp_size; } @@ -13762,7 +13771,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp_l, &device_param->kernel_local_mem_size_mp_l) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp_l, &device_param->kernel_dynamic_local_mem_size_mp_l) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp_l, &device_param->kernel_dynamic_local_mem_size_mp_l) == -1) return -1; device_param->kernel_preferred_wgs_multiple_mp_l = device_param->hip_warp_size; @@ -13774,7 +13783,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp_r, &device_param->kernel_local_mem_size_mp_r) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp_r, &device_param->kernel_dynamic_local_mem_size_mp_r) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp_r, &device_param->kernel_dynamic_local_mem_size_mp_r) == -1) return -1; device_param->kernel_preferred_wgs_multiple_mp_r = device_param->hip_warp_size; @@ -13795,7 +13804,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_local_mem_size_mp) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_dynamic_local_mem_size_mp) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_dynamic_local_mem_size_mp) == -1) return -1; device_param->kernel_preferred_wgs_multiple_mp = device_param->hip_warp_size; } @@ -13807,7 +13816,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_local_mem_size_mp) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_dynamic_local_mem_size_mp) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_dynamic_local_mem_size_mp) == -1) return -1; device_param->kernel_preferred_wgs_multiple_mp = device_param->hip_warp_size; } @@ -13830,7 +13839,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_amp, &device_param->kernel_local_mem_size_amp) == -1) return -1; - if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_amp, &device_param->kernel_dynamic_local_mem_size_amp) == -1) return -1; + //if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_amp, &device_param->kernel_dynamic_local_mem_size_amp) == -1) return -1; device_param->kernel_preferred_wgs_multiple_amp = device_param->hip_warp_size; }