Support multi-byte characters for TC/VC keyboard layout mapping tables

OpenCL/inc_truecrypt_keyboard.cl

@@ -1,22 +1,143 @@
-DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout)
+DECLSPEC int find_map (const u32 search, const int search_len, __local kb_layout_map_t *s_kb_layout_map, const int kb_layout_map_cnt)
 {
-  w[0] = (s_keyboard_layout[(w[0] >>  0) & 0xff] <<  0)
-       | (s_keyboard_layout[(w[0] >>  8) & 0xff] <<  8)
-       | (s_keyboard_layout[(w[0] >> 16) & 0xff] << 16)
-       | (s_keyboard_layout[(w[0] >> 24) & 0xff] << 24);
+  for (int idx = 0; idx < kb_layout_map_cnt; idx++)
+  {
+    const u32 src_char = s_kb_layout_map[idx].src_char;
+    const int src_len  = s_kb_layout_map[idx].src_len;
 
-  w[1] = (s_keyboard_layout[(w[1] >>  0) & 0xff] <<  0)
-       | (s_keyboard_layout[(w[1] >>  8) & 0xff] <<  8)
-       | (s_keyboard_layout[(w[1] >> 16) & 0xff] << 16)
-       | (s_keyboard_layout[(w[1] >> 24) & 0xff] << 24);
+    if (src_len == search_len)
+    {
+      const u32 mask = 0xffffffff >> ((4 - search_len) * 8);
 
-  w[2] = (s_keyboard_layout[(w[2] >>  0) & 0xff] <<  0)
-       | (s_keyboard_layout[(w[2] >>  8) & 0xff] <<  8)
-       | (s_keyboard_layout[(w[2] >> 16) & 0xff] << 16)
-       | (s_keyboard_layout[(w[2] >> 24) & 0xff] << 24);
+      if ((src_char & mask) == (search & mask)) return idx;
+    }
+  }
 
-  w[3] = (s_keyboard_layout[(w[3] >>  0) & 0xff] <<  0)
-       | (s_keyboard_layout[(w[3] >>  8) & 0xff] <<  8)
-       | (s_keyboard_layout[(w[3] >> 16) & 0xff] << 16)
-       | (s_keyboard_layout[(w[3] >> 24) & 0xff] << 24);
+  return -1;
+}
+
+DECLSPEC int keyboard_map (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int pw_len, __local kb_layout_map_t *s_kb_layout_map, const int kb_layout_map_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_map (src, src_len, s_kb_layout_map, kb_layout_map_cnt);
+
+      if (idx == -1) continue;
+
+      u32 dst_char = s_kb_layout_map[idx].dst_char;
+      int dst_len  = s_kb_layout_map[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;
 }

OpenCL/inc_types.cl

@@ -1303,14 +1303,25 @@ typedef struct krb5asrep
 
 } krb5asrep_t;
 
+typedef struct kb_layout_map
+{
+  u32 src_char;
+  int src_len;
+  u32 dst_char;
+  int dst_len;
+
+} kb_layout_map_t;
+
 typedef struct tc
 {
   u32 salt_buf[32];
   u32 data_buf[112];
   u32 keyfile_buf[16];
-  u32 keyboard_layout[256];
   u32 signature;
 
+  kb_layout_map_t kb_layout_map[256];
+  int             kb_layout_map_cnt;
+
 } tc_t;
 
 typedef struct pbkdf2_md5

OpenCL/m06211-pure.cl

@@ -68,11 +68,13 @@ __kernel void m06211_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -105,10 +107,9 @@ __kernel void m06211_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06212-pure.cl

@@ -68,11 +68,13 @@ __kernel void m06212_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -105,10 +107,9 @@ __kernel void m06212_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06213-pure.cl

@@ -68,11 +68,13 @@ __kernel void m06213_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -105,10 +107,9 @@ __kernel void m06213_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06221-pure.cl

@@ -92,11 +92,13 @@ __kernel void m06221_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -149,14 +151,9 @@ __kernel void m06221_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
   w7[2] = pws[gid].i[30];
   w7[3] = pws[gid].i[31];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
-  keyboard_map (w4, s_keyboard_layout);
-  keyboard_map (w5, s_keyboard_layout);
-  keyboard_map (w6, s_keyboard_layout);
-  keyboard_map (w7, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06222-pure.cl

@@ -92,11 +92,13 @@ __kernel void m06222_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -149,14 +151,9 @@ __kernel void m06222_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
   w7[2] = pws[gid].i[30];
   w7[3] = pws[gid].i[31];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
-  keyboard_map (w4, s_keyboard_layout);
-  keyboard_map (w5, s_keyboard_layout);
-  keyboard_map (w6, s_keyboard_layout);
-  keyboard_map (w7, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06223-pure.cl

@@ -92,11 +92,13 @@ __kernel void m06223_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -149,14 +151,9 @@ __kernel void m06223_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
   w7[2] = pws[gid].i[30];
   w7[3] = pws[gid].i[31];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
-  keyboard_map (w4, s_keyboard_layout);
-  keyboard_map (w5, s_keyboard_layout);
-  keyboard_map (w6, s_keyboard_layout);
-  keyboard_map (w7, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06231-pure.cl

@@ -130,11 +130,13 @@ __kernel void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   __local u32 s_Ch[8][256];
@@ -191,10 +193,9 @@ __kernel void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06232-pure.cl

@@ -130,11 +130,13 @@ __kernel void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   __local u32 s_Ch[8][256];
@@ -191,10 +193,9 @@ __kernel void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m06233-pure.cl

@@ -130,11 +130,13 @@ __kernel void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   __local u32 s_Ch[8][256];
@@ -191,10 +193,9 @@ __kernel void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m13751-pure.cl

@@ -76,11 +76,13 @@ __kernel void m13751_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -113,10 +115,9 @@ __kernel void m13751_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m13752-pure.cl

@@ -76,11 +76,13 @@ __kernel void m13752_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -113,10 +115,9 @@ __kernel void m13752_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m13753-pure.cl

@@ -76,11 +76,13 @@ __kernel void m13753_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -113,10 +115,9 @@ __kernel void m13753_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m13771-pure.cl

@@ -115,11 +115,13 @@ __kernel void m13771_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
   const u64 lid = get_local_id (0);
   const u64 lsz = get_local_size (0);
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   #ifdef REAL_SHM
@@ -174,10 +176,9 @@ __kernel void m13771_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m13772-pure.cl

@@ -115,11 +115,13 @@ __kernel void m13772_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
   const u64 lid = get_local_id (0);
   const u64 lsz = get_local_size (0);
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   #ifdef REAL_SHM
@@ -174,10 +176,9 @@ __kernel void m13772_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);

OpenCL/m13773-pure.cl

@@ -115,11 +115,13 @@ __kernel void m13773_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
   const u64 lid = get_local_id (0);
   const u64 lsz = get_local_size (0);
 
-  __local u32 s_keyboard_layout[256];
+  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+
+  __local kb_layout_map_t s_kb_layout_map[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
+    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
   }
 
   #ifdef REAL_SHM
@@ -174,10 +176,9 @@ __kernel void m13773_init (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
   w3[2] = pws[gid].i[14];
   w3[3] = pws[gid].i[15];
 
-  keyboard_map (w0, s_keyboard_layout);
-  keyboard_map (w1, s_keyboard_layout);
-  keyboard_map (w2, s_keyboard_layout);
-  keyboard_map (w3, s_keyboard_layout);
+  const u32 pw_len = pws[gid].pw_len;
+
+  keyboard_map (w0, w1, w2, w3, pw_len, s_kb_layout_map, kb_layout_map_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]);