diff --git a/OpenCL/inc_rp.cl b/OpenCL/inc_rp.cl
index c45a7c591..bad58a1f0 100644
--- a/OpenCL/inc_rp.cl
+++ b/OpenCL/inc_rp.cl
@@ -7,7 +7,21 @@
 #define MAYBE_UNUSED
 #endif
 
-u32 generate_cmask (const u32 value);
+#ifdef REAL_SHM
+#define COPY_PW(x)                \
+  __local pw_t s_pws[64];         \
+  s_pws[get_local_id (0)] = (x);
+#else
+#define COPY_PW(x)                \
+  pw_t pw = (x);
+#endif
+
+#ifdef REAL_SHM
+#define PASTE_PW s_pws[get_local_id(0)];
+#else
+#define PASTE_PW pw;
+#endif
+
 u32 generate_cmask (const u32 value)
 {
   const u32 rmask =  ((value & 0x40404040u) >> 1u)
@@ -19,7 +33,6 @@ u32 generate_cmask (const u32 value)
   return rmask & ~hmask & lmask;
 }
 
-void append_four_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst);
 void append_four_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst)
 {
   const int sd  = off_src / 4;
@@ -43,7 +56,6 @@ void append_four_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, cons
   buf_dst[dd + 1] |= t1;
 }
 
-void append_three_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst);
 void append_three_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst)
 {
   const int sd  = off_src / 4;
@@ -67,7 +79,6 @@ void append_three_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, con
   buf_dst[dd + 1] |= t1;
 }
 
-void append_two_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst);
 void append_two_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst)
 {
   const int sd  = off_src / 4;
@@ -91,7 +102,6 @@ void append_two_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const
   buf_dst[dd + 1] |= t1;
 }
 
-void append_one_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst);
 void append_one_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst)
 {
   const int sd  = off_src / 4;
@@ -111,7 +121,6 @@ void append_one_byte (const u32 *buf_src, const int off_src, u32 *buf_dst, const
   buf_dst[dd] |= t;
 }
 
-void append_block (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst, const int len);
 void append_block (const u32 *buf_src, const int off_src, u32 *buf_dst, const int off_dst, const int len)
 {
   int i;
@@ -131,7 +140,6 @@ void append_block (const u32 *buf_src, const int off_src, u32 *buf_dst, const in
   }
 }
 
-void exchange_byte (u32 *buf, const int off_src, const int off_dst);
 void exchange_byte (u32 *buf, const int off_src, const int off_dst)
 {
   u8 *ptr = (u8 *) buf;
@@ -171,7 +179,6 @@ void exchange_byte (u32 *buf, const int off_src, const int off_dst)
   */
 }
 
-int mangle_lrest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_lrest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
@@ -184,7 +191,6 @@ int mangle_lrest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf,
   return (len);
 }
 
-int mangle_lrest_ufirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_lrest_ufirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
@@ -201,7 +207,6 @@ int mangle_lrest_ufirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
   return (len);
 }
 
-int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
@@ -214,7 +219,6 @@ int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf,
   return (len);
 }
 
-int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
@@ -231,7 +235,6 @@ int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
   return (len);
 }
 
-int mangle_trest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_trest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
@@ -244,7 +247,6 @@ int mangle_trest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf,
   return (len);
 }
 
-int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -261,7 +263,6 @@ int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *b
   return (len);
 }
 
-int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int l = 0; l < len / 2; l++)
@@ -274,7 +275,6 @@ int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf
   return (len);
 }
 
