Workaround too much register use in -m 14600 on low-end GPU

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;