From 8511d9f04778d2551338368676400143c7874856 Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Sat, 6 May 2017 21:03:14 -0400 Subject: [PATCH] Completed kernel a0. --- OpenCL/m00670_a0.cl | 22 +++++++++++++++++----- 1 file changed, 17 insertions(+), 5 deletions(-) diff --git a/OpenCL/m00670_a0.cl b/OpenCL/m00670_a0.cl index 2d9c78d0a..ffb0fecb4 100644 --- a/OpenCL/m00670_a0.cl +++ b/OpenCL/m00670_a0.cl @@ -132,6 +132,17 @@ __kernel void m00670_m04 (__global pw_t *pws, __global const kernel_rule_t *rule const u32 pw_len = pws[gid].pw_len; + u32 iv[2] = { 0 }; + u32 plain[2] = { 0 }; + u32 position = esalt_bufs->position; + u32 plain_length = esalt_bufs->plain_length; + + iv[0] = esalt_bufs->iv[0]; + iv[1] = esalt_bufs->iv[1]; + + plain[0] = esalt_bufs->plain[0]; + plain[1] = esalt_bufs->plain[1]; + /** * loop */ @@ -145,13 +156,14 @@ __kernel void m00670_m04 (__global pw_t *pws, __global const kernel_rule_t *rule const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); - u64x digest[8]; + u32x digest[4] = { 0 }; + chacha20_transform (w0, w1, position, iv, plain, digest); - const u32x r0 = h32_from_64(digest[0]); - const u32x r1 = l32_from_64(digest[0]); - const u32x r2 = h32_from_64(digest[1]); - const u32x r3 = l32_from_64(digest[1]); + const u32x r0 = digest[0]; + const u32x r1 = digest[1]; + const u32x r2 = digest[2]; + const u32x r3 = digest[3]; COMPARE_M_SIMD(r0, r1, r2, r3); }