diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl
index 02db55a93..7a85722f8 100644
--- a/OpenCL/inc_common.cl
+++ b/OpenCL/inc_common.cl
@@ -236,6 +236,150 @@ DECLSPEC int is_valid_hex_32 (const u32 v)
   return 1;
 }
 
+DECLSPEC int find_keyboard_layout_map (const u32 search, const int search_len, __local keyboard_layout_mapping_t *s_keyboard_layout_mapping_buf, const int keyboard_layout_mapping_cnt)
+{
+  for (int idx = 0; idx < keyboard_layout_mapping_cnt; idx++)
+  {
+    const u32 src_char = s_keyboard_layout_mapping_buf[idx].src_char;
+    const int src_len  = s_keyboard_layout_mapping_buf[idx].src_len;
+
+    if (src_len == search_len)
+    {
+      const u32 mask = 0xffffffff >> ((4 - search_len) * 8);
+
+      if ((src_char & mask) == (search & mask)) return idx;
+    }
+  }
+
+  return -1;
+}
+
+DECLSPEC int execute_keyboard_layout_mapping (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int pw_len, __local keyboard_layout_mapping_t *s_keyboard_layout_mapping_buf, const int keyboard_layout_mapping_cnt)
+{
+  u32 out_buf[16] = { 0 };
+
+  u8 *out_ptr = (u8 *) out_buf;
+
+  int out_len = 0;
+
+  // TC/VC passwords are limited to 64
+
+  u32 w[16];
+
+  w[ 0] = w0[0];
+  w[ 1] = w0[1];
+  w[ 2] = w0[2];
+  w[ 3] = w0[3];
+  w[ 4] = w1[0];
+  w[ 5] = w1[1];
+  w[ 6] = w1[2];
+  w[ 7] = w1[3];
+  w[ 8] = w2[0];
+  w[ 9] = w2[1];
+  w[10] = w2[2];
+  w[11] = w2[3];
+  w[12] = w3[0];
+  w[13] = w3[1];
+  w[14] = w3[2];
+  w[15] = w3[3];
+
+  u8 *w_ptr = (u8 *) w;
+
+  int pw_pos = 0;
+
+  while (pw_pos < pw_len)
+  {
+    u32 src0 = 0;
+    u32 src1 = 0;
+    u32 src2 = 0;
+    u32 src3 = 0;
+
+    #define MIN(a,b) (((a) < (b)) ? (a) : (b))
+
+    const int rem = MIN (pw_len - pw_pos, 4);
+
+    #undef MIN
+
+    if (rem > 0) src0 = w_ptr[pw_pos + 0];
+    if (rem > 1) src1 = w_ptr[pw_pos + 1];
+    if (rem > 2) src2 = w_ptr[pw_pos + 2];
+    if (rem > 3) src3 = w_ptr[pw_pos + 3];
+
+    const u32 src = (src0 <<  0)
+                  | (src1 <<  8)
+                  | (src2 << 16)
+                  | (src3 << 24);
+
+    int src_len;
+
+    for (src_len = rem; src_len > 0; src_len--)
+    {
+      const int idx = find_keyboard_layout_map (src, src_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+
+      if (idx == -1) continue;
+
+      u32 dst_char = s_keyboard_layout_mapping_buf[idx].dst_char;
+      int dst_len  = s_keyboard_layout_mapping_buf[idx].dst_len;
+
+      switch (dst_len)
+      {
+        case 1:
+          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
+          break;
+        case 2:
+          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
+          out_ptr[out_len++] = (dst_char >>  8) & 0xff;
+          break;
+        case 3:
+          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
+          out_ptr[out_len++] = (dst_char >>  8) & 0xff;
+          out_ptr[out_len++] = (dst_char >> 16) & 0xff;
+          break;
+        case 4:
+          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
+          out_ptr[out_len++] = (dst_char >>  8) & 0xff;
+          out_ptr[out_len++] = (dst_char >> 16) & 0xff;
+          out_ptr[out_len++] = (dst_char >> 24) & 0xff;
+          break;
+      }
+
+      pw_pos += src_len;
+
+      break;
+    }
+
+    // not matched, keep original
+
+    if (src_len == 0)
+    {
+      out_ptr[out_len] = w_ptr[pw_pos];
+
+      out_len++;
+
+      pw_pos++;
+    }
+  }
+
+  w0[0] = out_buf[ 0];
+  w0[1] = out_buf[ 1];
+  w0[2] = out_buf[ 2];
+  w0[3] = out_buf[ 3];
+  w1[0] = out_buf[ 4];
+  w1[1] = out_buf[ 5];
+  w1[2] = out_buf[ 6];
+  w1[3] = out_buf[ 7];
+  w2[0] = out_buf[ 8];
+  w2[1] = out_buf[ 9];
+  w2[2] = out_buf[10];
+  w2[3] = out_buf[11];
+  w3[0] = out_buf[12];
+  w3[1] = out_buf[13];
+  w3[2] = out_buf[14];
+  w3[3] = out_buf[15];
+
+  return out_len;
+}
+
 /**
  * vector functions
  */
diff --git a/OpenCL/inc_truecrypt_keyboard.cl b/OpenCL/inc_truecrypt_keyboard.cl
deleted file mode 100644
index d0fb65fb9..000000000
--- a/OpenCL/inc_truecrypt_keyboard.cl
+++ /dev/null
@@ -1,143 +0,0 @@
-DECLSPEC int find_keyboard_layout_map (const u32 search, const int search_len, __local keyboard_layout_mapping_t *s_keyboard_layout_mapping_buf, const int keyboard_layout_mapping_cnt)
-{
-  for (int idx = 0; idx < keyboard_layout_mapping_cnt; idx++)
-  {
-    const u32 src_char = s_keyboard_layout_mapping_buf[idx].src_char;
-    const int src_len  = s_keyboard_layout_mapping_buf[idx].src_len;
-
-    if (src_len == search_len)
-    {
-      const u32 mask = 0xffffffff >> ((4 - search_len) * 8);
-
-      if ((src_char & mask) == (search & mask)) return idx;
-    }
-  }
-
-  return -1;
-}
-
-DECLSPEC int keyboard_map (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int pw_len, __local keyboard_layout_mapping_t *s_keyboard_layout_mapping_buf, const int keyboard_layout_mapping_cnt)
-{
-  u32 out_buf[16] = { 0 };
-
-  u8 *out_ptr = (u8 *) out_buf;
-
-  int out_len = 0;
-
-  // TC/VC passwords are limited to 64
-
-  u32 w[16];
-
-  w[ 0] = w0[0];
-  w[ 1] = w0[1];
-  w[ 2] = w0[2];
-  w[ 3] = w0[3];
-  w[ 4] = w1[0];
-  w[ 5] = w1[1];
-  w[ 6] = w1[2];
-  w[ 7] = w1[3];
-  w[ 8] = w2[0];
-  w[ 9] = w2[1];
-  w[10] = w2[2];
-  w[11] = w2[3];
-  w[12] = w3[0];
-  w[13] = w3[1];
-  w[14] = w3[2];
-  w[15] = w3[3];
-
-  u8 *w_ptr = (u8 *) w;
-
-  int pw_pos = 0;
-
-  while (pw_pos < pw_len)
-  {
-    u32 src0 = 0;
-    u32 src1 = 0;
-    u32 src2 = 0;
-    u32 src3 = 0;
-
-    #define MIN(a,b) (((a) < (b)) ? (a) : (b))
-
-    const int rem = MIN (pw_len - pw_pos, 4);
-
-    #undef MIN
-
-    if (rem > 0) src0 = w_ptr[pw_pos + 0];
-    if (rem > 1) src1 = w_ptr[pw_pos + 1];
-    if (rem > 2) src2 = w_ptr[pw_pos + 2];
-    if (rem > 3) src3 = w_ptr[pw_pos + 3];
-
-    const u32 src = (src0 <<  0)
-                  | (src1 <<  8)
-                  | (src2 << 16)
-                  | (src3 << 24);
-
-    int src_len;
-
-    for (src_len = rem; src_len > 0; src_len--)
-    {
-      const int idx = find_keyboard_layout_map (src, src_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
-
-      if (idx == -1) continue;
-
-      u32 dst_char = s_keyboard_layout_mapping_buf[idx].dst_char;
-      int dst_len  = s_keyboard_layout_mapping_buf[idx].dst_len;
-
-      switch (dst_len)
-      {
-        case 1:
-          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
-          break;
-        case 2:
-          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
-          out_ptr[out_len++] = (dst_char >>  8) & 0xff;
-          break;
-        case 3:
-          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
-          out_ptr[out_len++] = (dst_char >>  8) & 0xff;
-          out_ptr[out_len++] = (dst_char >> 16) & 0xff;
-          break;
-        case 4:
-          out_ptr[out_len++] = (dst_char >>  0) & 0xff;
-          out_ptr[out_len++] = (dst_char >>  8) & 0xff;
-          out_ptr[out_len++] = (dst_char >> 16) & 0xff;
-          out_ptr[out_len++] = (dst_char >> 24) & 0xff;
-          break;
-      }
-
-      pw_pos += src_len;
-
-      break;
-    }
-
-    // not matched, keep original
-
-    if (src_len == 0)
-    {
-      out_ptr[out_len] = w_ptr[pw_pos];
-
-      out_len++;
-
-      pw_pos++;
-    }
-  }
-
-  w0[0] = out_buf[ 0];
-  w0[1] = out_buf[ 1];
-  w0[2] = out_buf[ 2];
-  w0[3] = out_buf[ 3];
-  w1[0] = out_buf[ 4];
-  w1[1] = out_buf[ 5];
-  w1[2] = out_buf[ 6];
-  w1[3] = out_buf[ 7];
-  w2[0] = out_buf[ 8];
-  w2[1] = out_buf[ 9];
-  w2[2] = out_buf[10];
-  w2[3] = out_buf[11];
-  w3[0] = out_buf[12];
-  w3[1] = out_buf[13];
-  w3[2] = out_buf[14];
-  w3[3] = out_buf[15];
-
-  return out_len;
-}
diff --git a/OpenCL/m06211-pure.cl b/OpenCL/m06211-pure.cl
index 430739271..9026aa20c 100644
--- a/OpenCL/m06211-pure.cl
+++ b/OpenCL/m06211-pure.cl
@@ -18,7 +18,6 @@
 #include "inc_cipher_serpent.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 
@@ -109,7 +108,7 @@ __kernel void m06211_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06212-pure.cl b/OpenCL/m06212-pure.cl
index 18868ef7e..2c8f70332 100644
--- a/OpenCL/m06212-pure.cl
+++ b/OpenCL/m06212-pure.cl
@@ -18,7 +18,6 @@
 #include "inc_cipher_serpent.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 
@@ -109,7 +108,7 @@ __kernel void m06212_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06213-pure.cl b/OpenCL/m06213-pure.cl
index 0eb49b9ed..cb50addcc 100644
--- a/OpenCL/m06213-pure.cl
+++ b/OpenCL/m06213-pure.cl
@@ -18,7 +18,6 @@
 #include "inc_cipher_serpent.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 
@@ -109,7 +108,7 @@ __kernel void m06213_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06221-pure.cl b/OpenCL/m06221-pure.cl
index 83b4fb2fd..2b0df294b 100644
--- a/OpenCL/m06221-pure.cl
+++ b/OpenCL/m06221-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -153,7 +152,7 @@ __kernel void m06221_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06222-pure.cl b/OpenCL/m06222-pure.cl
index 861c02144..2c0ca53e8 100644
--- a/OpenCL/m06222-pure.cl
+++ b/OpenCL/m06222-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -153,7 +152,7 @@ __kernel void m06222_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06223-pure.cl b/OpenCL/m06223-pure.cl
index 08562cec3..2b05f0474 100644
--- a/OpenCL/m06223-pure.cl
+++ b/OpenCL/m06223-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -153,7 +152,7 @@ __kernel void m06223_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06231-pure.cl b/OpenCL/m06231-pure.cl
index df7300825..79412a7c6 100644
--- a/OpenCL/m06231-pure.cl
+++ b/OpenCL/m06231-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -195,7 +194,7 @@ __kernel void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06232-pure.cl b/OpenCL/m06232-pure.cl
index c39de9d67..535558087 100644
--- a/OpenCL/m06232-pure.cl
+++ b/OpenCL/m06232-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -195,7 +194,7 @@ __kernel void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m06233-pure.cl b/OpenCL/m06233-pure.cl
index 3a53eb600..6cb2a2f6e 100644
--- a/OpenCL/m06233-pure.cl
+++ b/OpenCL/m06233-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -195,7 +194,7 @@ __kernel void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m13751-pure.cl b/OpenCL/m13751-pure.cl
index 9468a6403..217502be9 100644
--- a/OpenCL/m13751-pure.cl
+++ b/OpenCL/m13751-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -117,7 +116,7 @@ __kernel void m13751_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m13752-pure.cl b/OpenCL/m13752-pure.cl
index dcd1957c1..92dc80413 100644
--- a/OpenCL/m13752-pure.cl
+++ b/OpenCL/m13752-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -117,7 +116,7 @@ __kernel void m13752_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m13753-pure.cl b/OpenCL/m13753-pure.cl
index e171e44ed..5947cc52d 100644
--- a/OpenCL/m13753-pure.cl
+++ b/OpenCL/m13753-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -117,7 +116,7 @@ __kernel void m13753_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m13771-pure.cl b/OpenCL/m13771-pure.cl
index 5c0e5134a..98857d936 100644
--- a/OpenCL/m13771-pure.cl
+++ b/OpenCL/m13771-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -178,7 +177,7 @@ __kernel void m13771_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m13772-pure.cl b/OpenCL/m13772-pure.cl
index cce25064d..a21607f8d 100644
--- a/OpenCL/m13772-pure.cl
+++ b/OpenCL/m13772-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -178,7 +177,7 @@ __kernel void m13772_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
diff --git a/OpenCL/m13773-pure.cl b/OpenCL/m13773-pure.cl
index 96d7c0595..fbea23e3f 100644
--- a/OpenCL/m13773-pure.cl
+++ b/OpenCL/m13773-pure.cl
@@ -19,7 +19,6 @@
 #include "inc_cipher_kuznyechik.cl"
 
 #include "inc_truecrypt_keyfile.cl"
-#include "inc_truecrypt_keyboard.cl"
 #include "inc_truecrypt_crc32.cl"
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
@@ -178,7 +177,7 @@ __kernel void m13773_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
 
   const u32 pw_len = pws[gid].pw_len;
 
-  keyboard_map (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
+  execute_keyboard_layout_mapping (w0, w1, w2, w3, pw_len, s_keyboard_layout_mapping_buf, keyboard_layout_mapping_cnt);
 
   w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
   w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);