-int mangle_dupeword (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_dupeword (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   const int out_len = len * 2;
@@ -286,7 +286,6 @@ int mangle_dupeword (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *bu
   return (out_len);
 }
 
-int mangle_dupeword_times (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_dupeword_times (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   const int out_len = (len * p0) + len;
@@ -300,7 +299,6 @@ int mangle_dupeword_times (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u
   return (out_len);
 }
 
-int mangle_reflect (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_reflect (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   const int out_len = len * 2;
@@ -319,7 +317,6 @@ int mangle_reflect (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf
   return out_len;
 }
 
-int mangle_append (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_append (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   const int out_len = len + 1;
@@ -331,7 +328,6 @@ int mangle_append (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf,
   return (out_len);
 }
 
-int mangle_prepend (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_prepend (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   const int out_len = len + 1;
@@ -348,7 +344,6 @@ int mangle_prepend (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf,
   return (out_len);
 }
 
-int mangle_rotate_left (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_rotate_left (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int l = 0, r = len - 1; r > l; r--)
@@ -359,7 +354,6 @@ int mangle_rotate_left (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
   return (len);
 }
 
-int mangle_rotate_right (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_rotate_right (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int l = 0, r = len - 1; l < r; l++)
@@ -370,7 +364,6 @@ int mangle_rotate_right (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
   return (len);
 }
 
-int mangle_delete_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_delete_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -385,13 +378,11 @@ int mangle_delete_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *bu
   return (len - 1);
 }
 
-int mangle_delete_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_delete_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   return mangle_delete_at (0, p1, buf, len);
 }
 
-int mangle_delete_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_delete_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (len == 0) return 0;
@@ -399,7 +390,6 @@ int mangle_delete_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *
   return mangle_delete_at (len - 1, p1, buf, len);
 }
 
-int mangle_extract (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_extract (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -419,7 +409,6 @@ int mangle_extract (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf,
   return (p1);
 }
 
-int mangle_omit (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_omit (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -439,7 +428,6 @@ int mangle_omit (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, co
   return (len - p1);
 }
 
-int mangle_insert (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_insert (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len + 1) return (len);
@@ -458,7 +446,6 @@ int mangle_insert (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf,
   return (out_len);
 }
 
-int mangle_overstrike (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_overstrike (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -468,7 +455,6 @@ int mangle_overstrike (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *b
   return (len);
 }
 
-int mangle_truncate_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_truncate_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -481,7 +467,6 @@ int mangle_truncate_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *
   return (p0);
 }
 
-int mangle_replace (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_replace (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   for (int pos = 0; pos < len; pos++)
@@ -494,7 +479,6 @@ int mangle_replace (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf,
   return (len);
 }
 
-int mangle_purgechar (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_purgechar (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   int out_len = 0;
@@ -516,7 +500,6 @@ int mangle_purgechar (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *bu
   return (out_len);
 }
 
-int mangle_dupechar_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_dupechar_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   const int out_len = len + p0;
@@ -533,7 +516,6 @@ int mangle_dupechar_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u
   return (out_len);
 }
 
-int mangle_dupechar_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_dupechar_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   const int out_len = len + p0;
@@ -551,7 +533,6 @@ int mangle_dupechar_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8
   return (out_len);
 }
 
-int mangle_dupechar_all (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_dupechar_all (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   const int out_len = len + len;
@@ -570,7 +551,6 @@ int mangle_dupechar_all (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8
   return (out_len);
 }
 
-int mangle_switch_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_switch_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   if (len < 2) return (len);
@@ -580,7 +560,6 @@ int mangle_switch_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
   return (len);
 }
 
-int mangle_switch_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_switch_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   if (len < 2) return (len);
@@ -590,7 +569,6 @@ int mangle_switch_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
   return (len);
 }
 
-int mangle_switch_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_switch_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -601,7 +579,6 @@ int mangle_switch_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *b
   return (len);
 }
 
-int mangle_chr_shiftl (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_chr_shiftl (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -611,7 +588,6 @@ int mangle_chr_shiftl (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *b
   return (len);
 }
 
-int mangle_chr_shiftr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_chr_shiftr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -621,7 +597,6 @@ int mangle_chr_shiftr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *b
   return (len);
 }
 
-int mangle_chr_incr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_chr_incr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -631,7 +606,6 @@ int mangle_chr_incr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf
   return (len);
 }
 
-int mangle_chr_decr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_chr_decr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -641,7 +615,6 @@ int mangle_chr_decr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf
   return (len);
 }
 
-int mangle_replace_np1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_replace_np1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if ((p0 + 1) >= len) return (len);
@@ -651,7 +624,6 @@ int mangle_replace_np1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *
   return (len);
 }
 
-int mangle_replace_nm1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_replace_nm1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 == 0) return (len);
@@ -663,7 +635,6 @@ int mangle_replace_nm1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *
   return (len);
 }
 
-int mangle_dupeblock_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_dupeblock_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -682,7 +653,6 @@ int mangle_dupeblock_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1,
   return (out_len);
 }
 
-int mangle_dupeblock_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
 int mangle_dupeblock_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
 {
   if (p0 >= len) return (len);
@@ -701,7 +671,6 @@ int mangle_dupeblock_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u
   return (out_len);
 }
 
-int mangle_title_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 int mangle_title_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   if ((len + 4) >= RP_PASSWORD_SIZE) return (len); // cheap way to not need to check for overflow of i + 1
@@ -727,7 +696,6 @@ int mangle_title_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *b
   return (len);
 }
 
-int apply_rule (const u32 name, MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int in_len);
 int apply_rule (const u32 name, MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int in_len)
 {
   int out_len = in_len;
@@ -779,7 +747,6 @@ int apply_rule (const u32 name, MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8
   return out_len;
 }
 
-int apply_rules (__constant const u32 *cmds, u32 *buf, const int in_len);
 int apply_rules (__constant const u32 *cmds, u32 *buf, const int in_len)
 {
   int out_len = in_len;
diff --git a/OpenCL/inc_rp.h b/OpenCL/inc_rp.h
index b0e37b28a..4694fb19e 100644
--- a/OpenCL/inc_rp.h
+++ b/OpenCL/inc_rp.h
@@ -58,18 +58,3 @@
 #define RULE_OP_MANGLE_TITLE            'E'
 
 #define RP_PASSWORD_SIZE 256
-
-#ifdef REAL_SHM
-#define COPY_PW(x)              \
-  __local pw_t s_pws[64];       \
-  s_pws[get_local_id(0)] = (x);
-#else
-#define COPY_PW(x)              \
-  pw_t pw = (x);
-#endif
-
-#ifdef REAL_SHM
-#define PASTE_PW s_pws[get_local_id(0)];
-#else
-#define PASTE_PW pw;
-#endif
diff --git a/OpenCL/inc_rp_optimized.cl b/OpenCL/inc_rp_optimized.cl
index 03b9039e7..6c81767a3 100644
--- a/OpenCL/inc_rp_optimized.cl
+++ b/OpenCL/inc_rp_optimized.cl
@@ -7,7 +7,6 @@
 #define MAYBE_UNUSED
 #endif
 
-u32 generate_cmask (const u32 value);
 u32 generate_cmask (const u32 value)
 {
   const u32 rmask =  ((value & 0x40404040u) >> 1u)
@@ -19,7 +18,6 @@ u32 generate_cmask (const u32 value)
   return rmask & ~hmask & lmask;
 }
 
-void truncate_right (u32 *buf0, u32 *buf1, const u32 offset);
 void truncate_right (u32 *buf0, u32 *buf1, const u32 offset)
 {
   const u32 tmp = (1u << ((offset & 3u) * 8u)) - 1u;
@@ -79,7 +77,6 @@ void truncate_right (u32 *buf0, u32 *buf1, const u32 offset)
   }
 }
 
-void truncate_left (u32 *buf0, u32 *buf1, const u32 offset);
 void truncate_left (u32 *buf0, u32 *buf1, const u32 offset)
 {
   const u32 tmp = ~((1u << ((offset & 3u) * 8u)) - 1u);
@@ -139,7 +136,6 @@ void truncate_left (u32 *buf0, u32 *buf1, const u32 offset)
   }
 }
 
-void lshift_block (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1);
 void lshift_block (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1)
 {
   out0[0] = hc_bytealign_S (in0[1], in0[0], 1);
@@ -152,7 +148,6 @@ void lshift_block (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1)
   out1[3] = hc_bytealign_S (     0, in1[3], 1);
 }
 
-void rshift_block (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1);
 void rshift_block (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1)
 {
   out1[3] = hc_bytealign_S (in1[3], in1[2], 3);
@@ -165,7 +160,6 @@ void rshift_block (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1)
   out0[0] = hc_bytealign_S (in0[0],      0, 3);
 }
 
-void lshift_block_N (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1, const u32 num);
 void lshift_block_N (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1, const u32 num)
 {
   switch (num)
@@ -461,7 +455,6 @@ void lshift_block_N (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1, const
   }
 }
 
-void rshift_block_N (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1, const u32 num);
 void rshift_block_N (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1, const u32 num)
 {
   switch (num)
@@ -757,7 +750,6 @@ void rshift_block_N (const u32 *in0, const u32 *in1, u32 *out0, u32 *out1, const
   }
 }
 
-void append_block1 (const u32 offset, u32 *buf0, u32 *buf1, const u32 src_r0);
 void append_block1 (const u32 offset, u32 *buf0, u32 *buf1, const u32 src_r0)
 {
   // this version works with 1 byte append only
@@ -779,7 +771,6 @@ void append_block1 (const u32 offset, u32 *buf0, u32 *buf1, const u32 src_r0)
 
 }
 
-void append_block8 (const u32 offset, u32 *buf0, u32 *buf1, const u32 *src_l0, const u32 *src_l1, const u32 *src_r0, const u32 *src_r1);
 void append_block8 (const u32 offset, u32 *buf0, u32 *buf1, const u32 *src_l0, const u32 *src_l1, const u32 *src_r0, const u32 *src_r1)
 {
   u32 s0 = 0;
@@ -1035,7 +1026,6 @@ void append_block8 (const u32 offset, u32 *buf0, u32 *buf1, const u32 *src_l0, c
   buf1[3] = src_l1[3] | s7;
 }
 
-void reverse_block (u32 *in0, u32 *in1, u32 *out0, u32 *out1, const u32 len);
 void reverse_block (u32 *in0, u32 *in1, u32 *out0, u32 *out1, const u32 len)
 {
   rshift_block_N (in0, in1, out0, out1, 32 - len);
@@ -1062,7 +1052,6 @@ void reverse_block (u32 *in0, u32 *in1, u32 *out0, u32 *out1, const u32 len)
   out1[3] = swap32_S (tib41[3]);
 }
 
-void exchange_byte (u32 *buf, const int off_src, const int off_dst);
 void exchange_byte (u32 *buf, const int off_src, const int off_dst)
 {
   u8 *ptr = (u8 *) buf;
@@ -1073,7 +1062,6 @@ void exchange_byte (u32 *buf, const int off_src, const int off_dst)
   ptr[off_dst] = tmp;
 }
 
-u32 rule_op_mangle_lrest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_lrest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   u32 t;
@@ -1090,7 +1078,6 @@ u32 rule_op_mangle_lrest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1,
   return in_len;
 }
 
-u32 rule_op_mangle_urest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_urest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   u32 t;
@@ -1107,7 +1094,6 @@ u32 rule_op_mangle_urest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1,
   return in_len;
 }
 
-u32 rule_op_mangle_lrest_ufirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_lrest_ufirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   u32 t;
@@ -1126,7 +1112,6 @@ u32 rule_op_mangle_lrest_ufirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u
   return in_len;
 }
 
-u32 rule_op_mangle_urest_lfirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_urest_lfirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   u32 t;
@@ -1145,7 +1130,6 @@ u32 rule_op_mangle_urest_lfirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u
   return in_len;
 }
 
-u32 rule_op_mangle_trest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_trest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   u32 t;
@@ -1162,7 +1146,6 @@ u32 rule_op_mangle_trest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1,
   return in_len;
 }
 
-u32 rule_op_mangle_toggle_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_toggle_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -1196,7 +1179,6 @@ u32 rule_op_mangle_toggle_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return (in_len);
 }
 
-u32 rule_op_mangle_reverse (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_reverse (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   reverse_block (buf0, buf1, buf0, buf1, in_len);
@@ -1204,7 +1186,6 @@ u32 rule_op_mangle_reverse (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1
   return in_len;
 }
 
-u32 rule_op_mangle_dupeword (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_dupeword (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ((in_len + in_len) >= 32) return in_len;
@@ -1218,7 +1199,6 @@ u32 rule_op_mangle_dupeword (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p
   return out_len;
 }
 
-u32 rule_op_mangle_dupeword_times (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_dupeword_times (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (((in_len * p0) + in_len) >= 32) return in_len;
@@ -1247,7 +1227,6 @@ u32 rule_op_mangle_dupeword_times (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const
   return out_len;
 }
 
-u32 rule_op_mangle_reflect (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_reflect (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ((in_len + in_len) >= 32) return in_len;
@@ -1266,7 +1245,6 @@ u32 rule_op_mangle_reflect (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1
   return out_len;
 }
 
-u32 rule_op_mangle_append (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_append (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ((in_len + 1) >= 32) return in_len;
@@ -1280,7 +1258,6 @@ u32 rule_op_mangle_append (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1,
   return out_len;
 }
 
-u32 rule_op_mangle_prepend (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_prepend (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ((in_len + 1) >= 32) return in_len;
@@ -1296,7 +1273,6 @@ u32 rule_op_mangle_prepend (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1
   return out_len;
 }
 
-u32 rule_op_mangle_rotate_left (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_rotate_left (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (in_len == 0) return in_len;
@@ -1312,7 +1288,6 @@ u32 rule_op_mangle_rotate_left (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u3
   return in_len;
 }
 
-u32 rule_op_mangle_rotate_right (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_rotate_right (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (in_len == 0) return in_len;
@@ -1355,7 +1330,6 @@ u32 rule_op_mangle_rotate_right (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u
   return in_len;
 }
 
-u32 rule_op_mangle_delete_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_delete_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (in_len == 0) return in_len;
@@ -1367,7 +1341,6 @@ u32 rule_op_mangle_delete_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u
   return in_len1;
 }
 
-u32 rule_op_mangle_delete_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_delete_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (in_len == 0) return in_len;
@@ -1388,7 +1361,6 @@ u32 rule_op_mangle_delete_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u3
   return in_len1;
 }
 
-u32 rule_op_mangle_delete_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_delete_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -1470,7 +1442,6 @@ u32 rule_op_mangle_delete_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return out_len;
 }
 
-u32 rule_op_mangle_extract (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_extract (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -1486,7 +1457,6 @@ u32 rule_op_mangle_extract (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1
   return out_len;
 }
 
-u32 rule_op_mangle_omit (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_omit (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -1579,7 +1549,6 @@ u32 rule_op_mangle_omit (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, M
   return out_len;
 }
 
-u32 rule_op_mangle_insert (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_insert (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 > in_len) return in_len;
@@ -1658,7 +1627,6 @@ u32 rule_op_mangle_insert (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1,
   return out_len;
 }
 
-u32 rule_op_mangle_overstrike (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_overstrike (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -1694,7 +1662,6 @@ u32 rule_op_mangle_overstrike (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return in_len;
 }
 
-u32 rule_op_mangle_truncate_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_truncate_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -1704,7 +1671,6 @@ u32 rule_op_mangle_truncate_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u3
   return p0;
 }
 
-u32 search_on_register (const u32 in, const u32 p0);
 u32 search_on_register (const u32 in, const u32 p0)
 {
   u32 r = 0;
@@ -1717,7 +1683,6 @@ u32 search_on_register (const u32 in, const u32 p0)
   return r;
 }
 
-u32 replace_on_register (const u32 in, const u32 r, const u32 p1);
 u32 replace_on_register (const u32 in, const u32 r, const u32 p1)
 {
   u32 out = in;
@@ -1730,7 +1695,6 @@ u32 replace_on_register (const u32 in, const u32 r, const u32 p1)
   return out;
 }
 
-u32 rule_op_mangle_replace (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_replace (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   const u32 r0 = search_on_register (buf0[0], p0);
@@ -1758,7 +1722,6 @@ u32 rule_op_mangle_replace (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1
   return in_len;
 }
 
-u32 rule_op_mangle_purgechar (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_purgechar (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   const u32 r0 = search_on_register (buf0[0], p0);
@@ -1813,7 +1776,6 @@ u32 rule_op_mangle_purgechar (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return out_len;
 }
 
-u32 rule_op_mangle_dupechar_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_dupechar_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ( in_len       ==  0) return in_len;
@@ -1849,7 +1811,6 @@ u32 rule_op_mangle_dupechar_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const
   return out_len;
 }
 
-u32 rule_op_mangle_dupechar_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_dupechar_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ( in_len       ==  0) return in_len;
@@ -1896,7 +1857,6 @@ u32 rule_op_mangle_dupechar_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const
   return out_len;
 }
 
-u32 rule_op_mangle_dupechar_all (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_dupechar_all (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ( in_len           ==  0) return in_len;
@@ -1930,7 +1890,6 @@ u32 rule_op_mangle_dupechar_all (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u
   return out_len;
 }
 
-u32 rule_op_mangle_switch_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_switch_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (in_len < 2) return in_len;
@@ -1940,7 +1899,6 @@ u32 rule_op_mangle_switch_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u
   return in_len;
 }
 
-u32 rule_op_mangle_switch_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_switch_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (in_len < 2) return in_len;
@@ -1970,7 +1928,6 @@ u32 rule_op_mangle_switch_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u3
   return in_len;
 }
 
-u32 rule_op_mangle_switch_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_switch_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -2001,7 +1958,6 @@ u32 rule_op_mangle_switch_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return in_len;
 }
 
-u32 rule_op_mangle_chr_shiftl (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_chr_shiftl (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -2036,7 +1992,6 @@ u32 rule_op_mangle_chr_shiftl (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return in_len;
 }
 
-u32 rule_op_mangle_chr_shiftr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_chr_shiftr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -2071,7 +2026,6 @@ u32 rule_op_mangle_chr_shiftr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return in_len;
 }
 
-u32 rule_op_mangle_chr_incr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_chr_incr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -2108,7 +2062,6 @@ u32 rule_op_mangle_chr_incr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p
   return in_len;
 }
 
-u32 rule_op_mangle_chr_decr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_chr_decr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 >= in_len) return in_len;
@@ -2145,7 +2098,6 @@ u32 rule_op_mangle_chr_decr (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p
   return in_len;
 }
 
-u32 rule_op_mangle_replace_np1 (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_replace_np1 (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if ((p0 + 1) >= in_len) return in_len;
@@ -2186,7 +2138,6 @@ u32 rule_op_mangle_replace_np1 (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u3
   return in_len;
 }
 
-u32 rule_op_mangle_replace_nm1 (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_replace_nm1 (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 == 0) return in_len;
@@ -2229,7 +2180,6 @@ u32 rule_op_mangle_replace_nm1 (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u3
   return in_len;
 }
 
-u32 rule_op_mangle_dupeblock_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_dupeblock_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 > in_len) return in_len;
@@ -2268,7 +2218,6 @@ u32 rule_op_mangle_dupeblock_first (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED cons
   return out_len;
 }
 
-u32 rule_op_mangle_dupeblock_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_dupeblock_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (p0 > in_len) return in_len;
@@ -2298,7 +2247,6 @@ u32 rule_op_mangle_dupeblock_last (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const
   return out_len;
 }
 
-u32 toggle_on_register (const u32 in, const u32 r);
 u32 toggle_on_register (const u32 in, const u32 r)
 {
   u32 out = in;
@@ -2313,7 +2261,6 @@ u32 toggle_on_register (const u32 in, const u32 r)
   return out;
 }
 
-u32 rule_op_mangle_title_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 u32 rule_op_mangle_title_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   if (in_len == 0) return in_len;
@@ -2354,7 +2301,6 @@ u32 rule_op_mangle_title_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32
   return in_len;
 }
 
-u32 apply_rule (const u32 name, const u32 p0, const u32 p1, u32 *buf0, u32 *buf1, const u32 in_len);
 u32 apply_rule (const u32 name, const u32 p0, const u32 p1, u32 *buf0, u32 *buf1, const u32 in_len)
 {
   u32 out_len = in_len;
@@ -2407,7 +2353,6 @@ u32 apply_rule (const u32 name, const u32 p0, const u32 p1, u32 *buf0, u32 *buf1
   return out_len;
 }
 
-u32 apply_rules (__constant const u32 *cmds, u32 *buf0, u32 *buf1, const u32 len);
 u32 apply_rules (__constant const u32 *cmds, u32 *buf0, u32 *buf1, const u32 len)
 {
   u32 out_len = len;
@@ -2426,7 +2371,6 @@ u32 apply_rules (__constant const u32 *cmds, u32 *buf0, u32 *buf1, const u32 len
   return out_len;
 }
 
-u32x apply_rules_vect (const u32 *pw_buf0, const u32 *pw_buf1, const u32 pw_len, __constant const kernel_rule_t *rules_buf, const u32 il_pos, u32x *buf0, u32x *buf1);
 u32x apply_rules_vect (const u32 *pw_buf0, const u32 *pw_buf1, const u32 pw_len, __constant const kernel_rule_t *rules_buf, const u32 il_pos, u32x *buf0, u32x *buf1)
 {
   #if VECT_SIZE == 1