diff --git a/include/interface.h b/include/interface.h
index 26397282a..3da17c2e4 100644
--- a/include/interface.h
+++ b/include/interface.h
@@ -1539,6 +1539,7 @@ void ascii_digest (char *out_buf, const uint salt_pos, const uint digest_pos, co
 
 int   hashconfig_init                (hashconfig_t *hashconfig, const user_options_t *user_options);
 void  hashconfig_destroy             (hashconfig_t *hashconfig);
+u32   hashconfig_kernel_thread_force (const hashconfig_t *hashconfig, const hc_device_param_t *device_param);
 uint  hashconfig_general_pw_min      (hashconfig_t *hashconfig);
 uint  hashconfig_general_pw_max      (hashconfig_t *hashconfig);
 void  hashconfig_general_defaults    (hashconfig_t *hashconfig, hashes_t *hashes, const user_options_t *user_options);
diff --git a/src/interface.c b/src/interface.c
index a02bb5f7f..7a9f4b78c 100644
--- a/src/interface.c
+++ b/src/interface.c
@@ -19919,6 +19919,39 @@ void hashconfig_destroy (hashconfig_t *hashconfig)
   myfree (hashconfig);
 }
 
+u32 hashconfig_kernel_thread_force (const hashconfig_t *hashconfig, const hc_device_param_t *device_param)
+{
+  u32 kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size);
+
+  if (hashconfig->hash_mode ==  8900) kernel_threads = 64; // Scrypt
+  if (hashconfig->hash_mode ==  9300) kernel_threads = 64; // Scrypt
+
+  if (device_param->device_type & CL_DEVICE_TYPE_CPU)
+  {
+    kernel_threads = KERNEL_THREADS_MAX_CPU;
+  }
+
+  if (hashconfig->hash_mode ==  1500) kernel_threads = 64; // DES
+  if (hashconfig->hash_mode ==  3000) kernel_threads = 64; // DES
+  if (hashconfig->hash_mode ==  3100) kernel_threads = 64; // DES
+  if (hashconfig->hash_mode ==  3200) kernel_threads = 8;  // Blowfish
+  if (hashconfig->hash_mode ==  7500) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode ==  8500) kernel_threads = 64; // DES
+  if (hashconfig->hash_mode ==  9000) kernel_threads = 8;  // Blowfish
+  if (hashconfig->hash_mode ==  9700) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode ==  9710) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode ==  9800) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode ==  9810) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode == 10400) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode == 10410) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode == 10500) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode == 13100) kernel_threads = 64; // RC4
+  if (hashconfig->hash_mode == 14000) kernel_threads = 64; // DES
+  if (hashconfig->hash_mode == 14100) kernel_threads = 64; // DES
+
+  return kernel_threads;
+}
+
 uint hashconfig_general_pw_min (hashconfig_t *hashconfig)
 {
   uint pw_min = PW_MIN;
diff --git a/src/opencl.c b/src/opencl.c
index e1032888d..919f3bbd6 100644
--- a/src/opencl.c
+++ b/src/opencl.c
@@ -2641,33 +2641,7 @@ int opencl_session_begin (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconf
      *                 there needs to be some upper limit, otherwise there's too much overhead
      */
 
-    uint kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size);
-
-    if (hashconfig->hash_mode ==  8900) kernel_threads = 64; // Scrypt
-    if (hashconfig->hash_mode ==  9300) kernel_threads = 64; // Scrypt
-
-    if (device_param->device_type & CL_DEVICE_TYPE_CPU)
-    {
-      kernel_threads = KERNEL_THREADS_MAX_CPU;
-    }
-
-    if (hashconfig->hash_mode ==  1500) kernel_threads = 64; // DES
-    if (hashconfig->hash_mode ==  3000) kernel_threads = 64; // DES
-    if (hashconfig->hash_mode ==  3100) kernel_threads = 64; // DES
-    if (hashconfig->hash_mode ==  3200) kernel_threads = 8;  // Blowfish
-    if (hashconfig->hash_mode ==  7500) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode ==  8500) kernel_threads = 64; // DES
-    if (hashconfig->hash_mode ==  9000) kernel_threads = 8;  // Blowfish
-    if (hashconfig->hash_mode ==  9700) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode ==  9710) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode ==  9800) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode ==  9810) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode == 10400) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode == 10410) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode == 10500) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode == 13100) kernel_threads = 64; // RC4
-    if (hashconfig->hash_mode == 14000) kernel_threads = 64; // DES
-    if (hashconfig->hash_mode == 14100) kernel_threads = 64; // DES
+    u32 kernel_threads = hashconfig_kernel_thread_force (hashconfig, device_param);
 
     device_param->kernel_threads = kernel_threads;