diff --git a/OpenCL/inc_rp.h b/OpenCL/inc_rp.h
index 25fd41cda..235fe4d45 100644
--- a/OpenCL/inc_rp.h
+++ b/OpenCL/inc_rp.h
@@ -59,17 +59,19 @@
 
 #define RP_PASSWORD_SIZE 256
 
-#ifdef IS_NV
-#define COPY_PW(x)        \
-  pw_t pw = (x);
+#ifdef REAL_SHM
+#define COPY_PW(x)              \
+  __local pw_t s_pws[64];       \
+  s_pws[get_local_id(0)] = (x); \
+  s_pws[get_local_id(0)].pw_len &= 255;
 #else
-#define COPY_PW(x)        \
-  __local pw_t s_pws[64]; \
-  s_pws[get_local_id(0)] = (x);
+#define COPY_PW(x)              \
+  pw_t pw = (x);                \
+  pw.pw_len &= 255;
 #endif
 
-#ifdef IS_NV
-#define PASTE_PW pw;
-#else
+#ifdef REAL_SHM
 #define PASTE_PW s_pws[get_local_id(0)];
+#else
+#define PASTE_PW pw;
 #endif
diff --git a/docs/changes.txt b/docs/changes.txt
index 6322d2cfe..7862e647c 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -48,8 +48,9 @@
 - Fixed missing call to WSACleanup() to cleanly shutdown windows sockets system
 - Fixed missing call to WSAStartup() and client indexing in order to start the brain server on Windows
 - Fixed out-of-bounds write in short-term memory of the brain server
-- Fixed status output of progress value when -S and -l are used in combination
 - Fixed output of --speed-only and --progress-only when fast hashes are used in combination with --slow-candidates
+- Fixed status output of progress value when -S and -l are used in combination
+- Fixed thread count maximum for pure kernels in straight attack mode
 
 ##
 ## Technical
diff --git a/src/opencl.c b/src/opencl.c
index acb3e921d..afb9b4aac 100644
--- a/src/opencl.c
+++ b/src/opencl.c
@@ -6482,7 +6482,22 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
      * now everything that depends on threads and accel, basically dynamic workload
      */
 
-    const u32 kernel_threads = hashconfig_get_kernel_threads (hashcat_ctx, device_param);
+    u32 kernel_threads = hashconfig_get_kernel_threads (hashcat_ctx, device_param);
+
+    // this is required because inside the kernels there is this:
+    // __local pw_t s_pws[64];
+
+    if (user_options->attack_mode == ATTACK_MODE_STRAIGHT)
+    {
+      if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
+      {
+        // not required
+      }
+      else
+      {
+        kernel_threads = MIN (kernel_threads, 64);
+      }
+    }
 
     device_param->kernel_threads = kernel_threads;