diff --git a/OpenCL/m08100_a0-pure.cl b/OpenCL/m08100_a0-pure.cl index 602e73dd7..bed40e229 100644 --- a/OpenCL/m08100_a0-pure.cl +++ b/OpenCL/m08100_a0-pure.cl @@ -30,6 +30,8 @@ __kernel void m08100_mxx (KERN_ATTR_RULES ()) * base */ + u32 z[16] = { 0 }; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -50,7 +52,9 @@ __kernel void m08100_mxx (KERN_ATTR_RULES ()) sha1_ctx_t ctx = ctx0; - sha1_update_swap (&ctx, tmp.i, tmp.pw_len + 1); + sha1_update_swap (&ctx, tmp.i, tmp.pw_len); + + sha1_update (&ctx, z, 1); sha1_final (&ctx); @@ -90,6 +94,8 @@ __kernel void m08100_sxx (KERN_ATTR_RULES ()) * base */ + u32 z[16] = { 0 }; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -110,7 +116,9 @@ __kernel void m08100_sxx (KERN_ATTR_RULES ()) sha1_ctx_t ctx = ctx0; - sha1_update_swap (&ctx, tmp.i, tmp.pw_len + 1); + sha1_update_swap (&ctx, tmp.i, tmp.pw_len); + + sha1_update (&ctx, z, 1); sha1_final (&ctx); diff --git a/OpenCL/m08100_a1-pure.cl b/OpenCL/m08100_a1-pure.cl index 395b55c8f..6214ce7ff 100644 --- a/OpenCL/m08100_a1-pure.cl +++ b/OpenCL/m08100_a1-pure.cl @@ -28,6 +28,8 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ()) * base */ + u32 z[16] = { 0 }; + sha1_ctx_t ctx0; sha1_init (&ctx0); @@ -44,7 +46,9 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ()) { sha1_ctx_t ctx = ctx0; - sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len + 1); + sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + + sha1_update (&ctx, z, 1); sha1_final (&ctx); @@ -57,6 +61,7 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ()) } } + __kernel void m08100_sxx (KERN_ATTR_BASIC ()) { /** @@ -84,6 +89,8 @@ __kernel void m08100_sxx (KERN_ATTR_BASIC ()) * base */ + u32 z[16] = { 0 }; + sha1_ctx_t ctx0; sha1_init (&ctx0); @@ -100,7 +107,9 @@ __kernel void m08100_sxx (KERN_ATTR_BASIC ()) { sha1_ctx_t ctx = ctx0; - sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len + 1); + sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + + sha1_update (&ctx, z, 1); sha1_final (&ctx); diff --git a/OpenCL/m08100_a3-pure.cl b/OpenCL/m08100_a3-pure.cl index 797646eb3..eefe97f74 100644 --- a/OpenCL/m08100_a3-pure.cl +++ b/OpenCL/m08100_a3-pure.cl @@ -28,6 +28,8 @@ __kernel void m08100_mxx (KERN_ATTR_VECTOR ()) * base */ + u32x z[16] = { 0 }; + const u32 pw_len = pws[gid].pw_len; u32x w[64] = { 0 }; @@ -61,7 +63,9 @@ __kernel void m08100_mxx (KERN_ATTR_VECTOR ()) sha1_init_vector_from_scalar (&ctx, &ctx0); - sha1_update_vector (&ctx, w, pw_len + 1); + sha1_update_vector (&ctx, w, pw_len); + + sha1_update_vector (&ctx, z, 1); sha1_final_vector (&ctx); @@ -101,6 +105,8 @@ __kernel void m08100_sxx (KERN_ATTR_VECTOR ()) * base */ + u32x z[16] = { 0 }; + const u32 pw_len = pws[gid].pw_len; u32x w[64] = { 0 }; @@ -134,7 +140,9 @@ __kernel void m08100_sxx (KERN_ATTR_VECTOR ()) sha1_init_vector_from_scalar (&ctx, &ctx0); - sha1_update_vector (&ctx, w, pw_len + 1); + sha1_update_vector (&ctx, w, pw_len); + + sha1_update_vector (&ctx, z, 1); sha1_final_vector (&ctx);