Prepare DNSSEC (NSEC3) optimized kernel for pure kernel version

OpenCL/m08300_a0-optimized.cl

@@ -77,7 +77,7 @@ __kernel void m08300_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
   domain_buf1[2] = salt_bufs[salt_pos].salt_buf_pc[ 6];
   domain_buf1[3] = 0;
 
-  const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
+  const u32 domain_len = salt_bufs[salt_pos].salt_len_pc;
 
   /**
    * loop
@@ -334,7 +334,7 @@ __kernel void m08300_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   domain_buf1[2] = salt_bufs[salt_pos].salt_buf_pc[ 6];
   domain_buf1[3] = 0;
 
-  const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
+  const u32 domain_len = salt_bufs[salt_pos].salt_len_pc;
 
   /**
    * digest

OpenCL/m08300_a1-optimized.cl

@@ -75,7 +75,7 @@ __kernel void m08300_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
   domain_buf1[2] = salt_bufs[salt_pos].salt_buf_pc[ 6];
   domain_buf1[3] = 0;
 
-  const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
+  const u32 domain_len = salt_bufs[salt_pos].salt_len_pc;
 
   /**
    * loop
@@ -392,7 +392,7 @@ __kernel void m08300_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   domain_buf1[2] = salt_bufs[salt_pos].salt_buf_pc[ 6];
   domain_buf1[3] = 0;
 
-  const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
+  const u32 domain_len = salt_bufs[salt_pos].salt_len_pc;
 
   /**
    * digest

OpenCL/m08300_a3-optimized.cl

@@ -54,7 +54,7 @@ void m08300m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
   domain_buf1[2] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 6]);
   domain_buf1[3] = 0;
 
-  const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
+  const u32 domain_len = salt_bufs[salt_pos].salt_len_pc;
 
   u32 s0[4];
   u32 s1[4];
@@ -262,7 +262,7 @@ void m08300s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
   domain_buf1[2] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 6]);
   domain_buf1[3] = 0;
 
-  const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
+  const u32 domain_len = salt_bufs[salt_pos].salt_len_pc;
 
   u32 s0[4];
   u32 s1[4];

src/interface.c

@@ -8004,7 +8004,7 @@ int nsec3_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUS
     }
   }
 
-  salt->salt_buf_pc[7] = domainbuf_len;
+  salt->salt_len_pc = domainbuf_len;
 
   // "real" salt
 
@@ -17612,7 +17612,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
 
     // domain
 
-    const u32 salt_pc_len = salt.salt_buf_pc[7]; // what a hack
+    const u32 salt_pc_len = salt.salt_len_pc;
 
     char domain_buf_c[33] = { 0 };