From d59474fded17b1c4eacda8dd8f0adbeb6a3a0a97 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Fri, 10 May 2019 17:27:15 +0200 Subject: [PATCH] Testwise unlock full thread count on NVidia --- src/backend.c | 58 +++++++++++++++++++++++++-------------------------- 1 file changed, 29 insertions(+), 29 deletions(-) diff --git a/src/backend.c b/src/backend.c index e7e31ae11..726594fe6 100644 --- a/src/backend.c +++ b/src/backend.c @@ -8782,7 +8782,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple1 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple1 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -8802,7 +8802,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple2 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple2 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -8822,7 +8822,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple3 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple3 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -8842,7 +8842,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple4 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple4 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -8867,7 +8867,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple1 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple1 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -8887,7 +8887,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple2 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple2 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -8907,7 +8907,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple3 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple3 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -8927,7 +8927,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple4 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple4 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -8956,7 +8956,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_tm = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_tm = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -8981,7 +8981,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple1 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple1 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9001,7 +9001,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple2 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple2 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9021,7 +9021,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple3 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple3 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9043,7 +9043,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple12 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple12 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9066,7 +9066,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple23 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple23 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9089,7 +9089,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_init2 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_init2 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9112,7 +9112,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_loop2 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_loop2 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9135,7 +9135,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_aux1 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_aux1 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9158,7 +9158,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_aux2 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_aux2 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9181,7 +9181,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_aux3 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_aux3 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9204,7 +9204,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_aux4 = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_aux4 = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9224,7 +9224,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_memset = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_memset = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9246,7 +9246,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_atinit = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_atinit = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9267,7 +9267,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_decompress = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_decompress = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9299,7 +9299,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_mp_l = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_mp_l = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9317,7 +9317,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_mp_r = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_mp_r = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; @@ -9341,7 +9341,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_mp = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_mp = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9359,7 +9359,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_mp = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_mp = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -9388,7 +9388,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; - device_param->kernel_preferred_wgs_multiple_amp = device_param->cuda_warp_size; + device_param->kernel_preferred_wgs_multiple_amp = device_param->device_maxworkgroup_size; if (CL_rc == -1) return -1; } @@ -10317,7 +10317,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) // this value should represent a reasonable amount of memory a host system has per GPU. // note we're allocating 3 blocks of that size. - const u64 PWS_SPACE = 4ull * 1024ull * 1024ull * 1024ull; + const u64 PWS_SPACE = 1024ull * 1024ull * 1024ull; // sometimes device_available_mem and device_maxmem_alloc reported back from the opencl runtime are a bit inaccurate. // let's add some extra space just to be sure.