arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-22 08:08:10 +00:00
Code Issues Releases Wiki Activity

Fix append_* function calls

Browse Source
This commit is contained in:
jsteube 2015-12-15 13:42:37 +01:00
parent 88284108b3
commit 50f39b3563
147 changed files with 1245 additions and 1243 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1840
OpenCL/common.c
View File

File diff suppressed because it is too large Load Diff

4
OpenCL/m00010_a0.cl
View File

@ -166,7 +166,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00010_m04 (__glo
w3[2] = pw_salt_len * 8;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* md5
@ -421,7 +421,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00010_s04 (__glo
w3[2] = pw_salt_len * 8;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* md5

4
OpenCL/m00010_a1.cl
View File

@ -198,7 +198,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00010_m04 (__glo
w3[2] = pw_salt_len * 8;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* md5
@ -487,7 +487,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00010_s04 (__glo
w3[2] = pw_salt_len * 8;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* md5

4
OpenCL/m00020_a0.cl
View File

@ -143,7 +143,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00020_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[2] = out_salt_len * 8;
@ -377,7 +377,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00020_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[2] = out_salt_len * 8;

4
OpenCL/m00020_a1.cl
View File

@ -189,7 +189,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00020_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;
@ -473,7 +473,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00020_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;

4
OpenCL/m00030_a0.cl
View File

@ -173,7 +173,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00030_m04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[2] = out_salt_len * 8;
@ -433,7 +433,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00030_s04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[2] = out_salt_len * 8;

4
OpenCL/m00030_a1.cl
View File

@ -219,7 +219,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00030_m04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;
@ -527,7 +527,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00030_s04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;

4
OpenCL/m00040_a0.cl
View File

@ -135,7 +135,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00040_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[2] = out_salt_len * 8;
@ -359,7 +359,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00040_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[2] = out_salt_len * 8;

4
OpenCL/m00040_a1.cl
View File

@ -183,7 +183,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00040_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;
@ -455,7 +455,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00040_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;

4
OpenCL/m00060_a0.cl
View File

@ -330,7 +330,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00060_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = w0[0];
w0_t[1] = w0[1];
@ -505,7 +505,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00060_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = w0[0];
w0_t[1] = w0[1];

4
OpenCL/m00060_a1.cl
View File

@ -382,7 +382,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00060_m04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = w0[0];
w0_t[1] = w0[1];
@ -611,7 +611,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00060_s04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = w0[0];
w0_t[1] = w0[1];

4
OpenCL/m00060_a3.cl
View File

@ -281,7 +281,7 @@ static void m00060m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
w0[0] = w0l | w0r;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = w0[0];
w0_t[1] = w0[1];
@ -401,7 +401,7 @@ static void m00060s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
w0[0] = w0l | w0r;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = w0[0];
w0_t[1] = w0[1];

4
OpenCL/m00100_a0.cl
View File

@ -90,7 +90,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00100_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1
@ -319,7 +319,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00100_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1

8
OpenCL/m00100_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00100_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -113,7 +113,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00100_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -335,7 +335,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00100_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -398,7 +398,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00100_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m00110_a0.cl
View File

@ -166,7 +166,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00110_m04 (__glo
w3[2] |= s3[2];
w3[3] |= s3[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha1
@ -471,7 +471,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00110_s04 (__glo
w3[2] |= s3[2];
w3[3] |= s3[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha1

4
OpenCL/m00110_a1.cl
View File

@ -198,7 +198,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00110_m04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha1
@ -537,7 +537,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00110_s04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha1

4
OpenCL/m00120_a0.cl
View File

@ -143,7 +143,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[3] = out_salt_len * 8;
@ -427,7 +427,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[3] = out_salt_len * 8;

4
OpenCL/m00120_a1.cl
View File

@ -189,7 +189,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[3] = pw_salt_len * 8;
@ -521,7 +521,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00120_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[3] = pw_salt_len * 8;

4
OpenCL/m00130_a0.cl
View File

@ -173,7 +173,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00130_m04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[3] = out_salt_len * 8;
@ -485,7 +485,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00130_s04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[3] = out_salt_len * 8;

4
OpenCL/m00130_a1.cl
View File

@ -219,7 +219,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00130_m04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[3] = pw_salt_len * 8;
@ -579,7 +579,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00130_s04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[3] = pw_salt_len * 8;

4
OpenCL/m00140_a0.cl
View File

@ -137,7 +137,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[3] = out_salt_len * 8;
@ -413,7 +413,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
w3_t[3] = out_salt_len * 8;

4
OpenCL/m00140_a1.cl
View File

@ -183,7 +183,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[3] = pw_salt_len * 8;
@ -507,7 +507,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00140_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[3] = pw_salt_len * 8;

4
OpenCL/m00160_a0.cl
View File

@ -362,7 +362,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00160_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);
@ -537,7 +537,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00160_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);

4
OpenCL/m00160_a1.cl
View File

@ -414,7 +414,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00160_m04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);
@ -643,7 +643,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00160_s04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);

4
OpenCL/m00190_a0.cl
View File

@ -90,7 +90,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00190_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1
@ -330,7 +330,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00190_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1

8
OpenCL/m00190_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00190_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -113,7 +113,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00190_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -352,7 +352,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00190_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -409,7 +409,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00190_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m00300_a0.cl
View File

@ -90,7 +90,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00300_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1
@ -444,7 +444,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00300_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1

8
OpenCL/m00300_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00300_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -113,7 +113,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00300_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -460,7 +460,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00300_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -523,7 +523,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00300_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m00900_a0.cl
View File

@ -90,7 +90,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00900_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;
@ -247,7 +247,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00900_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;

4
OpenCL/m00900_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00900_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -265,7 +265,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m00900_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m01000_a0.cl
View File

@ -90,7 +90,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];
@ -257,7 +257,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];

4
OpenCL/m01000_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -277,7 +277,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m01100_a0.cl
View File

@ -117,7 +117,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01100_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];
@ -387,7 +387,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01100_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];

