Prepare to use --keyboard-layout-mapping for algorithms other than TC/VC

OpenCL/inc_truecrypt_keyboard.cl

@@ -1,9 +1,9 @@
-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)
+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 < kb_layout_map_cnt; idx++)
+  for (int idx = 0; idx < keyboard_layout_mapping_cnt; idx++)
   {
-    const u32 src_char = s_kb_layout_map[idx].src_char;
-    const int src_len  = s_kb_layout_map[idx].src_len;
+    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)
     {
@@ -16,7 +16,7 @@ DECLSPEC int find_map (const u32 search, const int search_len, __local kb_layout
   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)
+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 };
 
@@ -76,12 +76,12 @@ DECLSPEC int keyboard_map (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int
 
     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);
+      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_kb_layout_map[idx].dst_char;
-      int dst_len  = s_kb_layout_map[idx].dst_len;
+      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)
       {

OpenCL/inc_types.cl

@@ -1303,14 +1303,14 @@ typedef struct krb5asrep
 
 } krb5asrep_t;
 
-typedef struct kb_layout_map
+typedef struct keyboard_layout_mapping
 {
   u32 src_char;
   int src_len;
   u32 dst_char;
   int dst_len;
 
-} kb_layout_map_t;
+} keyboard_layout_mapping_t;
 
 typedef struct tc
 {
@@ -1319,8 +1319,8 @@ typedef struct tc
   u32 keyfile_buf[16];
   u32 signature;
 
-  kb_layout_map_t kb_layout_map[256];
-  int             kb_layout_map_cnt;
+  keyboard_layout_mapping_t keyboard_layout_mapping_buf[256];
+  int                       keyboard_layout_mapping_cnt;
 
 } tc_t;
 

OpenCL/m06211-pure.cl

@@ -68,13 +68,13 @@ __kernel void m06211_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -109,7 +109,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06212-pure.cl

@@ -68,13 +68,13 @@ __kernel void m06212_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -109,7 +109,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06213-pure.cl

@@ -68,13 +68,13 @@ __kernel void m06213_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -109,7 +109,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06221-pure.cl

@@ -92,13 +92,13 @@ __kernel void m06221_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -153,7 +153,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06222-pure.cl

@@ -92,13 +92,13 @@ __kernel void m06222_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -153,7 +153,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06223-pure.cl

@@ -92,13 +92,13 @@ __kernel void m06223_init (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -153,7 +153,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06231-pure.cl

@@ -130,13 +130,13 @@ __kernel void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   __local u32 s_Ch[8][256];
@@ -195,7 +195,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06232-pure.cl

@@ -130,13 +130,13 @@ __kernel void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   __local u32 s_Ch[8][256];
@@ -195,7 +195,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m06233-pure.cl

@@ -130,13 +130,13 @@ __kernel void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   __local u32 s_Ch[8][256];
@@ -195,7 +195,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m13751-pure.cl

@@ -76,13 +76,13 @@ __kernel void m13751_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -117,7 +117,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m13752-pure.cl

@@ -76,13 +76,13 @@ __kernel void m13752_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -117,7 +117,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m13753-pure.cl

@@ -76,13 +76,13 @@ __kernel void m13753_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
    * keyboard layout shared
    */
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   barrier (CLK_LOCAL_MEM_FENCE);
@@ -117,7 +117,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m13771-pure.cl

@@ -115,13 +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);
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   #ifdef REAL_SHM
@@ -178,7 +178,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m13772-pure.cl

@@ -115,13 +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);
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   #ifdef REAL_SHM
@@ -178,7 +178,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);

OpenCL/m13773-pure.cl

@@ -115,13 +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);
 
-  const int kb_layout_map_cnt = esalt_bufs[digests_offset].kb_layout_map_cnt;
+  const int keyboard_layout_mapping_cnt = esalt_bufs[digests_offset].keyboard_layout_mapping_cnt;
 
-  __local kb_layout_map_t s_kb_layout_map[256];
+  __local keyboard_layout_mapping_t s_keyboard_layout_mapping_buf[256];
 
   for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
   {
-    s_kb_layout_map[i] = esalt_bufs[digests_offset].kb_layout_map[i];
+    s_keyboard_layout_mapping_buf[i] = esalt_bufs[digests_offset].keyboard_layout_mapping_buf[i];
   }
 
   #ifdef REAL_SHM
@@ -178,7 +178,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_kb_layout_map, kb_layout_map_cnt);
+  keyboard_map (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]);