diff --git a/OpenCL/m19000-pure.cl b/OpenCL/m19000-pure.cl
new file mode 100644
index 000000000..7e96edc63
--- /dev/null
+++ b/OpenCL/m19000-pure.cl
@@ -0,0 +1,93 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+//#define NEW_SIMD_CODE
+
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
+#include "inc_simd.cl"
+#include "inc_hash_md5.cl"
+
+#define COMPARE_S "inc_comp_single.cl"
+#define COMPARE_M "inc_comp_multi.cl"
+
+typedef struct qnx_md5_tmp
+{
+  md5_ctx_t md5_ctx;
+
+} qnx_md5_tmp_t;
+
+__kernel void m19000_init (KERN_ATTR_TMPS (qnx_md5_tmp_t))
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * init
+   */
+
+  md5_ctx_t md5_ctx;
+
+  md5_init (&md5_ctx);
+
+  md5_update_global (&md5_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
+
+  md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len);
+
+  tmps[gid].md5_ctx = md5_ctx;
+}
+
+__kernel void m19000_loop (KERN_ATTR_TMPS (qnx_md5_tmp_t))
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  md5_ctx_t md5_ctx = tmps[gid].md5_ctx;
+
+  for (u32 i = 0; i < loop_cnt; i++)
+  {
+    md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len);
+  }
+
+  tmps[gid].md5_ctx = md5_ctx;
+}
+
+__kernel void m19000_comp (KERN_ATTR_TMPS (qnx_md5_tmp_t))
+{
+  /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  md5_ctx_t md5_ctx = tmps[gid].md5_ctx;
+
+  md5_final (&md5_ctx);
+
+  const u32 r0 = md5_ctx.h[0];
+  const u32 r1 = md5_ctx.h[1];
+  const u32 r2 = md5_ctx.h[2];
+  const u32 r3 = md5_ctx.h[3];
+
+  #define il_pos 0
+
+  #include COMPARE_M
+}
diff --git a/OpenCL/m19100-pure.cl b/OpenCL/m19100-pure.cl
new file mode 100644
index 000000000..cdb4dff85
--- /dev/null
+++ b/OpenCL/m19100-pure.cl
@@ -0,0 +1,93 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+//#define NEW_SIMD_CODE
+
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
+#include "inc_simd.cl"
+#include "inc_hash_sha256.cl"
+
+#define COMPARE_S "inc_comp_single.cl"
+#define COMPARE_M "inc_comp_multi.cl"
+
+typedef struct qnx_sha256_tmp
+{
+  sha256_ctx_t sha256_ctx;
+
+} qnx_sha256_tmp_t;
+
+__kernel void m19100_init (KERN_ATTR_TMPS (qnx_sha256_tmp_t))
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * init
+   */
+
+  sha256_ctx_t sha256_ctx;
+
+  sha256_init (&sha256_ctx);
+
+  sha256_update_global_swap (&sha256_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
+
+  sha256_update_global_swap (&sha256_ctx, pws[gid].i, pws[gid].pw_len);
+
+  tmps[gid].sha256_ctx = sha256_ctx;
+}
+
+__kernel void m19100_loop (KERN_ATTR_TMPS (qnx_sha256_tmp_t))
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  sha256_ctx_t sha256_ctx = tmps[gid].sha256_ctx;
+
+  for (u32 i = 0; i < loop_cnt; i++)
+  {
+    sha256_update_global_swap (&sha256_ctx, pws[gid].i, pws[gid].pw_len);
+  }
+
+  tmps[gid].sha256_ctx = sha256_ctx;
+}
+
+__kernel void m19100_comp (KERN_ATTR_TMPS (qnx_sha256_tmp_t))
+{
+  /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  sha256_ctx_t sha256_ctx = tmps[gid].sha256_ctx;
+
+  sha256_final (&sha256_ctx);
+
+  const u32 r0 = swap32_S (sha256_ctx.h[0]);
+  const u32 r1 = swap32_S (sha256_ctx.h[1]);
+  const u32 r2 = swap32_S (sha256_ctx.h[2]);
+  const u32 r3 = swap32_S (sha256_ctx.h[3]);
+
+  #define il_pos 0
+
+  #include COMPARE_M
+}
diff --git a/OpenCL/m19200-pure.cl b/OpenCL/m19200-pure.cl
new file mode 100644
index 000000000..75250b009
--- /dev/null
+++ b/OpenCL/m19200-pure.cl
@@ -0,0 +1,455 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+//#define NEW_SIMD_CODE
+
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
+#include "inc_simd.cl"
+#include "inc_hash_sha512.cl"
+
+#define COMPARE_S "inc_comp_single.cl"
+#define COMPARE_M "inc_comp_multi.cl"
+
+typedef struct qnx_sha512_tmp
+{
+  sha512_ctx_t sha512_ctx;
+
+  u32 sav; // to trigger sha512 bug
+
+} qnx_sha512_tmp_t;
+
+DECLSPEC u32 sha512_update_128_qnxbug (sha512_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, u32 *w4, u32 *w5, u32 *w6, u32 *w7, const int len, u32 sav)
+{
+  MAYBE_VOLATILE const int pos = ctx->len & 127;
+
+  ctx->len += len;
+
+  if ((pos + len) < 128)
+  {
+    switch_buffer_by_offset_8x4_be_S (w0, w1, w2, w3, w4, w5, w6, w7, pos);
+
+    ctx->w0[0] |= w0[0];
+    ctx->w0[1] |= w0[1];
+    ctx->w0[2] |= w0[2];
+    ctx->w0[3] |= w0[3];
+    ctx->w1[0] |= w1[0];
+    ctx->w1[1] |= w1[1];
+    ctx->w1[2] |= w1[2];
+    ctx->w1[3] |= w1[3];
+    ctx->w2[0] |= w2[0];
+    ctx->w2[1] |= w2[1];
+    ctx->w2[2] |= w2[2];
+    ctx->w2[3] |= w2[3];
+    ctx->w3[0] |= w3[0];
+    ctx->w3[1] |= w3[1];
+    ctx->w3[2] |= w3[2];
+    ctx->w3[3] |= w3[3];
+    ctx->w4[0] |= w4[0];
+    ctx->w4[1] |= w4[1];
+    ctx->w4[2] |= w4[2];
+    ctx->w4[3] |= w4[3];
+    ctx->w5[0] |= w5[0];
+    ctx->w5[1] |= w5[1];
+    ctx->w5[2] |= w5[2];
+    ctx->w5[3] |= w5[3];
+    ctx->w6[0] |= w6[0];
+    ctx->w6[1] |= w6[1];
+    ctx->w6[2] |= w6[2];
+    ctx->w6[3] |= w6[3];
+    ctx->w7[0] |= w7[0];
+    ctx->w7[1] |= w7[1];
+    ctx->w7[2] |= w7[2];
+    ctx->w7[3] |= w7[3];
+  }
+  else
+  {
+    u32 c0[4] = { 0 };
+    u32 c1[4] = { 0 };
+    u32 c2[4] = { 0 };
+    u32 c3[4] = { 0 };
+    u32 c4[4] = { 0 };
+    u32 c5[4] = { 0 };
+    u32 c6[4] = { 0 };
+    u32 c7[4] = { 0 };
+
+    switch_buffer_by_offset_8x4_carry_be_S (w0, w1, w2, w3, w4, w5, w6, w7, c0, c1, c2, c3, c4, c5, c6, c7, pos);
+
+    ctx->w0[0] |= w0[0];
+    ctx->w0[1] |= w0[1];
+    ctx->w0[2] |= w0[2];
+    ctx->w0[3] |= w0[3];
+    ctx->w1[0] |= w1[0];
+    ctx->w1[1] |= w1[1];
+    ctx->w1[2] |= w1[2];
+    ctx->w1[3] |= w1[3];
+    ctx->w2[0] |= w2[0];
+    ctx->w2[1] |= w2[1];
+    ctx->w2[2] |= w2[2];
+    ctx->w2[3] |= w2[3];
+    ctx->w3[0] |= w3[0];
+    ctx->w3[1] |= w3[1];
+    ctx->w3[2] |= w3[2];
+    ctx->w3[3] |= w3[3];
+    ctx->w4[0] |= w4[0];
+    ctx->w4[1] |= w4[1];
+    ctx->w4[2] |= w4[2];
+    ctx->w4[3] |= w4[3];
+    ctx->w5[0] |= w5[0];
+    ctx->w5[1] |= w5[1];
+    ctx->w5[2] |= w5[2];
+    ctx->w5[3] |= w5[3];
+    ctx->w6[0] |= w6[0];
+    ctx->w6[1] |= w6[1];
+    ctx->w6[2] |= w6[2];
+    ctx->w6[3] |= w6[3];
+    ctx->w7[0] |= w7[0];
+    ctx->w7[1] |= w7[1];
+    ctx->w7[2] |= w7[2];
+    ctx->w7[3] |= w7[3];
+
+    sav = ctx->w7[1];
+
+    sha512_transform (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, ctx->h);
+
+    ctx->w0[0] = c0[0];
+    ctx->w0[1] = c0[1];
+    ctx->w0[2] = c0[2];
+    ctx->w0[3] = c0[3];
+    ctx->w1[0] = c1[0];
+    ctx->w1[1] = c1[1];
+    ctx->w1[2] = c1[2];
+    ctx->w1[3] = c1[3];
+    ctx->w2[0] = c2[0];
+    ctx->w2[1] = c2[1];
+    ctx->w2[2] = c2[2];
+    ctx->w2[3] = c2[3];
+    ctx->w3[0] = c3[0];
+    ctx->w3[1] = c3[1];
+    ctx->w3[2] = c3[2];
+    ctx->w3[3] = c3[3];
+    ctx->w4[0] = c4[0];
+    ctx->w4[1] = c4[1];
+    ctx->w4[2] = c4[2];
+    ctx->w4[3] = c4[3];
+    ctx->w5[0] = c5[0];
+    ctx->w5[1] = c5[1];
+    ctx->w5[2] = c5[2];
+    ctx->w5[3] = c5[3];
+    ctx->w6[0] = c6[0];
+    ctx->w6[1] = c6[1];
+    ctx->w6[2] = c6[2];
+    ctx->w6[3] = c6[3];
+    ctx->w7[0] = c7[0];
+    ctx->w7[1] = c7[1];
+    ctx->w7[2] = c7[2];
+    ctx->w7[3] = c7[3];
+  }
+
+  return sav;
+}
+
+DECLSPEC u32 sha512_update_global_swap_qnxbug (sha512_ctx_t *ctx, const __global u32 *w, const int len, u32 sav)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  int pos1;
+  int pos4;
+
+  for (pos1 = 0, pos4 = 0; pos1 < len - 128; pos1 += 128, pos4 += 32)
+  {
+    w0[0] = w[pos4 +  0];
+    w0[1] = w[pos4 +  1];
+    w0[2] = w[pos4 +  2];
+    w0[3] = w[pos4 +  3];
+    w1[0] = w[pos4 +  4];
+    w1[1] = w[pos4 +  5];
+    w1[2] = w[pos4 +  6];
+    w1[3] = w[pos4 +  7];
+    w2[0] = w[pos4 +  8];
+    w2[1] = w[pos4 +  9];
+    w2[2] = w[pos4 + 10];
+    w2[3] = w[pos4 + 11];
+    w3[0] = w[pos4 + 12];
+    w3[1] = w[pos4 + 13];
+    w3[2] = w[pos4 + 14];
+    w3[3] = w[pos4 + 15];
+    w4[0] = w[pos4 + 16];
+    w4[1] = w[pos4 + 17];
+    w4[2] = w[pos4 + 18];
+    w4[3] = w[pos4 + 19];
+    w5[0] = w[pos4 + 20];
+    w5[1] = w[pos4 + 21];
+    w5[2] = w[pos4 + 22];
+    w5[3] = w[pos4 + 23];
+    w6[0] = w[pos4 + 24];
+    w6[1] = w[pos4 + 25];
+    w6[2] = w[pos4 + 26];
+    w6[3] = w[pos4 + 27];
+    w7[0] = w[pos4 + 28];
+    w7[1] = w[pos4 + 29];
+    w7[2] = w[pos4 + 30];
+    w7[3] = w[pos4 + 31];
+
+    w0[0] = swap32_S (w0[0]);
+    w0[1] = swap32_S (w0[1]);
+    w0[2] = swap32_S (w0[2]);
+    w0[3] = swap32_S (w0[3]);
+    w1[0] = swap32_S (w1[0]);
+    w1[1] = swap32_S (w1[1]);
+    w1[2] = swap32_S (w1[2]);
+    w1[3] = swap32_S (w1[3]);
+    w2[0] = swap32_S (w2[0]);
+    w2[1] = swap32_S (w2[1]);
+    w2[2] = swap32_S (w2[2]);
+    w2[3] = swap32_S (w2[3]);
+    w3[0] = swap32_S (w3[0]);
+    w3[1] = swap32_S (w3[1]);
+    w3[2] = swap32_S (w3[2]);
+    w3[3] = swap32_S (w3[3]);
+    w4[0] = swap32_S (w4[0]);
+    w4[1] = swap32_S (w4[1]);
+    w4[2] = swap32_S (w4[2]);
+    w4[3] = swap32_S (w4[3]);
+    w5[0] = swap32_S (w5[0]);
+    w5[1] = swap32_S (w5[1]);
+    w5[2] = swap32_S (w5[2]);
+    w5[3] = swap32_S (w5[3]);
+    w6[0] = swap32_S (w6[0]);
+    w6[1] = swap32_S (w6[1]);
+    w6[2] = swap32_S (w6[2]);
+    w6[3] = swap32_S (w6[3]);
+    w7[0] = swap32_S (w7[0]);
+    w7[1] = swap32_S (w7[1]);
+    w7[2] = swap32_S (w7[2]);
+    w7[3] = swap32_S (w7[3]);
+
+    sav = sha512_update_128_qnxbug (ctx, w0, w1, w2, w3, w4, w5, w6, w7, 128, sav);
+  }
+
+  w0[0] = w[pos4 +  0];
+  w0[1] = w[pos4 +  1];
+  w0[2] = w[pos4 +  2];
+  w0[3] = w[pos4 +  3];
+  w1[0] = w[pos4 +  4];
+  w1[1] = w[pos4 +  5];
+  w1[2] = w[pos4 +  6];
+  w1[3] = w[pos4 +  7];
+  w2[0] = w[pos4 +  8];
+  w2[1] = w[pos4 +  9];
+  w2[2] = w[pos4 + 10];
+  w2[3] = w[pos4 + 11];
+  w3[0] = w[pos4 + 12];
+  w3[1] = w[pos4 + 13];
+  w3[2] = w[pos4 + 14];
+  w3[3] = w[pos4 + 15];
+  w4[0] = w[pos4 + 16];
+  w4[1] = w[pos4 + 17];
+  w4[2] = w[pos4 + 18];
+  w4[3] = w[pos4 + 19];
+  w5[0] = w[pos4 + 20];
+  w5[1] = w[pos4 + 21];
+  w5[2] = w[pos4 + 22];
+  w5[3] = w[pos4 + 23];
+  w6[0] = w[pos4 + 24];
+  w6[1] = w[pos4 + 25];
+  w6[2] = w[pos4 + 26];
+  w6[3] = w[pos4 + 27];
+  w7[0] = w[pos4 + 28];
+  w7[1] = w[pos4 + 29];
+  w7[2] = w[pos4 + 30];
+  w7[3] = w[pos4 + 31];
+
+  w0[0] = swap32_S (w0[0]);
+  w0[1] = swap32_S (w0[1]);
+  w0[2] = swap32_S (w0[2]);
+  w0[3] = swap32_S (w0[3]);
+  w1[0] = swap32_S (w1[0]);
+  w1[1] = swap32_S (w1[1]);
+  w1[2] = swap32_S (w1[2]);
+  w1[3] = swap32_S (w1[3]);
+  w2[0] = swap32_S (w2[0]);
+  w2[1] = swap32_S (w2[1]);
+  w2[2] = swap32_S (w2[2]);
+  w2[3] = swap32_S (w2[3]);
+  w3[0] = swap32_S (w3[0]);
+  w3[1] = swap32_S (w3[1]);
+  w3[2] = swap32_S (w3[2]);
+  w3[3] = swap32_S (w3[3]);
+  w4[0] = swap32_S (w4[0]);
+  w4[1] = swap32_S (w4[1]);
+  w4[2] = swap32_S (w4[2]);
+  w4[3] = swap32_S (w4[3]);
+  w5[0] = swap32_S (w5[0]);
+  w5[1] = swap32_S (w5[1]);
+  w5[2] = swap32_S (w5[2]);
+  w5[3] = swap32_S (w5[3]);
+  w6[0] = swap32_S (w6[0]);
+  w6[1] = swap32_S (w6[1]);
+  w6[2] = swap32_S (w6[2]);
+  w6[3] = swap32_S (w6[3]);
+  w7[0] = swap32_S (w7[0]);
+  w7[1] = swap32_S (w7[1]);
+  w7[2] = swap32_S (w7[2]);
+  w7[3] = swap32_S (w7[3]);
+
+  sav = sha512_update_128_qnxbug (ctx, w0, w1, w2, w3, w4, w5, w6, w7, len - pos1, sav);
+
+  return sav;
+}
+
+DECLSPEC void sha512_final_qnxbug (sha512_ctx_t *ctx, u32 sav)
+{
+  MAYBE_VOLATILE const int pos = ctx->len & 127;
+
+  append_0x80_8x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, pos ^ 3);
+
+  if (pos >= 112)
+  {
+    sav = ctx->w7[1];
+
+    sha512_transform (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, ctx->h);
+
+    ctx->w0[0] = 0;
+    ctx->w0[1] = 0;
+    ctx->w0[2] = 0;
+    ctx->w0[3] = 0;
+    ctx->w1[0] = 0;
+    ctx->w1[1] = 0;
+    ctx->w1[2] = 0;
+    ctx->w1[3] = 0;
+    ctx->w2[0] = 0;
+    ctx->w2[1] = 0;
+    ctx->w2[2] = 0;
+    ctx->w2[3] = 0;
+    ctx->w3[0] = 0;
+    ctx->w3[1] = 0;
+    ctx->w3[2] = 0;
+    ctx->w3[3] = 0;
+    ctx->w4[0] = 0;
+    ctx->w4[1] = 0;
+    ctx->w4[2] = 0;
+    ctx->w4[3] = 0;
+    ctx->w5[0] = 0;
+    ctx->w5[1] = 0;
+    ctx->w5[2] = 0;
+    ctx->w5[3] = 0;
+    ctx->w6[0] = 0;
+    ctx->w6[1] = 0;
+    ctx->w6[2] = 0;
+    ctx->w6[3] = 0;
+    ctx->w7[0] = 0;
+    ctx->w7[1] = 0;
+    ctx->w7[2] = 0;
+    ctx->w7[3] = 0;
+  }
+
+  ctx->w7[1] = sav;
+  ctx->w7[2] = 0;
+  ctx->w7[3] = ctx->len * 8;
+
+  sha512_transform (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, ctx->h);
+}
+
+__kernel void m19200_init (KERN_ATTR_TMPS (qnx_sha512_tmp_t))
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * init
+   */
+
+  sha512_ctx_t sha512_ctx;
+
+  sha512_init (&sha512_ctx);
+
+  sha512_update_global_swap (&sha512_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
+
+  sha512_update_global_swap (&sha512_ctx, pws[gid].i, pws[gid].pw_len);
+
+  tmps[gid].sha512_ctx = sha512_ctx;
+  tmps[gid].sav = 0;
+}
+
+__kernel void m19200_loop (KERN_ATTR_TMPS (qnx_sha512_tmp_t))
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  sha512_ctx_t sha512_ctx = tmps[gid].sha512_ctx;
+  u32 sav = tmps[gid].sav;
+
+  for (u32 i = 0; i < loop_cnt; i++)
+  {
+    sav = sha512_update_global_swap_qnxbug (&sha512_ctx, pws[gid].i, pws[gid].pw_len, sav);
+  }
+
+  tmps[gid].sha512_ctx = sha512_ctx;
+  tmps[gid].sav = sav;
+}
+
+__kernel void m19200_comp (KERN_ATTR_TMPS (qnx_sha512_tmp_t))
+{
+  /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  sha512_ctx_t sha512_ctx = tmps[gid].sha512_ctx;
+
+  sha512_final (&sha512_ctx);
+
+  const u32 r0 = l32_from_64_S (swap64_S (sha512_ctx.h[0]));
+  const u32 r1 = h32_from_64_S (swap64_S (sha512_ctx.h[0]));
+  const u32 r2 = l32_from_64_S (swap64_S (sha512_ctx.h[1]));
+  const u32 r3 = h32_from_64_S (swap64_S (sha512_ctx.h[1]));
+
+  #define il_pos 0
+
+  #include COMPARE_M
+
+  // we should also handle the buggy qnx sha512 implementation
+  // see https://github.com/magnumripper/JohnTheRipper/blob/bleeding-jumbo/src/sha2.c#L578-L595
+
+  sha512_ctx_t sha512_ctx2 = tmps[gid].sha512_ctx;
+  u32 sav = tmps[gid].sav;
+
+  if (sha512_ctx2.len >= 116)
+  {
+    sha512_final_qnxbug (&sha512_ctx2, sav);
+
+    const u32 r0 = l32_from_64_S (swap64_S (sha512_ctx2.h[0]));
+    const u32 r1 = h32_from_64_S (swap64_S (sha512_ctx2.h[0]));
+    const u32 r2 = l32_from_64_S (swap64_S (sha512_ctx2.h[1]));
+    const u32 r3 = h32_from_64_S (swap64_S (sha512_ctx2.h[1]));
+
+    #include COMPARE_M
+  }
+}
diff --git a/docs/changes.txt b/docs/changes.txt
index ee040c275..579b157a1 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -10,6 +10,9 @@
 - Added hash-mode 18700 Java Object hashCode()
 - Added hash-mode 18800 Blockchain, My Wallet, Second Password (SHA256)
 - Added hash-mode 18900 Android Backup
