Fix more append_* functions in kernels

This commit is contained in:
jsteube
2015-12-15 16:50:21 +01:00
parent 531910f8c0
commit 2283d5c843
171 changed files with 7401 additions and 7356 deletions

View File

@@ -1559,22 +1559,22 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_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))) m06232_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))) m06232_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))) m06232_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];