From dd14b798c3697f06ca07f90ae718ddbaabffaf33 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 8 Jul 2017 23:57:27 +0200 Subject: [PATCH] Vectorized WPA/WPA2 PMK OpenCL kernel --- OpenCL/m02501.cl | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/OpenCL/m02501.cl b/OpenCL/m02501.cl index 49503914b..275667494 100644 --- a/OpenCL/m02501.cl +++ b/OpenCL/m02501.cl @@ -68,14 +68,14 @@ __kernel void m02501_init (__global pw_t *pws, __global const kernel_rule_t *rul out_ptr[i] = hex_to_u8 (in_ptr + j); } - tmps[gid].out[0] = swap32 (out[0]); - tmps[gid].out[1] = swap32 (out[1]); - tmps[gid].out[2] = swap32 (out[2]); - tmps[gid].out[3] = swap32 (out[3]); - tmps[gid].out[4] = swap32 (out[4]); - tmps[gid].out[5] = swap32 (out[5]); - tmps[gid].out[6] = swap32 (out[6]); - tmps[gid].out[7] = swap32 (out[7]); + tmps[gid].out[0] = swap32_S (out[0]); + tmps[gid].out[1] = swap32_S (out[1]); + tmps[gid].out[2] = swap32_S (out[2]); + tmps[gid].out[3] = swap32_S (out[3]); + tmps[gid].out[4] = swap32_S (out[4]); + tmps[gid].out[5] = swap32_S (out[5]); + tmps[gid].out[6] = swap32_S (out[6]); + tmps[gid].out[7] = swap32_S (out[7]); } __kernel void m02501_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const wpa_t *wpa_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)