|
|
|
@ -1027,8 +1027,6 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
|
|
|
|
|
// md5
|
|
|
|
|
|
|
|
|
|
u32 tmp2;
|
|
|
|
|
|
|
|
|
|
u32 a = MD5M_A;
|
|
|
|
|
u32 b = MD5M_B;
|
|
|
|
|
u32 c = MD5M_C;
|
|
|
|
@ -1068,22 +1066,22 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -1168,22 +1166,22 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -1282,22 +1280,22 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -1391,22 +1389,22 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -1495,22 +1493,22 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -1846,8 +1844,6 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
|
|
|
|
|
// md5
|
|
|
|
|
|
|
|
|
|
u32 tmp2;
|
|
|
|
|
|
|
|
|
|
u32 a = MD5M_A;
|
|
|
|
|
u32 b = MD5M_B;
|
|
|
|
|
u32 c = MD5M_C;
|
|
|
|
@ -1887,22 +1883,22 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -1987,22 +1983,22 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -2101,22 +2097,22 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -2210,22 +2206,22 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
@ -2314,22 +2310,22 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|
|
|
|
MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
|
|
|
|
|
MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H1, a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H2, d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H1, c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H2, b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
|
|
|
|
|
MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
|
|
|
|
|
MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
|
|
|
|
|
MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
|
|
|
|
|
MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
|
|
|
|
|
|
|
|
|
|
MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
|
|
|
|
|
MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
|
|
|
|
|