diff --git a/src/interface.c b/src/interface.c
index b737ecbfa..76d784e41 100644
--- a/src/interface.c
+++ b/src/interface.c
@@ -25076,6 +25076,7 @@ u32 hashconfig_forced_kernel_threads (hashcat_ctx_t *hashcat_ctx)
 
 u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param_t *device_param)
 {
+  const hashconfig_t   *hashconfig   = hashcat_ctx->hashconfig;
   const user_options_t *user_options = hashcat_ctx->user_options;
 
   // a kernel can force a fixed value
@@ -25096,27 +25097,25 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p
   {
     if (device_param->device_vendor_id == VENDOR_ID_NV)
     {
-      kernel_threads = KERNEL_THREADS_MAX_GPU_NV;
-
-      switch (user_options->workload_profile)
+      if (hashconfig->attack_exec == ATTACK_EXEC_OUTSIDE_KERNEL)
       {
-        case 1: kernel_threads *= 1; break;
-        case 2: kernel_threads *= 2; break;
-        case 3: kernel_threads *= 4; break;
-        case 4: kernel_threads *= 8; break;
+        if (user_options->workload_profile < 4)
+        {
+          kernel_threads = KERNEL_THREADS_MAX_GPU_NV;
+        }
+        else
+        {
+          kernel_threads = device_param->device_maxworkgroup_size;
+        }
+      }
+      else
+      {
+        kernel_threads = device_param->device_maxworkgroup_size;
       }
     }
     else if (device_param->device_vendor_id == VENDOR_ID_AMD)
     {
       kernel_threads = KERNEL_THREADS_MAX_GPU_AMD;
-
-      switch (user_options->workload_profile)
-      {
-        case 1: kernel_threads *= 1; break;
-        case 2: kernel_threads *= 1; break;
-        case 3: kernel_threads *= 2; break;
-        case 4: kernel_threads *= 4; break;
-      }
     }
     else
     {
diff --git a/src/opencl.c b/src/opencl.c
index 1dc885aac..a99c1a91f 100644
--- a/src/opencl.c
+++ b/src/opencl.c
@@ -2266,6 +2266,13 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co
 
       if (user_options->speed_only == true)
       {
+        for (int i = 0; i < 16; i++)
+        {
+          const int rc = choose_kernel (hashcat_ctx, device_param, highest_pw_len, pws_cnt, fast_iteration, salt_pos);
+
+          if (rc == -1) return -1;
+        }
+
         hc_timer_set (&device_param->timer_speed);
       }
 
@@ -3306,10 +3313,8 @@ int opencl_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
               #if defined (__linux__)
               if (device_param->is_rocm == false)
               {
-                // ROCm is so much better, give the user some hint.
-                // They can still use --force to use AMDGPU-Pro.
+                // ROCm is so much better, we should give the user some hint and remove this block
 
-                /*
                 // AMDGPU-PRO Driver 16.40 and higher
                 if (atoi (device_param->driver_version) >= 2117) amd_warn = false;
                 // AMDGPU-PRO Driver 16.50 is known to be broken
@@ -3319,7 +3324,6 @@ int opencl_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime)
                 // AMDGPU-PRO Driver 17.10 is known to be broken
                 if (atoi (device_param->driver_version) == 2348) amd_warn = true;
                 // AMDGPU-PRO Driver 17.20 (2416) is fine, doesn't need check will match >= 2117
-                */
               }
               else
               {
@@ -3675,6 +3679,19 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
 
     #endif // __APPLE__
 
+    if (device_param->platform_vendor_id == VENDOR_ID_AMD)
+    {
+      if (device_param->is_rocm == false)
+      {
+        if ((user_options->hash_mode ==  7900)
+         || (user_options->hash_mode == 10700)
+         || (user_options->hash_mode == 13731))
+        {
+          skipped_temp = true;
+        }
+      }
+    }
+
     if ((skipped_temp == true) && (user_options->force == false))
     {
       event_log_warning (hashcat_ctx, "* Device #%u: Skipping unstable hash-mode %u for this device.", device_id + 1, user_options->hash_mode);