diff --git a/OpenCL/m12150-pure.cl b/OpenCL/m12150-pure.cl
index 5b837c13d..e6d8f56f0 100644
--- a/OpenCL/m12150-pure.cl
+++ b/OpenCL/m12150-pure.cl
@@ -31,7 +31,7 @@ typedef struct shiro1_sha512_tmp
   u64 out[8];  // Final output hash
 } shiro1_sha512_tmp_t;
 
-KERNEL_FQ void m12150_init (KERN_ATTR_TMPS_ESALT (shiro1_sha512_tmp_t, shiro1_sha512_t))
+KERNEL_FQ void m12150_init (KERN_ATTR_TMPS (shiro1_sha512_tmp_t))
 {
   const u32 gid = get_global_id (0);
 
@@ -53,7 +53,7 @@ KERNEL_FQ void m12150_init (KERN_ATTR_TMPS_ESALT (shiro1_sha512_tmp_t, shiro1_sh
   }
 }
 
-KERNEL_FQ void m12150_loop(KERN_ATTR_TMPS_ESALT(shiro1_sha512_tmp_t, shiro1_sha512_t)) {
+KERNEL_FQ void m12150_loop(KERN_ATTR_TMPS (shiro1_sha512_tmp_t)) {
   const u32 gid = get_global_id(0);
 
   if (gid >= GID_CNT) return;
@@ -115,7 +115,7 @@ KERNEL_FQ void m12150_loop(KERN_ATTR_TMPS_ESALT(shiro1_sha512_tmp_t, shiro1_sha5
   }
 }
 
-KERNEL_FQ void m12150_comp (KERN_ATTR_TMPS_ESALT (shiro1_sha512_tmp_t, shiro1_sha512_t))
+KERNEL_FQ void m12150_comp (KERN_ATTR_TMPS (shiro1_sha512_tmp_t))
 {
   const u64 gid = get_global_id (0);
 
diff --git a/src/modules/module_12150.c b/src/modules/module_12150.c
index 91d1c1e6f..82fb91abe 100644
--- a/src/modules/module_12150.c
+++ b/src/modules/module_12150.c
@@ -32,20 +32,20 @@ static const u32   SALT_TYPE      = SALT_TYPE_GENERIC;
 static const char *ST_PASS        = "admin123";
 static const char *ST_HASH        = "$shiro1$SHA-512$1024$NE+wqQq/TmjZMvfI7ENh/g==$V4yPw8T64UQ6GfJfxYq2hLsVrBY8D1v+bktfOxGdt4b/9BthpWPNUy/CBk6V9iA0nHpzYzJFWO8v/tZFtES8CA==";
 
-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;         }
+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 shiro1_sha512
 {
@@ -59,7 +59,7 @@ typedef struct shiro1_sha512_tmp
   u64 out[8];  // Final output hash
 } shiro1_sha512_tmp_t;
 
-static const char *SIGNATURE_SHIRO1_SHA512 = "$shiro1$";
+static const char *SIGNATURE_SHIRO1_SHA512 = "$shiro1$SHA-512$";
 
 u64 module_esalt_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)
 {
@@ -85,71 +85,62 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con
 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;
-  shiro1_sha512_t *shiro_sha512 = (shiro1_sha512_t *) esalt_buf;
 
   hc_token_t token;
 
   memset (&token, 0, sizeof (hc_token_t));
 
-  token.token_cnt  = 5;
+  token.token_cnt  = 4;
 
   token.signatures_cnt    = 1;
   token.signatures_buf[0] = SIGNATURE_SHIRO1_SHA512;
 
-  token.len[0]     = 8;
+  token.len[0]     = 16;
   token.attr[0]    = TOKEN_ATTR_FIXED_LENGTH
                     | TOKEN_ATTR_VERIFY_SIGNATURE;
 
   token.sep[1]     = '$';
-  token.len[1]     = 7;
-  token.attr[1]    = TOKEN_ATTR_FIXED_LENGTH;                    
-
-  token.sep[2]     = '$';
-  token.len_min[2] = 1;
-  token.len_max[2] = 999999;
-  token.attr[2]    = TOKEN_ATTR_VERIFY_LENGTH
+  token.len_min[1] = 1;
+  token.len_max[1] = 999999;
+  token.attr[1]    = TOKEN_ATTR_VERIFY_LENGTH
                     | TOKEN_ATTR_VERIFY_DIGIT;
 
-  token.sep[3]     = '$';
-  token.len_min[3] = ((SALT_MIN * 8) / 6) + 0;
-  token.len_max[3] = ((SALT_MAX * 8) / 6) + 3;
-  token.attr[3]    = TOKEN_ATTR_VERIFY_LENGTH
+  token.sep[2]     = '$';
+  token.len_min[2] = ((SALT_MIN * 8) / 6) + 0;
+  token.len_max[2] = ((SALT_MAX * 8) / 6) + 3;
+  token.attr[2]    = TOKEN_ATTR_VERIFY_LENGTH
                     | TOKEN_ATTR_VERIFY_BASE64A;
 
-  token.sep[4]     = '$';
-  token.len_min[4] = 16;
-  token.len_max[4] = 256;
-  token.attr[4]    = TOKEN_ATTR_VERIFY_LENGTH
+  token.sep[3]     = '$';
+  token.len_min[3] = 16;
+  token.len_max[3] = 256;
+  token.attr[3]    = TOKEN_ATTR_VERIFY_LENGTH
                     | TOKEN_ATTR_VERIFY_BASE64A;
 
   const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token);
 
-  if (rc_tokenizer != PARSER_OK) { return (rc_tokenizer); }
+  if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
 
   u8  tmp_buf[512];
   int tmp_len;
 
   // iterations
-  const u8 *iter_pos = token.buf[2];
+  const u8 *iter_pos = token.buf[1];
   const u32 iter = hc_strtoul ((const char *) iter_pos, NULL, 10);
-  shiro_sha512->iterations = iter - 1;
   salt->salt_iter = iter - 1;
 
   // salt
-  const u8 *salt_pos = token.buf[3];
-  const int salt_len = token.len[3];
+  const u8 *salt_pos = token.buf[2];
+  const int salt_len = token.len[2];
   memset (tmp_buf, 0, sizeof (tmp_buf));
   tmp_len = base64_decode (base64_to_int, salt_pos, salt_len, tmp_buf);
-  memcpy (shiro_sha512->salt_buf, tmp_buf, tmp_len);
+  if (tmp_len > SALT_MAX) return (PARSER_SALT_LENGTH);
+  memcpy (salt->salt_buf, tmp_buf, tmp_len);
   salt->salt_len = tmp_len;
-  salt->salt_buf[0] = shiro_sha512->salt_buf[0];
-  salt->salt_buf[1] = shiro_sha512->salt_buf[1];
-  salt->salt_buf[2] = shiro_sha512->salt_buf[2];
-  salt->salt_buf[3] = shiro_sha512->salt_buf[3];
 
   // hash
-  const u8 *hash_pos = token.buf[4];
-  const int hash_len = token.len[4];
+  const u8 *hash_pos = token.buf[3];
+  const int hash_len = token.len[3];
   memset (tmp_buf, 0, sizeof (tmp_buf));
   tmp_len = base64_decode (base64_to_int, hash_pos, hash_len, tmp_buf);
   if (tmp_len < 64) return (PARSER_HASH_LENGTH);