diff --git a/OpenCL/inc_luks_aes.cl b/OpenCL/inc_luks_aes.cl
index 3b8499820..53ecf6812 100644
--- a/OpenCL/inc_luks_aes.cl
+++ b/OpenCL/inc_luks_aes.cl
@@ -31,7 +31,7 @@ static void aes256_decrypt_cbc (const u32 *ks1, const u32 *in, u32 *out, u32 *es
   essiv[3] = in[3];
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -65,7 +65,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128 (__global u32 *in, u32 *out, co
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -102,7 +102,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1 (__global u32 *in, u32
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -162,7 +162,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha1_final (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -199,7 +199,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -259,7 +259,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha256_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -296,7 +296,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -356,7 +356,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_sha512_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -393,7 +393,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -453,7 +453,7 @@ static void luks_decrypt_sector_aes_cbc_essiv128_mk_ripemd160_final (__global u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -487,7 +487,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256 (__global u32 *in, u32 *out, co
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -524,7 +524,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1 (__global u32 *in, u32
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -584,7 +584,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha1_final (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -621,7 +621,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -681,7 +681,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha256_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -718,7 +718,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -778,7 +778,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_sha512_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -815,7 +815,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -877,7 +877,7 @@ static void luks_decrypt_sector_aes_cbc_essiv256_mk_ripemd160_final (__global u3
 
 // cbc-plain
 
-static void luks_decrypt_sector_aes_cbc_plain128 (__global u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -907,7 +907,7 @@ static void luks_decrypt_sector_aes_cbc_plain128 (__global u32 *in, u32 *out, co
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -940,7 +940,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1 (__global u32 *in, u32
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -996,7 +996,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha1_final (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1029,7 +1029,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1085,7 +1085,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha256_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1118,7 +1118,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1174,7 +1174,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_sha512_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1207,7 +1207,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1263,7 +1263,7 @@ static void luks_decrypt_sector_aes_cbc_plain128_mk_ripemd160_final (__global u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1293,7 +1293,7 @@ static void luks_decrypt_sector_aes_cbc_plain256 (__global u32 *in, u32 *out, co
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1326,7 +1326,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1 (__global u32 *in, u32
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1382,7 +1382,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha1_final (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1415,7 +1415,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1471,7 +1471,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha256_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1504,7 +1504,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1560,7 +1560,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_sha512_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1593,7 +1593,7 @@ static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_cbc_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1695,7 +1695,7 @@ static void aes256_decrypt_xts (const u32 *ks1, const u32 *in, u32 *out, u32 *T,
   xts_mul2 (T, T);
 }
 
-static void luks_decrypt_sector_aes_xts_plain256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1729,7 +1729,7 @@ static void luks_decrypt_sector_aes_xts_plain256 (__global u32 *in, u32 *out, co
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1766,7 +1766,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha1 (__global u32 *in, u32
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1826,7 +1826,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha1_final (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1863,7 +1863,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha256 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1923,7 +1923,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha256_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1960,7 +1960,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha512 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2020,7 +2020,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_sha512_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2057,7 +2057,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2117,7 +2117,7 @@ static void luks_decrypt_sector_aes_xts_plain256_mk_ripemd160_final (__global u3
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2151,7 +2151,7 @@ static void luks_decrypt_sector_aes_xts_plain512 (__global u32 *in, u32 *out, co
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2188,7 +2188,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha1 (__global u32 *in, u32
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2248,7 +2248,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha1_final (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2285,7 +2285,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha256 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2345,7 +2345,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha256_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2382,7 +2382,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha512 (__global u32 *in, u3
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2442,7 +2442,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_sha512_final (__global u32 *
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2479,7 +2479,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, const u32 sector, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2541,7 +2541,7 @@ static void luks_decrypt_sector_aes_xts_plain512_mk_ripemd160_final (__global u3
 
 // luks helper
 
-static void luks_af_sha1_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_af_sha1_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -2932,7 +2932,7 @@ static void luks_af_sha1_then_aes_decrypt (__global luks_t *luks_bufs, __global
   }
 }
 
-static void luks_af_sha256_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_af_sha256_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -3291,7 +3291,7 @@ static void luks_af_sha256_then_aes_decrypt (__global luks_t *luks_bufs, __globa
   }
 }
 
-static void luks_af_sha512_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_af_sha512_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -3669,7 +3669,7 @@ static void luks_af_sha512_then_aes_decrypt (__global luks_t *luks_bufs, __globa
   }
 }
 
-static void luks_af_ripemd160_then_aes_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+static void luks_af_ripemd160_then_aes_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
diff --git a/OpenCL/inc_luks_serpent.cl b/OpenCL/inc_luks_serpent.cl
index 00068776b..9f0dcf2cf 100644
--- a/OpenCL/inc_luks_serpent.cl
+++ b/OpenCL/inc_luks_serpent.cl
@@ -31,7 +31,7 @@ static void serpent256_decrypt_cbc (const u32 *ks1, const u32 *in, u32 *out, u32
   essiv[3] = in[3];
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -65,7 +65,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -102,7 +102,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -162,7 +162,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -199,7 +199,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -259,7 +259,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -296,7 +296,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -356,7 +356,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -393,7 +393,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -453,7 +453,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv128_mk_ripemd160_final (__globa
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -487,7 +487,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -524,7 +524,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -584,7 +584,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -621,7 +621,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -681,7 +681,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -718,7 +718,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -778,7 +778,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -815,7 +815,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -877,7 +877,7 @@ static void luks_decrypt_sector_serpent_cbc_essiv256_mk_ripemd160_final (__globa
 
 // cbc-plain
 
-static void luks_decrypt_sector_serpent_cbc_plain128 (__global u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128 (__global const u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -907,7 +907,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -940,7 +940,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -996,7 +996,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1029,7 +1029,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1085,7 +1085,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1118,7 +1118,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1174,7 +1174,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1207,7 +1207,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1263,7 +1263,7 @@ static void luks_decrypt_sector_serpent_cbc_plain128_mk_ripemd160_final (__globa
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256 (__global u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1293,7 +1293,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1326,7 +1326,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1382,7 +1382,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1415,7 +1415,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1471,7 +1471,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1504,7 +1504,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1560,7 +1560,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1593,7 +1593,7 @@ static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_cbc_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1695,7 +1695,7 @@ static void serpent256_decrypt_xts (const u32 *ks1, const u32 *in, u32 *out, u32
   xts_mul2 (T, T);
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1729,7 +1729,7 @@ static void luks_decrypt_sector_serpent_xts_plain256 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1766,7 +1766,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1826,7 +1826,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1863,7 +1863,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1923,7 +1923,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1960,7 +1960,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2020,7 +2020,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2057,7 +2057,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2117,7 +2117,7 @@ static void luks_decrypt_sector_serpent_xts_plain256_mk_ripemd160_final (__globa
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512 (__global u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512 (__global const u32 *in, u32 *out, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2151,7 +2151,7 @@ static void luks_decrypt_sector_serpent_xts_plain512 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2188,7 +2188,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2248,7 +2248,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2285,7 +2285,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2345,7 +2345,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2382,7 +2382,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2442,7 +2442,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2479,7 +2479,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
+static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *ks1, const u32 *ks2, volatile const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2541,7 +2541,7 @@ static void luks_decrypt_sector_serpent_xts_plain512_mk_ripemd160_final (__globa
 
 // luks helper
 
-static void luks_af_sha1_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_sha1_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -2900,7 +2900,7 @@ static void luks_af_sha1_then_serpent_decrypt (__global luks_t *luks_bufs, __glo
   }
 }
 
-static void luks_af_sha256_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_sha256_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -3259,7 +3259,7 @@ static void luks_af_sha256_then_serpent_decrypt (__global luks_t *luks_bufs, __g
   }
 }
 
-static void luks_af_sha512_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_sha512_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -3637,7 +3637,7 @@ static void luks_af_sha512_then_serpent_decrypt (__global luks_t *luks_bufs, __g
   }
 }
 
-static void luks_af_ripemd160_then_serpent_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_ripemd160_then_serpent_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
diff --git a/OpenCL/inc_luks_twofish.cl b/OpenCL/inc_luks_twofish.cl
index 38800a872..e17ac1c1b 100644
--- a/OpenCL/inc_luks_twofish.cl
+++ b/OpenCL/inc_luks_twofish.cl
@@ -31,7 +31,7 @@ static void twofish256_decrypt_cbc (const u32 *sk1, const u32 *lk1, const u32 *i
   essiv[3] = in[3];
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -65,7 +65,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -102,7 +102,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -162,7 +162,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -199,7 +199,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -259,7 +259,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -296,7 +296,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -356,7 +356,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -393,7 +393,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -453,7 +453,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv128_mk_ripemd160_final (__globa
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -487,7 +487,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -524,7 +524,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -584,7 +584,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -621,7 +621,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -681,7 +681,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -718,7 +718,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -778,7 +778,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -815,7 +815,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -877,7 +877,7 @@ static void luks_decrypt_sector_twofish_cbc_essiv256_mk_ripemd160_final (__globa
 
 // cbc-plain
 
-static void luks_decrypt_sector_twofish_cbc_plain128 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -907,7 +907,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -940,7 +940,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -996,7 +996,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1029,7 +1029,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1085,7 +1085,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1118,7 +1118,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1174,7 +1174,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1207,7 +1207,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1263,7 +1263,7 @@ static void luks_decrypt_sector_twofish_cbc_plain128_mk_ripemd160_final (__globa
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1293,7 +1293,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1326,7 +1326,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1382,7 +1382,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1415,7 +1415,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1471,7 +1471,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1504,7 +1504,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1560,7 +1560,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1593,7 +1593,7 @@ static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
+static void luks_decrypt_sector_twofish_cbc_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1695,7 +1695,7 @@ static void twofish256_decrypt_xts (const u32 *sk1, const u32 *lk1, const u32 *i
   xts_mul2 (T, T);
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1729,7 +1729,7 @@ static void luks_decrypt_sector_twofish_xts_plain256 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1766,7 +1766,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1826,7 +1826,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1863,7 +1863,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1923,7 +1923,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -1960,7 +1960,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2020,7 +2020,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2057,7 +2057,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2117,7 +2117,7 @@ static void luks_decrypt_sector_twofish_xts_plain256_mk_ripemd160_final (__globa
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512 (__global u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512 (__global const u32 *in, u32 *out, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2151,7 +2151,7 @@ static void luks_decrypt_sector_twofish_xts_plain512 (__global u32 *in, u32 *out
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2188,7 +2188,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1 (__global u32 *in,
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2248,7 +2248,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha1_final (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2285,7 +2285,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2345,7 +2345,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha256_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2382,7 +2382,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512 (__global u32 *in
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2442,7 +2442,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_sha512_final (__global u
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160 (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160 (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2479,7 +2479,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160 (__global u32
   }
 }
 
-static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160_final (__global u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
+static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160_final (__global const u32 *in, u32 *mk, const u32 *sk1, const u32 *lk1, const u32 *sk2, const u32 *lk2, const u32 sector)
 {
   u32 S[4] = { sector, 0, 0, 0 };
 
@@ -2541,7 +2541,7 @@ static void luks_decrypt_sector_twofish_xts_plain512_mk_ripemd160_final (__globa
 
 // luks helper
 
-static void luks_af_sha1_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_sha1_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -2900,7 +2900,7 @@ static void luks_af_sha1_then_twofish_decrypt (__global luks_t *luks_bufs, __glo
   }
 }
 
-static void luks_af_sha256_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_sha256_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -3259,7 +3259,7 @@ static void luks_af_sha256_then_twofish_decrypt (__global luks_t *luks_bufs, __g
   }
 }
 
-static void luks_af_sha512_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_sha512_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
@@ -3637,7 +3637,7 @@ static void luks_af_sha512_then_twofish_decrypt (__global luks_t *luks_bufs, __g
   }
 }
 
-static void luks_af_ripemd160_then_twofish_decrypt (__global luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
+static void luks_af_ripemd160_then_twofish_decrypt (__global const luks_t *luks_bufs, __global luks_tmp_t *tmps, u32 *pt_buf)
 {
   const u32 key_size    = luks_bufs->key_size;
   const u32 cipher_mode = luks_bufs->cipher_mode;
diff --git a/OpenCL/m14612.cl b/OpenCL/m14612.cl
index 501a5daf5..992429ca2 100644
--- a/OpenCL/m14612.cl
+++ b/OpenCL/m14612.cl
@@ -61,7 +61,7 @@ void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[
   sha1_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -135,7 +135,7 @@ __kernel void m14612_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14612_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14612_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -222,7 +222,7 @@ __kernel void m14612_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14612_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14612_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
diff --git a/OpenCL/m14613.cl b/OpenCL/m14613.cl
index e99737a09..d328e3cb6 100644
--- a/OpenCL/m14613.cl
+++ b/OpenCL/m14613.cl
@@ -61,7 +61,7 @@ void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[
   sha1_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -135,7 +135,7 @@ __kernel void m14613_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14613_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14613_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -222,7 +222,7 @@ __kernel void m14613_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14613_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14613_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
diff --git a/OpenCL/m14621.cl b/OpenCL/m14621.cl
index 3c83d7cb6..622945efd 100644
--- a/OpenCL/m14621.cl
+++ b/OpenCL/m14621.cl
@@ -67,7 +67,7 @@ void hmac_sha256_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipa
   sha256_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -153,7 +153,7 @@ __kernel void m14621_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14621_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14621_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -261,7 +261,7 @@ __kernel void m14621_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14621_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14621_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
   const u64 lid = get_local_id (0);
diff --git a/OpenCL/m14622.cl b/OpenCL/m14622.cl
index 03ff787c6..a66ec0e93 100644
--- a/OpenCL/m14622.cl
+++ b/OpenCL/m14622.cl
@@ -67,7 +67,7 @@ void hmac_sha256_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipa
   sha256_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -153,7 +153,7 @@ __kernel void m14622_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14622_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14622_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -261,7 +261,7 @@ __kernel void m14622_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14622_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14622_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
diff --git a/OpenCL/m14623.cl b/OpenCL/m14623.cl
index 7859fa949..79937b817 100644
--- a/OpenCL/m14623.cl
+++ b/OpenCL/m14623.cl
@@ -67,7 +67,7 @@ void hmac_sha256_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipa
   sha256_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -153,7 +153,7 @@ __kernel void m14623_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14623_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14623_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -261,7 +261,7 @@ __kernel void m14623_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14623_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14623_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
diff --git a/OpenCL/m14631.cl b/OpenCL/m14631.cl
index bc1c01bab..0d4a7f84f 100644
--- a/OpenCL/m14631.cl
+++ b/OpenCL/m14631.cl
@@ -83,7 +83,7 @@ void hmac_sha512_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[
   sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest);
 }
 
-__kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -189,7 +189,7 @@ __kernel void m14631_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14631_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14631_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -317,7 +317,7 @@ __kernel void m14631_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14631_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14631_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
   const u64 lid = get_local_id (0);
diff --git a/OpenCL/m14632.cl b/OpenCL/m14632.cl
index 38d2ef114..86f273ee9 100644
--- a/OpenCL/m14632.cl
+++ b/OpenCL/m14632.cl
@@ -83,7 +83,7 @@ void hmac_sha512_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[
   sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest);
 }
 
-__kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -189,7 +189,7 @@ __kernel void m14632_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14632_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14632_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -317,7 +317,7 @@ __kernel void m14632_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14632_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14632_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
diff --git a/OpenCL/m14633.cl b/OpenCL/m14633.cl
index 5aabd1007..ae2e98a85 100644
--- a/OpenCL/m14633.cl
+++ b/OpenCL/m14633.cl
@@ -83,7 +83,7 @@ void hmac_sha512_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[
   sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest);
 }
 
-__kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -189,7 +189,7 @@ __kernel void m14633_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14633_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14633_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -317,7 +317,7 @@ __kernel void m14633_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14633_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14633_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
diff --git a/OpenCL/m14641.cl b/OpenCL/m14641.cl
index 12d3f7215..240c1fb43 100644
--- a/OpenCL/m14641.cl
+++ b/OpenCL/m14641.cl
@@ -61,7 +61,7 @@ void hmac_ripemd160_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x
   ripemd160_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -135,7 +135,7 @@ __kernel void m14641_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14641_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14641_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -222,7 +222,7 @@ __kernel void m14641_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14641_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14641_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
   const u64 lid = get_local_id (0);
diff --git a/OpenCL/m14642.cl b/OpenCL/m14642.cl
index 7bb6caf41..5bad9397d 100644
--- a/OpenCL/m14642.cl
+++ b/OpenCL/m14642.cl
@@ -61,7 +61,7 @@ void hmac_ripemd160_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x
   ripemd160_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -135,7 +135,7 @@ __kernel void m14642_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14642_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14642_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -222,7 +222,7 @@ __kernel void m14642_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14642_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14642_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
diff --git a/OpenCL/m14643.cl b/OpenCL/m14643.cl
index 0ef4addbf..0bb93ccee 100644
--- a/OpenCL/m14643.cl
+++ b/OpenCL/m14643.cl
@@ -61,7 +61,7 @@ void hmac_ripemd160_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x
   ripemd160_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   /**
    * base
@@ -135,7 +135,7 @@ __kernel void m14643_init (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14643_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14643_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);
 
@@ -222,7 +222,7 @@ __kernel void m14643_loop (__global pw_t *pws, __global const kernel_rule_t *rul
   }
 }
 
-__kernel void m14643_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
+__kernel void m14643_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global luks_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const luks_t *luks_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
 {
   const u64 gid = get_global_id (0);