+- Added hash-mode 19000 QNX /etc/shadow (MD5)
+- Added hash-mode 19100 QNX /etc/shadow (SHA256)
+- Added hash-mode 19200 QNX /etc/shadow (SHA512)
 
 ##
 ## Bugs
diff --git a/src/modules/module_19000.c b/src/modules/module_19000.c
new file mode 100644
index 000000000..935c995f0
--- /dev/null
+++ b/src/modules/module_19000.c
@@ -0,0 +1,257 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+#include "common.h"
+#include "types.h"
+#include "modules.h"
+#include "bitops.h"
+#include "convert.h"
+#include "shared.h"
+#include "cpu_md5.h"
+#include "memory.h"
+
+static const u32   ATTACK_EXEC    = ATTACK_EXEC_OUTSIDE_KERNEL;
+static const u32   DGST_POS0      = 0;
+static const u32   DGST_POS1      = 1;
+static const u32   DGST_POS2      = 2;
+static const u32   DGST_POS3      = 3;
+static const u32   DGST_SIZE      = DGST_SIZE_8_8;
+static const u32   HASH_CATEGORY  = HASH_CATEGORY_OS;
+static const char *HASH_NAME      = "QNX /etc/shadow (MD5)";
+static const u64   KERN_TYPE      = 19000;
+static const u32   OPTI_TYPE      = OPTI_TYPE_ZERO_BYTE;
+static const u64   OPTS_TYPE      = OPTS_TYPE_PT_GENERATE_LE;
+static const u32   SALT_TYPE      = SALT_TYPE_EMBEDDED;
+static const char *ST_PASS        = "hashcat";
+static const char *ST_HASH        = "@m@75f6f129f9c9e77b6b1b78f791ed764a@8741857532330050";
+
+u32         module_attack_exec    (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC;     }
+u32         module_dgst_pos0      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0;       }
+u32         module_dgst_pos1      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS1;       }
+u32         module_dgst_pos2      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS2;       }
+u32         module_dgst_pos3      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS3;       }
+u32         module_dgst_size      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_SIZE;       }
+u32         module_hash_category  (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_CATEGORY;   }
+const char *module_hash_name      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_NAME;       }
+u64         module_kern_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return KERN_TYPE;       }
+u32         module_opti_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTI_TYPE;       }
+u64         module_opts_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTS_TYPE;       }
+u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return SALT_TYPE;       }
+const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
+const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
+
+typedef struct md5_ctx
+{
+  u32 h[4];
+
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  int len;
+
+} md5_ctx_t;
+
+typedef struct qnx_md5_tmp
+{
+  md5_ctx_t md5_ctx;
+
+} qnx_md5_tmp_t;
+
+static const int ROUNDS_QNX = 1000;
+
+u64 module_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
+{
+  const u64 tmp_size = (const u64) sizeof (qnx_md5_tmp_t);
+
+  return tmp_size;
+}
+
+int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
+{
+  u32 *digest = (u32 *) digest_buf;
+
+  token_t token;
+
+  token.token_cnt  = 4;
+
+  token.sep[0]     = '@';
+  token.len_min[0] = 0;
+  token.len_max[0] = 0;
+  token.attr[0]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  token.sep[1]     = '@';
+  token.len_min[1] = 1;
+  token.len_max[1] = 8;
+  token.attr[1]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  token.sep[2]     = '@';
+  token.len_min[2] = 32;
+  token.len_max[2] = 128;
+  token.attr[2]    = TOKEN_ATTR_VERIFY_LENGTH
+                   | TOKEN_ATTR_VERIFY_HEX;
+
+  token.sep[3]     = '@';
+  token.len_min[3] = 8;
+  token.len_max[3] = 16;
+  token.attr[3]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token);
+
+  if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
+
+  // check hash type
+
+  if (token.buf[1][0] != 'm') return (PARSER_SIGNATURE_UNMATCHED);
+
+  // check iter
+
+  u32 iter = ROUNDS_QNX;
+
+  if (token.len[1] > 1)
+  {
+    if (token.buf[1][1] != ',') return (PARSER_SEPARATOR_UNMATCHED);
+
+    iter = hc_strtoul ((const char *) token.buf[1] + 2, NULL, 10);
+  }
+
+  // iter++; the additinal round is added in the init kernel
+
+  salt->salt_iter = iter;
+
+  // digest
+
+  if (token.len[2] != 32) return (PARSER_HASH_LENGTH);
+
+  digest[0] = hex_to_u32 ((const u8 *) token.buf[2] +  0);
+  digest[1] = hex_to_u32 ((const u8 *) token.buf[2] +  8);
+  digest[2] = hex_to_u32 ((const u8 *) token.buf[2] + 16);
+  digest[3] = hex_to_u32 ((const u8 *) token.buf[2] + 24);
+
+  // salt
+
+  if ((token.len[3] == 8) || (token.len[3] == 16))
+  {
+    memcpy (salt->salt_buf, token.buf[3], token.len[3]);
+
+    salt->salt_len = token.len[3];
+  }
+  else
+  {
+    return (PARSER_SALT_LENGTH);
+  }
+
+  return (PARSER_OK);
+}
+
+int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size)
+{
+  const u32 *digest = (const u32 *) digest_buf;
+
+  int line_len = 0;
+
+  const int iter = salt->salt_iter;
+
+  if (iter == ROUNDS_QNX)
+  {
+    line_buf[line_len++] = '@';
+    line_buf[line_len++] = 'm';
+    line_buf[line_len++] = '@';
+  }
+  else
+  {
+    line_buf[line_len++] = '@';
+    line_buf[line_len++] = 'm';
+    line_buf[line_len++] = ',';
+
+    line_len += snprintf (line_buf + line_len, line_size - line_len, "%u", iter);
+
+    line_buf[line_len++] = '@';
+  }
+
+  u32_to_hex (digest[0], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[1], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[2], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[3], (u8 *) line_buf + line_len); line_len += 8;
+
+  line_buf[line_len++] = '@';
+
+  memcpy (line_buf + line_len, salt->salt_buf, salt->salt_len);
+
+  line_len += salt->salt_len;
+
+  line_buf[line_len] = 0;
+
+  return line_len;
+}
+
+void module_init (module_ctx_t *module_ctx)
+{
+  module_ctx->module_context_size             = MODULE_CONTEXT_SIZE_CURRENT;
+  module_ctx->module_interface_version        = MODULE_INTERFACE_VERSION_CURRENT;
+
+  module_ctx->module_attack_exec              = module_attack_exec;
+  module_ctx->module_benchmark_esalt          = MODULE_DEFAULT;
+  module_ctx->module_benchmark_hook_salt      = MODULE_DEFAULT;
+  module_ctx->module_benchmark_mask           = MODULE_DEFAULT;
+  module_ctx->module_benchmark_salt           = MODULE_DEFAULT;
+  module_ctx->module_build_plain_postprocess  = MODULE_DEFAULT;
+  module_ctx->module_deep_comp_kernel         = MODULE_DEFAULT;
+  module_ctx->module_dgst_pos0                = module_dgst_pos0;
+  module_ctx->module_dgst_pos1                = module_dgst_pos1;
+  module_ctx->module_dgst_pos2                = module_dgst_pos2;
+  module_ctx->module_dgst_pos3                = module_dgst_pos3;
+  module_ctx->module_dgst_size                = module_dgst_size;
+  module_ctx->module_dictstat_disable         = MODULE_DEFAULT;
+  module_ctx->module_esalt_size               = MODULE_DEFAULT;
+  module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
+  module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_save         = MODULE_DEFAULT;
+  module_ctx->module_hash_decode_outfile      = MODULE_DEFAULT;
+  module_ctx->module_hash_decode_zero_hash    = MODULE_DEFAULT;
+  module_ctx->module_hash_decode              = module_hash_decode;
+  module_ctx->module_hash_encode_status       = MODULE_DEFAULT;
+  module_ctx->module_hash_encode              = module_hash_encode;
+  module_ctx->module_hash_init_selftest       = MODULE_DEFAULT;
+  module_ctx->module_hash_mode                = MODULE_DEFAULT;
+  module_ctx->module_hash_category            = module_hash_category;
+  module_ctx->module_hash_name                = module_hash_name;
+  module_ctx->module_hlfmt_disable            = MODULE_DEFAULT;
+  module_ctx->module_hook12                   = MODULE_DEFAULT;
+  module_ctx->module_hook23                   = MODULE_DEFAULT;
+  module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
+  module_ctx->module_hook_size                = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
+  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_min       = MODULE_DEFAULT;
+  module_ctx->module_kern_type                = module_kern_type;
+  module_ctx->module_kern_type_dynamic        = MODULE_DEFAULT;
+  module_ctx->module_opti_type                = module_opti_type;
+  module_ctx->module_opts_type                = module_opts_type;
+  module_ctx->module_outfile_check_disable    = MODULE_DEFAULT;
+  module_ctx->module_outfile_check_nocomp     = MODULE_DEFAULT;
+  module_ctx->module_potfile_disable          = MODULE_DEFAULT;
+  module_ctx->module_potfile_keep_all_hashes  = MODULE_DEFAULT;
+  module_ctx->module_pwdump_column            = MODULE_DEFAULT;
+  module_ctx->module_pw_max                   = MODULE_DEFAULT;
+  module_ctx->module_pw_min                   = MODULE_DEFAULT;
+  module_ctx->module_salt_max                 = MODULE_DEFAULT;
+  module_ctx->module_salt_min                 = MODULE_DEFAULT;
+  module_ctx->module_salt_type                = module_salt_type;
+  module_ctx->module_separator                = MODULE_DEFAULT;
+  module_ctx->module_st_hash                  = module_st_hash;
+  module_ctx->module_st_pass                  = module_st_pass;
+  module_ctx->module_tmp_size                 = module_tmp_size;
+  module_ctx->module_unstable_warning         = MODULE_DEFAULT;
+  module_ctx->module_warmup_disable           = MODULE_DEFAULT;
+}
diff --git a/src/modules/module_19100.c b/src/modules/module_19100.c
new file mode 100644
index 000000000..9c9085c87
--- /dev/null
+++ b/src/modules/module_19100.c
@@ -0,0 +1,265 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+#include "common.h"
+#include "types.h"
+#include "modules.h"
+#include "bitops.h"
+#include "convert.h"
+#include "shared.h"
+#include "cpu_md5.h"
+#include "memory.h"
+
+static const u32   ATTACK_EXEC    = ATTACK_EXEC_OUTSIDE_KERNEL;
+static const u32   DGST_POS0      = 0;
+static const u32   DGST_POS1      = 1;
+static const u32   DGST_POS2      = 2;
+static const u32   DGST_POS3      = 3;
+static const u32   DGST_SIZE      = DGST_SIZE_8_8;
+static const u32   HASH_CATEGORY  = HASH_CATEGORY_OS;
+static const char *HASH_NAME      = "QNX /etc/shadow (SHA256)";
+static const u64   KERN_TYPE      = 19100;
+static const u32   OPTI_TYPE      = OPTI_TYPE_ZERO_BYTE;
+static const u64   OPTS_TYPE      = OPTS_TYPE_PT_GENERATE_LE;
+static const u32   SALT_TYPE      = SALT_TYPE_EMBEDDED;
+static const char *ST_PASS        = "hashcat";
+static const char *ST_HASH        = "@s@0b365cab7e17ee1e7e1a90078501cc1aa85888d6da34e2f5b04f5c614b882a93@5498317092471604";
+
+u32         module_attack_exec    (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC;     }
+u32         module_dgst_pos0      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0;       }
+u32         module_dgst_pos1      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS1;       }
+u32         module_dgst_pos2      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS2;       }
+u32         module_dgst_pos3      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS3;       }
+u32         module_dgst_size      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_SIZE;       }
+u32         module_hash_category  (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_CATEGORY;   }
+const char *module_hash_name      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_NAME;       }
+u64         module_kern_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return KERN_TYPE;       }
+u32         module_opti_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTI_TYPE;       }
+u64         module_opts_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTS_TYPE;       }
+u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return SALT_TYPE;       }
+const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
+const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
+
+typedef struct sha256_ctx
+{
+  u32 h[8];
+
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  int len;
+
+} sha256_ctx_t;
+
+typedef struct qnx_sha256_tmp
+{
+  sha256_ctx_t sha256_ctx;
+
+} qnx_sha256_tmp_t;
+
+static const int ROUNDS_QNX = 1000;
+
+u64 module_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
+{
+  const u64 tmp_size = (const u64) sizeof (qnx_sha256_tmp_t);
+
+  return tmp_size;
+}
+
+int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
+{
+  u32 *digest = (u32 *) digest_buf;
+
+  token_t token;
+
+  token.token_cnt  = 4;
+
+  token.sep[0]     = '@';
+  token.len_min[0] = 0;
+  token.len_max[0] = 0;
+  token.attr[0]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  token.sep[1]     = '@';
+  token.len_min[1] = 1;
+  token.len_max[1] = 8;
+  token.attr[1]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  token.sep[2]     = '@';
+  token.len_min[2] = 32;
+  token.len_max[2] = 128;
+  token.attr[2]    = TOKEN_ATTR_VERIFY_LENGTH
+                   | TOKEN_ATTR_VERIFY_HEX;
+
+  token.sep[3]     = '@';
+  token.len_min[3] = 8;
+  token.len_max[3] = 16;
+  token.attr[3]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token);
+
+  if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
+
+  // check hash type
+
+  if (token.buf[1][0] != 's') return (PARSER_SIGNATURE_UNMATCHED);
+
+  // check iter
+
+  u32 iter = ROUNDS_QNX;
+
+  if (token.len[1] > 1)
+  {
+    if (token.buf[1][1] != ',') return (PARSER_SEPARATOR_UNMATCHED);
+
+    iter = hc_strtoul ((const char *) token.buf[1] + 2, NULL, 10);
+  }
+
+  // iter++; the additinal round is added in the init kernel
+
+  salt->salt_iter = iter;
+
+  // digest
+
+  if (token.len[2] != 64) return (PARSER_HASH_LENGTH);
+
+  digest[0] = hex_to_u32 ((const u8 *) token.buf[2] +  0);
+  digest[1] = hex_to_u32 ((const u8 *) token.buf[2] +  8);
+  digest[2] = hex_to_u32 ((const u8 *) token.buf[2] + 16);
+  digest[3] = hex_to_u32 ((const u8 *) token.buf[2] + 24);
+  digest[4] = hex_to_u32 ((const u8 *) token.buf[2] + 32);
+  digest[5] = hex_to_u32 ((const u8 *) token.buf[2] + 40);
+  digest[6] = hex_to_u32 ((const u8 *) token.buf[2] + 48);
+  digest[7] = hex_to_u32 ((const u8 *) token.buf[2] + 56);
+
+  // salt
+
+  if ((token.len[3] == 8) || (token.len[3] == 16))
+  {
+    memcpy (salt->salt_buf, token.buf[3], token.len[3]);
+
+    salt->salt_len = token.len[3];
+  }
+  else
+  {
+    return (PARSER_SALT_LENGTH);
+  }
+
+  return (PARSER_OK);
+}
+
+int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size)
+{
+  const u32 *digest = (const u32 *) digest_buf;
+
+  int line_len = 0;
+
+  const int iter = salt->salt_iter;
+
+  if (iter == ROUNDS_QNX)
+  {
+    line_buf[line_len++] = '@';
+    line_buf[line_len++] = 's';
+    line_buf[line_len++] = '@';
+  }
+  else
+  {
+    line_buf[line_len++] = '@';
+    line_buf[line_len++] = 's';
+    line_buf[line_len++] = ',';
+
+    line_len += snprintf (line_buf + line_len, line_size - line_len, "%u", iter);
+
+    line_buf[line_len++] = '@';
+  }
+
+  u32_to_hex (digest[0], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[1], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[2], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[3], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[4], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[5], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[6], (u8 *) line_buf + line_len); line_len += 8;
+  u32_to_hex (digest[7], (u8 *) line_buf + line_len); line_len += 8;
+
+  line_buf[line_len++] = '@';
+
+  memcpy (line_buf + line_len, salt->salt_buf, salt->salt_len);
+
+  line_len += salt->salt_len;
+
+  line_buf[line_len] = 0;
+
+  return line_len;
+}
+
+void module_init (module_ctx_t *module_ctx)
+{
+  module_ctx->module_context_size             = MODULE_CONTEXT_SIZE_CURRENT;
+  module_ctx->module_interface_version        = MODULE_INTERFACE_VERSION_CURRENT;
+
+  module_ctx->module_attack_exec              = module_attack_exec;
+  module_ctx->module_benchmark_esalt          = MODULE_DEFAULT;
+  module_ctx->module_benchmark_hook_salt      = MODULE_DEFAULT;
+  module_ctx->module_benchmark_mask           = MODULE_DEFAULT;
+  module_ctx->module_benchmark_salt           = MODULE_DEFAULT;
+  module_ctx->module_build_plain_postprocess  = MODULE_DEFAULT;
+  module_ctx->module_deep_comp_kernel         = MODULE_DEFAULT;
+  module_ctx->module_dgst_pos0                = module_dgst_pos0;
+  module_ctx->module_dgst_pos1                = module_dgst_pos1;
+  module_ctx->module_dgst_pos2                = module_dgst_pos2;
+  module_ctx->module_dgst_pos3                = module_dgst_pos3;
+  module_ctx->module_dgst_size                = module_dgst_size;
+  module_ctx->module_dictstat_disable         = MODULE_DEFAULT;
+  module_ctx->module_esalt_size               = MODULE_DEFAULT;
+  module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
+  module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_save         = MODULE_DEFAULT;
+  module_ctx->module_hash_decode_outfile      = MODULE_DEFAULT;
+  module_ctx->module_hash_decode_zero_hash    = MODULE_DEFAULT;
+  module_ctx->module_hash_decode              = module_hash_decode;
+  module_ctx->module_hash_encode_status       = MODULE_DEFAULT;
+  module_ctx->module_hash_encode              = module_hash_encode;
+  module_ctx->module_hash_init_selftest       = MODULE_DEFAULT;
+  module_ctx->module_hash_mode                = MODULE_DEFAULT;
+  module_ctx->module_hash_category            = module_hash_category;
+  module_ctx->module_hash_name                = module_hash_name;
+  module_ctx->module_hlfmt_disable            = MODULE_DEFAULT;
+  module_ctx->module_hook12                   = MODULE_DEFAULT;
+  module_ctx->module_hook23                   = MODULE_DEFAULT;
+  module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
+  module_ctx->module_hook_size                = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
+  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_min       = MODULE_DEFAULT;
+  module_ctx->module_kern_type                = module_kern_type;
+  module_ctx->module_kern_type_dynamic        = MODULE_DEFAULT;
+  module_ctx->module_opti_type                = module_opti_type;
+  module_ctx->module_opts_type                = module_opts_type;
+  module_ctx->module_outfile_check_disable    = MODULE_DEFAULT;
+  module_ctx->module_outfile_check_nocomp     = MODULE_DEFAULT;
+  module_ctx->module_potfile_disable          = MODULE_DEFAULT;
+  module_ctx->module_potfile_keep_all_hashes  = MODULE_DEFAULT;
+  module_ctx->module_pwdump_column            = MODULE_DEFAULT;
+  module_ctx->module_pw_max                   = MODULE_DEFAULT;
+  module_ctx->module_pw_min                   = MODULE_DEFAULT;
+  module_ctx->module_salt_max                 = MODULE_DEFAULT;
+  module_ctx->module_salt_min                 = MODULE_DEFAULT;
+  module_ctx->module_salt_type                = module_salt_type;
+  module_ctx->module_separator                = MODULE_DEFAULT;
+  module_ctx->module_st_hash                  = module_st_hash;
+  module_ctx->module_st_pass                  = module_st_pass;
+  module_ctx->module_tmp_size                 = module_tmp_size;
+  module_ctx->module_unstable_warning         = MODULE_DEFAULT;
+  module_ctx->module_warmup_disable           = MODULE_DEFAULT;
+}
diff --git a/src/modules/module_19200.c b/src/modules/module_19200.c
new file mode 100644
index 000000000..17be0ba32
--- /dev/null
+++ b/src/modules/module_19200.c
@@ -0,0 +1,271 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+#include "common.h"
+#include "types.h"
+#include "modules.h"
+#include "bitops.h"
+#include "convert.h"
+#include "shared.h"
+#include "cpu_md5.h"
+#include "memory.h"
+
+static const u32   ATTACK_EXEC    = ATTACK_EXEC_OUTSIDE_KERNEL;
+static const u32   DGST_POS0      = 0;
+static const u32   DGST_POS1      = 1;
+static const u32   DGST_POS2      = 2;
+static const u32   DGST_POS3      = 3;
+static const u32   DGST_SIZE      = DGST_SIZE_8_8;
+static const u32   HASH_CATEGORY  = HASH_CATEGORY_OS;
+static const char *HASH_NAME      = "QNX /etc/shadow (SHA512)";
+static const u64   KERN_TYPE      = 19200;
+static const u32   OPTI_TYPE      = OPTI_TYPE_ZERO_BYTE;
+static const u64   OPTS_TYPE      = OPTS_TYPE_PT_GENERATE_LE;
+static const u32   SALT_TYPE      = SALT_TYPE_EMBEDDED;
+static const char *ST_PASS        = "hashcat";
+static const char *ST_HASH        = "@S@715df9e94c097805dd1e13c6a40f331d02ce589765a2100ec7435e76b978d5efc364ce10870780622cee003c9951bd92ec1020c924b124cfff7e0fa1f73e3672@2257314490293159";
+
+u32         module_attack_exec    (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC;     }
+u32         module_dgst_pos0      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0;       }
+u32         module_dgst_pos1      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS1;       }
+u32         module_dgst_pos2      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS2;       }
+u32         module_dgst_pos3      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS3;       }
+u32         module_dgst_size      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_SIZE;       }
+u32         module_hash_category  (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_CATEGORY;   }
+const char *module_hash_name      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_NAME;       }
+u64         module_kern_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return KERN_TYPE;       }
+u32         module_opti_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTI_TYPE;       }
+u64         module_opts_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTS_TYPE;       }
+u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return SALT_TYPE;       }
+const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
+const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
+
+typedef struct sha512_ctx
+{
+  u64 h[8];
+
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  int len;
+
+} sha512_ctx_t;
+
+typedef struct qnx_sha512_tmp
+{
+  sha512_ctx_t sha512_ctx;
+
+  u32 sav; // to trigger sha512 bug
+
+} qnx_sha512_tmp_t;
+
+static const int ROUNDS_QNX = 1000;
+
+u64 module_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
+{
+  const u64 tmp_size = (const u64) sizeof (qnx_sha512_tmp_t);
+
+  return tmp_size;
+}
+
+int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
+{
+  u64 *digest = (u64 *) digest_buf;
+
+  token_t token;
+
+  token.token_cnt  = 4;
+
+  token.sep[0]     = '@';
+  token.len_min[0] = 0;
+  token.len_max[0] = 0;
+  token.attr[0]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  token.sep[1]     = '@';
+  token.len_min[1] = 1;
+  token.len_max[1] = 8;
+  token.attr[1]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  token.sep[2]     = '@';
+  token.len_min[2] = 32;
+  token.len_max[2] = 128;
+  token.attr[2]    = TOKEN_ATTR_VERIFY_LENGTH
+                   | TOKEN_ATTR_VERIFY_HEX;
+
+  token.sep[3]     = '@';
+  token.len_min[3] = 8;
+  token.len_max[3] = 16;
+  token.attr[3]    = TOKEN_ATTR_VERIFY_LENGTH;
+
+  const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token);
+
+  if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
+
+  // check hash type
+
+  if (token.buf[1][0] != 'S') return (PARSER_SIGNATURE_UNMATCHED);
+
+  // check iter
+
+  u32 iter = ROUNDS_QNX;
+
+  if (token.len[1] > 1)
+  {
+    if (token.buf[1][1] != ',') return (PARSER_SEPARATOR_UNMATCHED);
+
+    iter = hc_strtoul ((const char *) token.buf[1] + 2, NULL, 10);
+  }
+
+  // iter++; the additinal round is added in the init kernel
+
+  salt->salt_iter = iter;
+
+  // digest
+
+  if (token.len[2] != 128) return (PARSER_HASH_LENGTH);
+
+  digest[0] = hex_to_u64 ((const u8 *) token.buf[2] +   0);
+  digest[1] = hex_to_u64 ((const u8 *) token.buf[2] +  16);
+  digest[2] = hex_to_u64 ((const u8 *) token.buf[2] +  32);
+  digest[3] = hex_to_u64 ((const u8 *) token.buf[2] +  48);
+  digest[4] = hex_to_u64 ((const u8 *) token.buf[2] +  64);
+  digest[5] = hex_to_u64 ((const u8 *) token.buf[2] +  80);
+  digest[6] = hex_to_u64 ((const u8 *) token.buf[2] +  96);
+  digest[7] = hex_to_u64 ((const u8 *) token.buf[2] + 112);
+
+  // salt
+
+  if ((token.len[3] == 8) || (token.len[3] == 16))
+  {
+    memcpy (salt->salt_buf, token.buf[3], token.len[3]);
+
+    salt->salt_len = token.len[3];
+  }
+  else
+  {
+    return (PARSER_SALT_LENGTH);
+  }
+
+  return (PARSER_OK);
+}
+
+int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size)
+{
+  const u64 *digest = (const u64 *) digest_buf;
+
+  int line_len = 0;
+
+  const int iter = salt->salt_iter;
+
+  if (iter == ROUNDS_QNX)
+  {
+    line_buf[line_len++] = '@';
+    line_buf[line_len++] = 'S';
+    line_buf[line_len++] = '@';
+  }
+  else
+  {
+    line_buf[line_len++] = '@';
+    line_buf[line_len++] = 'S';
+    line_buf[line_len++] = ',';
+
+    line_len += snprintf (line_buf + line_len, line_size - line_len, "%u", iter);
+
+    line_buf[line_len++] = '@';
+  }
+
+  u64_to_hex (digest[0], (u8 *) line_buf + line_len); line_len += 16;
+  u64_to_hex (digest[1], (u8 *) line_buf + line_len); line_len += 16;
+  u64_to_hex (digest[2], (u8 *) line_buf + line_len); line_len += 16;
+  u64_to_hex (digest[3], (u8 *) line_buf + line_len); line_len += 16;
+  u64_to_hex (digest[4], (u8 *) line_buf + line_len); line_len += 16;
+  u64_to_hex (digest[5], (u8 *) line_buf + line_len); line_len += 16;
+  u64_to_hex (digest[6], (u8 *) line_buf + line_len); line_len += 16;
+  u64_to_hex (digest[7], (u8 *) line_buf + line_len); line_len += 16;
+
+  line_buf[line_len++] = '@';
+
+  memcpy (line_buf + line_len, salt->salt_buf, salt->salt_len);
+
+  line_len += salt->salt_len;
+
+  line_buf[line_len] = 0;
+
+  return line_len;
+}
+
+void module_init (module_ctx_t *module_ctx)
+{
+  module_ctx->module_context_size             = MODULE_CONTEXT_SIZE_CURRENT;
+  module_ctx->module_interface_version        = MODULE_INTERFACE_VERSION_CURRENT;
+
+  module_ctx->module_attack_exec              = module_attack_exec;
+  module_ctx->module_benchmark_esalt          = MODULE_DEFAULT;
+  module_ctx->module_benchmark_hook_salt      = MODULE_DEFAULT;
+  module_ctx->module_benchmark_mask           = MODULE_DEFAULT;
+  module_ctx->module_benchmark_salt           = MODULE_DEFAULT;
+  module_ctx->module_build_plain_postprocess  = MODULE_DEFAULT;
+  module_ctx->module_deep_comp_kernel         = MODULE_DEFAULT;
+  module_ctx->module_dgst_pos0                = module_dgst_pos0;
+  module_ctx->module_dgst_pos1                = module_dgst_pos1;
+  module_ctx->module_dgst_pos2                = module_dgst_pos2;
+  module_ctx->module_dgst_pos3                = module_dgst_pos3;
+  module_ctx->module_dgst_size                = module_dgst_size;
+  module_ctx->module_dictstat_disable         = MODULE_DEFAULT;
+  module_ctx->module_esalt_size               = MODULE_DEFAULT;
+  module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
+  module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;
+  module_ctx->module_hash_binary_save         = MODULE_DEFAULT;
+  module_ctx->module_hash_decode_outfile      = MODULE_DEFAULT;
+  module_ctx->module_hash_decode_zero_hash    = MODULE_DEFAULT;
+  module_ctx->module_hash_decode              = module_hash_decode;
+  module_ctx->module_hash_encode_status       = MODULE_DEFAULT;
+  module_ctx->module_hash_encode              = module_hash_encode;
+  module_ctx->module_hash_init_selftest       = MODULE_DEFAULT;
+  module_ctx->module_hash_mode                = MODULE_DEFAULT;
+  module_ctx->module_hash_category            = module_hash_category;
+  module_ctx->module_hash_name                = module_hash_name;
+  module_ctx->module_hlfmt_disable            = MODULE_DEFAULT;
+  module_ctx->module_hook12                   = MODULE_DEFAULT;
+  module_ctx->module_hook23                   = MODULE_DEFAULT;
+  module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
+  module_ctx->module_hook_size                = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
+  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_min       = MODULE_DEFAULT;
+  module_ctx->module_kern_type                = module_kern_type;
+  module_ctx->module_kern_type_dynamic        = MODULE_DEFAULT;
+  module_ctx->module_opti_type                = module_opti_type;
+  module_ctx->module_opts_type                = module_opts_type;
+  module_ctx->module_outfile_check_disable    = MODULE_DEFAULT;
+  module_ctx->module_outfile_check_nocomp     = MODULE_DEFAULT;
+  module_ctx->module_potfile_disable          = MODULE_DEFAULT;
+  module_ctx->module_potfile_keep_all_hashes  = MODULE_DEFAULT;
+  module_ctx->module_pwdump_column            = MODULE_DEFAULT;
+  module_ctx->module_pw_max                   = MODULE_DEFAULT;
+  module_ctx->module_pw_min                   = MODULE_DEFAULT;
+  module_ctx->module_salt_max                 = MODULE_DEFAULT;
+  module_ctx->module_salt_min                 = MODULE_DEFAULT;
+  module_ctx->module_salt_type                = module_salt_type;
+  module_ctx->module_separator                = MODULE_DEFAULT;
+  module_ctx->module_st_hash                  = module_st_hash;
+  module_ctx->module_st_pass                  = module_st_pass;
+  module_ctx->module_tmp_size                 = module_tmp_size;
+  module_ctx->module_unstable_warning         = MODULE_DEFAULT;
+  module_ctx->module_warmup_disable           = MODULE_DEFAULT;
+}
diff --git a/tools/test_modules/m19000.pm b/tools/test_modules/m19000.pm
new file mode 100644
index 000000000..7a87e50e6
--- /dev/null
+++ b/tools/test_modules/m19000.pm
@@ -0,0 +1,65 @@
+#!/usr/bin/env perl
+
+##
+## Author......: See docs/credits.txt
+## License.....: MIT
+##
+
+use strict;
+use warnings;
+
+use Digest::MD5 qw (md5_hex);
+
+sub module_constraints { [[0, 256], [16, 16], [-1, -1], [-1, -1], [-1, -1]] }
+
+sub module_generate_hash
+{
+  my $word = shift;
+  my $salt = shift;
+  my $iter = shift // 1000;
+
+  my $data;
+
+  $data .= $salt;
+  $data .= $word x $iter;
+  $data .= $word;
+
+  my $digest = md5_hex ($data);
+
+  my $hash;
+
+  if ($iter == 1000)
+  {
+    $hash = sprintf ("\@m\@%s\@%s", $digest, $salt);
+  }
+  else
+  {
+    $hash = sprintf ("\@m,%u\@%s\@%s", $iter, $digest, $salt);
+  }
+
+  return $hash;
+}
+
+sub module_verify_hash
+{
+  my $line = shift;
+
+  my ($hash, $word) = split (':', $line);
+
+  return unless defined $hash;
+  return unless defined $word;
+
+  my (undef, $tag, $digest, $salt) = split (/\@/, $hash);
+
+  my ($type, $iter) = split (/\,/, $tag);
+
+  return unless ($type eq "m");
+
+  my $word_packed = pack_if_HEX_notation ($word);
+
+  my $new_hash = module_generate_hash ($word_packed, $salt, $iter);
+
+  return ($new_hash, $word);
+}
+
+1;
diff --git a/tools/test_modules/m19100.pm b/tools/test_modules/m19100.pm
new file mode 100644
index 000000000..23380bcf8
--- /dev/null
+++ b/tools/test_modules/m19100.pm
@@ -0,0 +1,65 @@
+#!/usr/bin/env perl
+
+##
+## Author......: See docs/credits.txt
+## License.....: MIT
+##
+
+use strict;
+use warnings;
+
+use Digest::SHA qw (sha256_hex);
+
+sub module_constraints { [[0, 256], [16, 16], [-1, -1], [-1, -1], [-1, -1]] }
+
+sub module_generate_hash
+{
+  my $word = shift;
+  my $salt = shift;
+  my $iter = shift // 1000;
+
+  my $data;
+
+  $data .= $salt;
+  $data .= $word x $iter;
+  $data .= $word;
+
+  my $digest = sha256_hex ($data);
+
+  my $hash;
+
+  if ($iter == 1000)
+  {
+    $hash = sprintf ("\@s\@%s\@%s", $digest, $salt);
+  }
+  else
+  {
+    $hash = sprintf ("\@s,%u\@%s\@%s", $iter, $digest, $salt);
+  }
+
+  return $hash;
+}
+
+sub module_verify_hash
+{
+  my $line = shift;
+
+  my ($hash, $word) = split (':', $line);
+
+  return unless defined $hash;
+  return unless defined $word;
+
+  my (undef, $tag, $digest, $salt) = split (/\@/, $hash);
+
+  my ($type, $iter) = split (/\,/, $tag);
+
+  return unless ($type eq "s");
+
+  my $word_packed = pack_if_HEX_notation ($word);
+
+  my $new_hash = module_generate_hash ($word_packed, $salt, $iter);
+
+  return ($new_hash, $word);
+}
+
+1;
diff --git a/tools/test_modules/m19200.pm b/tools/test_modules/m19200.pm
new file mode 100644
index 000000000..19f5edfed
--- /dev/null
+++ b/tools/test_modules/m19200.pm
@@ -0,0 +1,65 @@
+#!/usr/bin/env perl
+
+##
+## Author......: See docs/credits.txt
+## License.....: MIT
+##
+
+use strict;
+use warnings;
+
+use Digest::SHA qw (sha512_hex);
+
+sub module_constraints { [[0, 256], [16, 16], [-1, -1], [-1, -1], [-1, -1]] }
+
+sub module_generate_hash
+{
+  my $word = shift;
+  my $salt = shift;
+  my $iter = shift // 1000;
+
+  my $data;
+
+  $data .= $salt;
+  $data .= $word x $iter;
+  $data .= $word;
+
+  my $digest = sha512_hex ($data);
+
+  my $hash;
+
+  if ($iter == 1000)
+  {
+    $hash = sprintf ("\@S\@%s\@%s", $digest, $salt);
+  }
+  else
+  {
+    $hash = sprintf ("\@S,%u\@%s\@%s", $iter, $digest, $salt);
+  }
+
+  return $hash;
+}
+
+sub module_verify_hash
+{
+  my $line = shift;
+
+  my ($hash, $word) = split (':', $line);
+
+  return unless defined $hash;
+  return unless defined $word;
+
+  my (undef, $tag, $digest, $salt) = split (/\@/, $hash);
+
+  my ($type, $iter) = split (/\,/, $tag);
+
+  return unless ($type eq "S");
+
+  my $word_packed = pack_if_HEX_notation ($word);
+
+  my $new_hash = module_generate_hash ($word_packed, $salt, $iter);
+
+  return ($new_hash, $word);
+}
+
+1;