Fix more append_* functions in kernels

[hashcat.git] / OpenCL / m06233.cl

diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl
index fb67506..864dcda 100644 (file)
--- a/OpenCL/m06233.cl
+++ b/OpenCL/m06233.cl
@@ -1559,22 +1559,22 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06233_init (__gl
 
   u32 salt_buf1[16];
 
-  salt_buf1[ 0] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 0]);
-  salt_buf1[ 1] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 1]);
-  salt_buf1[ 2] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 2]);
-  salt_buf1[ 3] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 3]);
-  salt_buf1[ 4] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 4]);
-  salt_buf1[ 5] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 5]);
-  salt_buf1[ 6] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 6]);
-  salt_buf1[ 7] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 7]);
-  salt_buf1[ 8] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 8]);
-  salt_buf1[ 9] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 9]);
-  salt_buf1[10] = swap_workaround (esalt_bufs[salt_pos].salt_buf[10]);
-  salt_buf1[11] = swap_workaround (esalt_bufs[salt_pos].salt_buf[11]);
-  salt_buf1[12] = swap_workaround (esalt_bufs[salt_pos].salt_buf[12]);
-  salt_buf1[13] = swap_workaround (esalt_bufs[salt_pos].salt_buf[13]);
-  salt_buf1[14] = swap_workaround (esalt_bufs[salt_pos].salt_buf[14]);
-  salt_buf1[15] = swap_workaround (esalt_bufs[salt_pos].salt_buf[15]);
+  salt_buf1[ 0] = swap32 (esalt_bufs[salt_pos].salt_buf[ 0]);
+  salt_buf1[ 1] = swap32 (esalt_bufs[salt_pos].salt_buf[ 1]);
+  salt_buf1[ 2] = swap32 (esalt_bufs[salt_pos].salt_buf[ 2]);
+  salt_buf1[ 3] = swap32 (esalt_bufs[salt_pos].salt_buf[ 3]);
+  salt_buf1[ 4] = swap32 (esalt_bufs[salt_pos].salt_buf[ 4]);
+  salt_buf1[ 5] = swap32 (esalt_bufs[salt_pos].salt_buf[ 5]);
+  salt_buf1[ 6] = swap32 (esalt_bufs[salt_pos].salt_buf[ 6]);
+  salt_buf1[ 7] = swap32 (esalt_bufs[salt_pos].salt_buf[ 7]);
+  salt_buf1[ 8] = swap32 (esalt_bufs[salt_pos].salt_buf[ 8]);
+  salt_buf1[ 9] = swap32 (esalt_bufs[salt_pos].salt_buf[ 9]);
+  salt_buf1[10] = swap32 (esalt_bufs[salt_pos].salt_buf[10]);
+  salt_buf1[11] = swap32 (esalt_bufs[salt_pos].salt_buf[11]);
+  salt_buf1[12] = swap32 (esalt_bufs[salt_pos].salt_buf[12]);
+  salt_buf1[13] = swap32 (esalt_bufs[salt_pos].salt_buf[13]);
+  salt_buf1[14] = swap32 (esalt_bufs[salt_pos].salt_buf[14]);
+  salt_buf1[15] = swap32 (esalt_bufs[salt_pos].salt_buf[15]);
 
   u32 salt_buf2[16];
 
