From 2aa529999266f76b7a4d1a88fb4943432333e67b Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Mon, 5 Aug 2019 15:23:00 +0200 Subject: [PATCH] Workaround too much register use in -m 14600 on low-end GPU --- src/modules/module_14600.c | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/src/modules/module_14600.c b/src/modules/module_14600.c index 9b8fcc357..74c962be4 100644 --- a/src/modules/module_14600.c +++ b/src/modules/module_14600.c @@ -178,6 +178,17 @@ typedef struct luks_tmp } luks_tmp_t; +u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + // the module requires a lot of registers for key schedulers on _comp kernel. + // it's possible, if using too many threads, there's not enough registers available, typically ending with misleading error message: + // cuLaunchKernel(): out of memory + + const u32 kernel_threads_max = 64; + + return kernel_threads_max; +} + void *module_benchmark_esalt (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { luks_t *luks = (luks_t *) hcmalloc (sizeof (luks_t)); @@ -628,7 +639,7 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_kernel_accel_min = MODULE_DEFAULT; module_ctx->module_kernel_loops_max = MODULE_DEFAULT; module_ctx->module_kernel_loops_min = MODULE_DEFAULT; - module_ctx->module_kernel_threads_max = MODULE_DEFAULT; + module_ctx->module_kernel_threads_max = module_kernel_threads_max; module_ctx->module_kernel_threads_min = MODULE_DEFAULT; module_ctx->module_kern_type = module_kern_type; module_ctx->module_kern_type_dynamic = module_kern_type_dynamic;