4
OpenCL/m01100_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01100_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -380,7 +380,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01100_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m01400_a0.cl
View File

@ -90,7 +90,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01400_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* SHA256
@ -288,7 +288,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01400_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* SHA256

8
OpenCL/m01400_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01400_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -113,7 +113,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01400_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -304,7 +304,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01400_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -361,7 +361,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01400_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m01410_a0.cl
View File

@ -166,7 +166,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01410_m04 (__glo
w3[2] |= s3[2];
w3[3] |= s3[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha256
@ -440,7 +440,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01410_s04 (__glo
w3[2] |= s3[2];
w3[3] |= s3[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha256

4
OpenCL/m01410_a1.cl
View File

@ -192,7 +192,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01410_m04 (__glo
w3[2] = wordl3[2] | wordr3[2] | s3[2];
w3[3] = wordl3[3] | wordr3[3] | s3[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha256
@ -494,7 +494,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01410_s04 (__glo
w3[2] = wordl3[2] | wordr3[2] | s3[2];
w3[3] = wordl3[3] | wordr3[3] | s3[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha256

4
OpenCL/m01420_a0.cl
View File

@ -127,7 +127,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01420_m04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha256
@ -362,7 +362,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01420_s04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha256

4
OpenCL/m01420_a1.cl
View File

@ -167,7 +167,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01420_m04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha256
@ -444,7 +444,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01420_s04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha256

4
OpenCL/m01430_a0.cl
View File

@ -171,7 +171,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01430_m04 (__glo
w3_t2[2] |= s3[2];
w3_t2[3] |= s3[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
/**
* sha256
@ -450,7 +450,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01430_s04 (__glo
w3_t2[2] |= s3[2];
w3_t2[3] |= s3[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
/**
* sha256

4
OpenCL/m01430_a1.cl
View File

@ -211,7 +211,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01430_m04 (__glo
w3_t2[2] |= s3[2];
w3_t2[3] |= s3[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
/**
* sha256
@ -532,7 +532,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01430_s04 (__glo
w3_t2[2] |= s3[2];
w3_t2[3] |= s3[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
/**
* sha256

4
OpenCL/m01440_a0.cl
View File

@ -129,7 +129,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01440_m04 (__glo
w1_t2[2] |= salt_buf1[2];
w1_t2[3] |= salt_buf1[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
/**
* sha256
@ -366,7 +366,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01440_s04 (__glo
w1_t2[2] |= salt_buf1[2];
w1_t2[3] |= salt_buf1[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, out_salt_len);
/**
* sha256

4
OpenCL/m01440_a1.cl
View File

@ -175,7 +175,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01440_m04 (__glo
w1_t2[2] |= salt_buf1[2];
w1_t2[3] |= salt_buf1[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
/**
* sha256
@ -460,7 +460,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01440_s04 (__glo
w1_t2[2] |= salt_buf1[2];
w1_t2[3] |= salt_buf1[3];
append_0x80_4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
append_0x80_4x4 (w0_t2, w1_t2, w2_t2, w3_t2, pw_salt_len);
/**
* sha256

4
OpenCL/m01460_a0.cl
View File

@ -352,7 +352,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01460_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);
@ -527,7 +527,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01460_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);

4
OpenCL/m01460_a1.cl
View File

@ -404,7 +404,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01460_m04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);
@ -633,7 +633,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01460_s04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);

6
OpenCL/m01600.cl
View File

@ -1750,7 +1750,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01600_init (__gl
block_len += pw_len;
append_0x80_4 (block0, block1, block2, block3, block_len);
append_0x80_4x4 (block0, block1, block2, block3, block_len);
block3[2] = block_len * 8;
@ -1821,7 +1821,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01600_init (__gl
block_len++;
}
append_0x80_4 (block0, block1, block2, block3, block_len);
append_0x80_4x4 (block0, block1, block2, block3, block_len);
block3[2] = block_len * 8;
@ -1864,7 +1864,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01600_loop (__gl
w0_x80[2] = w0[2];
w0_x80[3] = w0[3];
append_0x80_1 (w0_x80, pw_len);
append_0x80_1x4 (w0_x80, pw_len);
/**
* salt

4
OpenCL/m01700_a0.cl
View File

@ -211,7 +211,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01700_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* SHA512
@ -350,7 +350,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01700_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* SHA512

8
OpenCL/m01700_a1.cl
View File

@ -189,7 +189,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01700_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -234,7 +234,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01700_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -366,7 +366,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01700_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -423,7 +423,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01700_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m01710_a0.cl
View File

@ -287,7 +287,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01710_m04 (__glo
w3[2] |= s3[2];
w3[3] |= s3[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha512
@ -502,7 +502,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01710_s04 (__glo
w3[2] |= s3[2];
w3[3] |= s3[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha512

4
OpenCL/m01710_a1.cl
View File

@ -313,7 +313,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01710_m04 (__glo
w3[2] = wordl3[2] | wordr3[2] | s3[2];
w3[3] = wordl3[3] | wordr3[3] | s3[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha512
@ -556,7 +556,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01710_s04 (__glo
w3[2] = wordl3[2] | wordr3[2] | s3[2];
w3[3] = wordl3[3] | wordr3[3] | s3[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha512

4
OpenCL/m01720_a0.cl
View File

@ -248,7 +248,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01720_m04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha512
@ -424,7 +424,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01720_s04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, out_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, out_salt_len);
/**
* sha512

4
OpenCL/m01720_a1.cl
View File

@ -288,7 +288,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01720_m04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha512
@ -506,7 +506,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01720_s04 (__glo
w1[2] |= salt_buf1[2];
w1[3] |= salt_buf1[3];
append_0x80_4 (w0, w1, w2, w3, pw_salt_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_salt_len);
/**
* sha512

4
OpenCL/m01730_a0.cl
View File

@ -292,7 +292,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01730_m04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
/**
* sha512
@ -507,7 +507,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01730_s04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
/**
* sha512

4
OpenCL/m01730_a1.cl
View File

@ -332,7 +332,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01730_m04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
/**
* sha512
@ -589,7 +589,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01730_s04 (__glo
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
/**
* sha512

4
OpenCL/m01740_a0.cl
View File

@ -250,7 +250,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01740_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
/**
* sha512
@ -423,7 +423,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01740_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, out_salt_len);
/**
* sha512

4
OpenCL/m01740_a1.cl
View File

@ -296,7 +296,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01740_m04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
/**
* sha512
@ -517,7 +517,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01740_s04 (__glo
w1_t[2] |= salt_buf1[2];
w1_t[3] |= salt_buf1[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
/**
* sha512

4
OpenCL/m01760_a0.cl
View File

@ -383,7 +383,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01760_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);
@ -559,7 +559,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01760_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);

4
OpenCL/m01760_a1.cl
View File

@ -435,7 +435,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01760_m04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);
@ -665,7 +665,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01760_s04 (__glo
w3[2] = 0;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, pw_len);
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
w0_t[0] = swap_workaround (w0[0]);
w0_t[1] = swap_workaround (w0[1]);

8
OpenCL/m02100.cl
View File

@ -373,7 +373,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02100_init (__gl
* generate dcc
*/
append_0x80_1 (w0, pw_len);
append_0x80_1x4 (w0, pw_len);
make_unicode (w0, w0, w1);
@ -405,7 +405,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02100_init (__gl
w3[2] = (16 + salt_len) * 8;
w3[3] = 0;
append_0x80_4 (w0, w1, w2, w3, 16 + salt_len);
append_0x80_4x4 (w0, w1, w2, w3, 16 + salt_len);
digest_md4[0] = MD4M_A;
digest_md4[1] = MD4M_B;
@ -473,8 +473,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02100_init (__gl
w3[2] = 0;
w3[3] = (64 + salt_len + 4) * 8;
append_0x01_4 (w0, w1, w2, w3, salt_len + 3);
append_0x80_4 (w0, w1, w2, w3, salt_len + 4);
append_0x01_4x4 (w0, w1, w2, w3, salt_len + 3);
append_0x80_4x4 (w0, w1, w2, w3, salt_len + 4);
w0[0] = swap_workaround (w0[0]);
w0[1] = swap_workaround (w0[1]);

6
OpenCL/m02500.cl
View File

@ -545,11 +545,11 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02500_init (__gl
w3[3] = 0;
if (j == 1)
append_0x01_3 (w0, w1, w2, salt_len + 3);
append_0x01_3x4 (w0, w1, w2, salt_len + 3);
else
append_0x02_3 (w0, w1, w2, salt_len + 3);
append_0x02_3x4 (w0, w1, w2, salt_len + 3);
append_0x80_3 (w0, w1, w2, salt_len + 4);
append_0x80_3x4 (w0, w1, w2, salt_len + 4);
w0[0] = swap_workaround (w0[0]);
w0[1] = swap_workaround (w0[1]);

4
OpenCL/m02610_a0.cl
View File

@ -152,7 +152,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02610_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;
@ -480,7 +480,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02610_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;

4
OpenCL/m02610_a1.cl
View File

@ -78,7 +78,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02610_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -446,7 +446,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02610_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m02710_a0.cl
View File

@ -154,7 +154,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02710_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;
@ -567,7 +567,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02710_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;

4
OpenCL/m02710_a1.cl
View File

@ -82,7 +82,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02710_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -535,7 +535,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02710_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m02810_a0.cl
View File

@ -157,7 +157,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02810_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;
@ -569,7 +569,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02810_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;

4
OpenCL/m02810_a1.cl
View File

@ -78,7 +78,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02810_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -530,7 +530,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02810_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m03710_a0.cl
View File

@ -171,7 +171,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;
@ -553,7 +553,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;

4
OpenCL/m03710_a1.cl
View File

@ -78,7 +78,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -500,7 +500,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m03800_a0.cl
View File

@ -262,7 +262,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03800_m04 (__glo
const u32 pw_salt_len = salt_len + out_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;
@ -613,7 +613,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03800_s04 (__glo
const u32 pw_salt_len = salt_len + out_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;

4
OpenCL/m03800_a1.cl
View File

@ -313,7 +313,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03800_m04 (__glo
const u32 pw_salt_len = salt_len + pw_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;
@ -717,7 +717,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03800_s04 (__glo
const u32 pw_salt_len = salt_len + pw_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;

4
OpenCL/m03800_a3.cl
View File

@ -184,7 +184,7 @@ static void m03800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;
@ -447,7 +447,7 @@ static void m03800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
w3_t[2] |= s3[2];
w3_t[3] |= s3[3];
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
w3_t[2] = pw_salt_len * 8;

4
OpenCL/m04310_a0.cl
View File

@ -152,7 +152,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04310_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;
@ -480,7 +480,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04310_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;

4
OpenCL/m04310_a1.cl
View File

@ -78,7 +78,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04310_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -446,7 +446,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04310_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m04400_a0.cl
View File

@ -135,7 +135,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1
@ -501,7 +501,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1

8
OpenCL/m04400_a1.cl
View File

@ -78,7 +78,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -158,7 +158,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -488,7 +488,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -580,7 +580,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m04500_a0.cl
View File

@ -135,7 +135,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1
@ -537,7 +537,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* sha1

8
OpenCL/m04500_a1.cl
View File

@ -78,7 +78,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -158,7 +158,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -517,7 +517,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -615,7 +615,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m04700_a0.cl
View File

@ -136,7 +136,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* md5
@ -490,7 +490,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
/**
* md5

8
OpenCL/m04700_a1.cl
View File

@ -79,7 +79,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -159,7 +159,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -471,7 +471,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -569,7 +569,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x80_2 (wordr0, wordr1, pw_r_len);
append_0x80_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m04900_a0.cl
View File

@ -200,7 +200,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_m04 (__glo
const u32 pw_salt_len = salt_len + out_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
u32 w0 = swap_workaround (w0_t[0]);
u32 w1 = swap_workaround (w0_t[1]);
@ -538,7 +538,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_s04 (__glo
const u32 pw_salt_len = salt_len + out_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
u32 w0 = swap_workaround (w0_t[0]);
u32 w1 = swap_workaround (w0_t[1]);

4
OpenCL/m04900_a1.cl
View File

@ -252,7 +252,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_m04 (__glo
const u32 pw_salt_len = salt_len + pw_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
u32 w0 = swap_workaround (w0_t[0]);
u32 w1 = swap_workaround (w0_t[1]);
@ -644,7 +644,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04900_s04 (__glo
const u32 pw_salt_len = salt_len + pw_len + salt_len;
append_0x80_4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
append_0x80_4x4 (w0_t, w1_t, w2_t, w3_t, pw_salt_len);
u32 w0 = swap_workaround (w0_t[0]);
u32 w1 = swap_workaround (w0_t[1]);

4
OpenCL/m04900_a3.cl
View File

@ -119,7 +119,7 @@ static void m04900m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
salt_buf3[2] |= salt_buf3_t[2];
salt_buf3[3] |= salt_buf3_t[3];
append_0x80_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_salt_len);
append_0x80_4x4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_salt_len);
/**
* loop
@ -432,7 +432,7 @@ static void m04900s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
salt_buf3[2] |= salt_buf3_t[2];
salt_buf3[3] |= salt_buf3_t[3];
append_0x80_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_salt_len);
append_0x80_4x4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_salt_len);
/**
* loop

4
OpenCL/m05000_a0.cl
View File

@ -162,7 +162,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x01_2 (w0, w1, out_len);
append_0x01_2x4 (w0, w1, out_len);
u64 st[25];
@ -373,7 +373,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x01_2 (w0, w1, out_len);
append_0x01_2x4 (w0, w1, out_len);
u64 st[25];

8
OpenCL/m05000_a1.cl
View File

@ -130,7 +130,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x01_2 (wordl0, wordl1, pw_l_len);
append_0x01_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -185,7 +185,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x01_2 (wordr0, wordr1, pw_r_len);
append_0x01_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}
@ -385,7 +385,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x01_2 (wordl0, wordl1, pw_l_len);
append_0x01_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -452,7 +452,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05000_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
{
append_0x01_2 (wordr0, wordr1, pw_r_len);
append_0x01_2x4 (wordr0, wordr1, pw_r_len);
switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
}

4
OpenCL/m05100_a0.cl
View File

@ -90,7 +90,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05100_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;
@ -289,7 +289,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05100_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
w3[2] = out_len * 8;

4
OpenCL/m05100_a1.cl
View File

@ -68,7 +68,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05100_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -319,7 +319,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05100_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

2
OpenCL/m05200.cl
View File

@ -233,7 +233,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05200_init (__gl
const u32 block_len = pw_len + salt_len;
append_0x80_4 (w0, w1, w2, w3, block_len);
append_0x80_4x4 (w0, w1, w2, w3, block_len);
/**
* init

4
OpenCL/m05500_a0.cl
View File

@ -645,7 +645,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];
@ -891,7 +891,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];

4
OpenCL/m05500_a1.cl
View File

@ -581,7 +581,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -869,7 +869,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m05600_a0.cl
View File

@ -381,7 +381,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05600_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];
@ -662,7 +662,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05600_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 w0_t[4];
u32 w1_t[4];

4
OpenCL/m05600_a1.cl
View File

@ -330,7 +330,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05600_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -653,7 +653,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05600_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m06000_a0.cl
View File

@ -287,7 +287,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06000_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 wl[16];
@ -415,7 +415,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06000_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 wl[16];

4
OpenCL/m06000_a1.cl
View File

@ -265,7 +265,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06000_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -435,7 +435,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06000_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

4
OpenCL/m06100_a0.cl
View File

@ -1432,7 +1432,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_m04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 wl[16];
@ -1594,7 +1594,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_s04 (__glo
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
append_0x80_2 (w0, w1, out_len);
append_0x80_2x4 (w0, w1, out_len);
u32 wl[16];

4
OpenCL/m06100_a1.cl
View File

@ -1385,7 +1385,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_m04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}
@ -1589,7 +1589,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_s04 (__glo
if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
{
append_0x80_2 (wordl0, wordl1, pw_l_len);
append_0x80_2x4 (wordl0, wordl1, pw_l_len);
switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
}

6
OpenCL/m06300.cl
View File

@ -1649,7 +1649,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_init (__gl
block_len += pw_len;
append_0x80_4 (block0, block1, block2, block3, block_len);
append_0x80_4x4 (block0, block1, block2, block3, block_len);
block3[2] = block_len * 8;
@ -1715,7 +1715,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_init (__gl
block_len++;
}
append_0x80_4 (block0, block1, block2, block3, block_len);
append_0x80_4x4 (block0, block1, block2, block3, block_len);
block3[2] = block_len * 8;
@ -1758,7 +1758,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_loop (__gl
w0_x80[2] = w0[2];
w0_x80[3] = w0[3];
append_0x80_1 (w0_x80, pw_len);
append_0x80_1x4 (w0_x80, pw_len);
/**
* salt

4
OpenCL/m06400.cl
View File

@ -308,9 +308,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06400_init (__gl
salt_buf3[2] = 0;
salt_buf3[3] = 0;
append_0x01_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 3);
append_0x01_4x4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 3);
append_0x80_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 4);
append_0x80_4x4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 4);
/**
* pads

4
OpenCL/m06500.cl
View File

@ -312,9 +312,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06500_init (__gl
salt_buf3[2] = 0;
salt_buf3[3] = 0;
append_0x01_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 3);
append_0x01_4x4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 3);
append_0x80_4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 4);
append_0x80_4x4 (salt_buf0, salt_buf1, salt_buf2, salt_buf3, salt_len + 4);
/**
* pads

4
OpenCL/m06600.cl
View File

@ -1174,8 +1174,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06600_init (__gl
w3[2] = 0;
w3[3] = 0;
append_0x01_1 (w0, salt_len + 3);
append_0x80_1 (w0, salt_len + 4);
append_0x01_1x4 (w0, salt_len + 3);
append_0x80_1x4 (w0, salt_len + 4);
w0[0] = swap_workaround (w0[0]);
w0[1] = swap_workaround (w0[1]);

Some files were not shown because too many files have changed in this diff Show More