@@ -1599,22 +1599,22 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06233_init (__gl
 
   u32 w[16];
 
-  w[ 0] = swap_workaround (w0[0]);
-  w[ 1] = swap_workaround (w0[1]);
-  w[ 2] = swap_workaround (w0[2]);
-  w[ 3] = swap_workaround (w0[3]);
-  w[ 4] = swap_workaround (w1[0]);
-  w[ 5] = swap_workaround (w1[1]);
-  w[ 6] = swap_workaround (w1[2]);
-  w[ 7] = swap_workaround (w1[3]);
-  w[ 8] = swap_workaround (w2[0]);
-  w[ 9] = swap_workaround (w2[1]);
-  w[10] = swap_workaround (w2[2]);
-  w[11] = swap_workaround (w2[3]);
-  w[12] = swap_workaround (w3[0]);
-  w[13] = swap_workaround (w3[1]);
-  w[14] = swap_workaround (w3[2]);
-  w[15] = swap_workaround (w3[3]);
+  w[ 0] = swap32 (w0[0]);
+  w[ 1] = swap32 (w0[1]);
+  w[ 2] = swap32 (w0[2]);
+  w[ 3] = swap32 (w0[3]);
+  w[ 4] = swap32 (w1[0]);
+  w[ 5] = swap32 (w1[1]);
+  w[ 6] = swap32 (w1[2]);
+  w[ 7] = swap32 (w1[3]);
+  w[ 8] = swap32 (w2[0]);
+  w[ 9] = swap32 (w2[1]);
+  w[10] = swap32 (w2[2]);
+  w[11] = swap32 (w2[3]);
+  w[12] = swap32 (w3[0]);
+  w[13] = swap32 (w3[1]);
+  w[14] = swap32 (w3[2]);
+  w[15] = swap32 (w3[3]);
 
   u32 ipad[16];
   u32 opad[16];
@@ -1914,25 +1914,25 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06233_comp (__gl
 
   u32 ukey1[8];
 
-  ukey1[0] = swap_workaround (tmps[gid].out[ 0]);
-  ukey1[1] = swap_workaround (tmps[gid].out[ 1]);
-  ukey1[2] = swap_workaround (tmps[gid].out[ 2]);
-  ukey1[3] = swap_workaround (tmps[gid].out[ 3]);
-  ukey1[4] = swap_workaround (tmps[gid].out[ 4]);
-  ukey1[5] = swap_workaround (tmps[gid].out[ 5]);
-  ukey1[6] = swap_workaround (tmps[gid].out[ 6]);
-  ukey1[7] = swap_workaround (tmps[gid].out[ 7]);
+  ukey1[0] = swap32 (tmps[gid].out[ 0]);
+  ukey1[1] = swap32 (tmps[gid].out[ 1]);
+  ukey1[2] = swap32 (tmps[gid].out[ 2]);
+  ukey1[3] = swap32 (tmps[gid].out[ 3]);
+  ukey1[4] = swap32 (tmps[gid].out[ 4]);
+  ukey1[5] = swap32 (tmps[gid].out[ 5]);
+  ukey1[6] = swap32 (tmps[gid].out[ 6]);
+  ukey1[7] = swap32 (tmps[gid].out[ 7]);
 
   u32 ukey2[8];
 
-  ukey2[0] = swap_workaround (tmps[gid].out[ 8]);
-  ukey2[1] = swap_workaround (tmps[gid].out[ 9]);
-  ukey2[2] = swap_workaround (tmps[gid].out[10]);
-  ukey2[3] = swap_workaround (tmps[gid].out[11]);
-  ukey2[4] = swap_workaround (tmps[gid].out[12]);
-  ukey2[5] = swap_workaround (tmps[gid].out[13]);
-  ukey2[6] = swap_workaround (tmps[gid].out[14]);
-  ukey2[7] = swap_workaround (tmps[gid].out[15]);
+  ukey2[0] = swap32 (tmps[gid].out[ 8]);
+  ukey2[1] = swap32 (tmps[gid].out[ 9]);
+  ukey2[2] = swap32 (tmps[gid].out[10]);
+  ukey2[3] = swap32 (tmps[gid].out[11]);
+  ukey2[4] = swap32 (tmps[gid].out[12]);
+  ukey2[5] = swap32 (tmps[gid].out[13]);
+  ukey2[6] = swap32 (tmps[gid].out[14]);
+  ukey2[7] = swap32 (tmps[gid].out[15]);
 
   u32 data[4];
 
@@ -1993,25 +1993,25 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06233_comp (__gl
 
   u32 ukey3[8];
 
-  ukey3[0] = swap_workaround (tmps[gid].out[16]);
-  ukey3[1] = swap_workaround (tmps[gid].out[17]);
-  ukey3[2] = swap_workaround (tmps[gid].out[18]);
-  ukey3[3] = swap_workaround (tmps[gid].out[19]);
-  ukey3[4] = swap_workaround (tmps[gid].out[20]);
-  ukey3[5] = swap_workaround (tmps[gid].out[21]);
-  ukey3[6] = swap_workaround (tmps[gid].out[22]);
-  ukey3[7] = swap_workaround (tmps[gid].out[23]);
+  ukey3[0] = swap32 (tmps[gid].out[16]);
+  ukey3[1] = swap32 (tmps[gid].out[17]);
+  ukey3[2] = swap32 (tmps[gid].out[18]);
+  ukey3[3] = swap32 (tmps[gid].out[19]);
+  ukey3[4] = swap32 (tmps[gid].out[20]);
+  ukey3[5] = swap32 (tmps[gid].out[21]);
+  ukey3[6] = swap32 (tmps[gid].out[22]);
+  ukey3[7] = swap32 (tmps[gid].out[23]);
 
   u32 ukey4[8];
 
-  ukey4[0] = swap_workaround (tmps[gid].out[24]);
-  ukey4[1] = swap_workaround (tmps[gid].out[25]);
-  ukey4[2] = swap_workaround (tmps[gid].out[26]);
-  ukey4[3] = swap_workaround (tmps[gid].out[27]);
-  ukey4[4] = swap_workaround (tmps[gid].out[28]);
-  ukey4[5] = swap_workaround (tmps[gid].out[29]);
-  ukey4[6] = swap_workaround (tmps[gid].out[30]);
-  ukey4[7] = swap_workaround (tmps[gid].out[31]);
+  ukey4[0] = swap32 (tmps[gid].out[24]);
+  ukey4[1] = swap32 (tmps[gid].out[25]);
+  ukey4[2] = swap32 (tmps[gid].out[26]);
+  ukey4[3] = swap32 (tmps[gid].out[27]);
+  ukey4[4] = swap32 (tmps[gid].out[28]);
+  ukey4[5] = swap32 (tmps[gid].out[29]);
+  ukey4[6] = swap32 (tmps[gid].out[30]);
+  ukey4[7] = swap32 (tmps[gid].out[31]);
 
   {
     tmp[0] = data[0];
@@ -2066,25 +2066,25 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06233_comp (__gl
 
   u32 ukey5[8];
 
-  ukey5[0] = swap_workaround (tmps[gid].out[32]);
-  ukey5[1] = swap_workaround (tmps[gid].out[33]);
-  ukey5[2] = swap_workaround (tmps[gid].out[34]);
-  ukey5[3] = swap_workaround (tmps[gid].out[35]);
-  ukey5[4] = swap_workaround (tmps[gid].out[36]);
-  ukey5[5] = swap_workaround (tmps[gid].out[37]);
-  ukey5[6] = swap_workaround (tmps[gid].out[38]);
-  ukey5[7] = swap_workaround (tmps[gid].out[39]);
+  ukey5[0] = swap32 (tmps[gid].out[32]);
+  ukey5[1] = swap32 (tmps[gid].out[33]);
+  ukey5[2] = swap32 (tmps[gid].out[34]);
+  ukey5[3] = swap32 (tmps[gid].out[35]);
+  ukey5[4] = swap32 (tmps[gid].out[36]);
+  ukey5[5] = swap32 (tmps[gid].out[37]);
+  ukey5[6] = swap32 (tmps[gid].out[38]);
+  ukey5[7] = swap32 (tmps[gid].out[39]);
 
   u32 ukey6[8];
 
-  ukey6[0] = swap_workaround (tmps[gid].out[40]);
-  ukey6[1] = swap_workaround (tmps[gid].out[41]);
-  ukey6[2] = swap_workaround (tmps[gid].out[42]);
-  ukey6[3] = swap_workaround (tmps[gid].out[43]);
-  ukey6[4] = swap_workaround (tmps[gid].out[44]);
-  ukey6[5] = swap_workaround (tmps[gid].out[45]);
-  ukey6[6] = swap_workaround (tmps[gid].out[46]);
-  ukey6[7] = swap_workaround (tmps[gid].out[47]);
+  ukey6[0] = swap32 (tmps[gid].out[40]);
+  ukey6[1] = swap32 (tmps[gid].out[41]);
+  ukey6[2] = swap32 (tmps[gid].out[42]);
+  ukey6[3] = swap32 (tmps[gid].out[43]);
+  ukey6[4] = swap32 (tmps[gid].out[44]);
+  ukey6[5] = swap32 (tmps[gid].out[45]);
+  ukey6[6] = swap32 (tmps[gid].out[46]);
+  ukey6[7] = swap32 (tmps[gid].out[47]);
 
   {
     tmp[0] = data[0];