Converted to new SIMD: -m 6900 -a 0

diff --git a/OpenCL/m06900_a0.cl b/OpenCL/m06900_a0.cl
index 1b15632..f313b80 100644 (file)
--- a/OpenCL/m06900_a0.cl
+++ b/OpenCL/m06900_a0.cl
@@ -7,6 +7,8 @@
 
 #define _GOST_
 
+#define NEW_SIMD_CODE
+
 #include "include/constants.h"
 #include "include/kernel_vendor.h"
 
@@ -20,9 +22,7 @@
 #include "OpenCL/common.c"
 #include "include/rp_kernel.h"
 #include "OpenCL/rp.c"
-
-#define COMPARE_S "OpenCL/check_single_comp4.c"
-#define COMPARE_M "OpenCL/check_multi_comp4.c"
+#include "OpenCL/simd.c"
 
 __constant u32 c_tables[4][256] =
 {
@@ -304,7 +304,7 @@ __constant u32 c_tables[4][256] =
 
 #define _round(k1,k2,tbl)                 \
 {                                         \
-  u32 t;                                  \
+  u32x t;                                 \
   t = (k1) + r;                           \
   l ^= BOX (((t >>  0) & 0xff), 0, tbl) ^ \
        BOX (((t >>  8) & 0xff), 1, tbl) ^ \
@@ -319,8 +319,8 @@ __constant u32 c_tables[4][256] =
 
 #define R(k,h,s,i,t)      \
 {                         \
-  u32 r;                  \
-  u32 l;                  \
+  u32x r;                 \
+  u32x l;                 \
   r = h[i + 0];           \
   l = h[i + 1];           \
   _round (k[0], k[1], t);  \
@@ -389,8 +389,8 @@ __constant u32 c_tables[4][256] =
 
 #define A(x)        \
 {                   \
-  u32 l;          \
-  u32 r;          \
+  u32x l;           \
+  u32x r;           \
   l = x[0] ^ x[2];  \
   r = x[1] ^ x[3];  \
   x[0] = x[2];      \
@@ -405,8 +405,8 @@ __constant u32 c_tables[4][256] =
 
 #define AA(x)       \
 {                   \
-  u32 l;          \
-  u32 r;          \
+  u32x l;           \
+  u32x r;           \
   l    = x[0];      \
   r    = x[2];      \
   x[0] = x[4];      \
@@ -664,8 +664,8 @@ __constant u32 c_tables[4][256] =
 
 #define PASS0(h,s,u,v,t)  \
 {                         \
-  u32 k[8];             \
-  u32 w[8];             \
+  u32x k[8];              \
+  u32x w[8];              \
   X (w, u, v);            \
   P (k, w);               \
   R (k, h, s, 0, t);      \
@@ -675,8 +675,8 @@ __constant u32 c_tables[4][256] =
 
 #define PASS2(h,s,u,v,t)  \
 {                         \
-  u32 k[8];             \
-  u32 w[8];             \
+  u32x k[8];              \
+  u32x w[8];              \
   X (w, u, v);            \
   P (k, w);               \
   R (k, h, s, 2, t);      \
@@ -687,8 +687,8 @@ __constant u32 c_tables[4][256] =
 
 #define PASS4(h,s,u,v,t)  \
 {                         \
-  u32 k[8];             \
-  u32 w[8];             \
+  u32x k[8];              \
+  u32x w[8];              \
   X (w, u, v);            \
   P (k, w);               \
   R (k, h, s, 4, t);      \
@@ -698,8 +698,8 @@ __constant u32 c_tables[4][256] =
 
 #define PASS6(h,s,u,v,t)  \
 {                         \
-  u32 k[8];             \
-  u32 w[8];             \
+  u32x k[8];              \
+  u32x w[8];              \
   X (w, u, v);            \
   P (k, w);               \
   R (k, h, s, 6, t);      \
@@ -757,41 +757,18 @@ __kernel void m06900_m04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
    * loop
    */
 
-  for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
+  for (u32 il_pos = 0; il_pos < rules_cnt; il_pos += VECT_SIZE)
   {
-    u32 w0[4];
-
-    w0[0] = pw_buf0[0];
-    w0[1] = pw_buf0[1];
-    w0[2] = pw_buf0[2];
-    w0[3] = pw_buf0[3];
-
-    u32 w1[4];
-
-    w1[0] = pw_buf1[0];
-    w1[1] = pw_buf1[1];
-    w1[2] = pw_buf1[2];
-    w1[3] = pw_buf1[3];
-
-    u32 w2[4];
-
-    w2[0] = 0;
-    w2[1] = 0;
-    w2[2] = 0;
-    w2[3] = 0;
-
-    u32 w3[4];
-
-    w3[0] = 0;
-    w3[1] = 0;
-    w3[2] = 0;
-    w3[3] = 0;
+    u32x w0[4] = { 0 };
+    u32x w1[4] = { 0 };
+    u32x w2[4] = { 0 };
+    u32x w3[4] = { 0 };
 
-    const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
+    const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
     u32 w14 = out_len * 8;
 
-    u32 data[8];
+    u32x data[8];
 
     data[0] = w0[0];
     data[1] = w0[1];
@@ -802,7 +779,7 @@ __kernel void m06900_m04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     data[6] = w1[2];
     data[7] = w1[3];
 
-    u32 state[16];
+    u32x state[16];
 
     state[ 0] = 0;
     state[ 1] = 0;
@@ -821,8 +798,8 @@ __kernel void m06900_m04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     state[14] = data[6];
     state[15] = data[7];
 
-    u32 state_m[8];
-    u32 data_m[8];
+    u32x state_m[8];
+    u32x data_m[8];
 
     /* gost1 */
 
@@ -844,9 +821,9 @@ __kernel void m06900_m04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     data_m[6] = data[6];
     data_m[7] = data[7];
 
-    u32 tmp[8];
+    u32x tmp[8];
 
-    if (out_len > 0)
+    if (pw_len > 0)
     {
       PASS0 (state, tmp, state_m, data_m, s_tables);
       PASS2 (state, tmp, state_m, data_m, s_tables);
@@ -936,12 +913,7 @@ __kernel void m06900_m04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
 
     /* store */
 
-    const u32 r0 = state[0];
-    const u32 r1 = state[1];
-    const u32 r2 = state[2];
-    const u32 r3 = state[3];
-
-    #include COMPARE_M
+    COMPARE_M_SIMD (state[0], state[1], state[2], state[3]);
   }
 }
 
@@ -1017,41 +989,18 @@ __kernel void m06900_s04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
    * loop
    */
 
-  for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
+  for (u32 il_pos = 0; il_pos < rules_cnt; il_pos += VECT_SIZE)
   {
-    u32 w0[4];
-
-    w0[0] = pw_buf0[0];
-    w0[1] = pw_buf0[1];
-    w0[2] = pw_buf0[2];
-    w0[3] = pw_buf0[3];
-
-    u32 w1[4];
-
-    w1[0] = pw_buf1[0];
-    w1[1] = pw_buf1[1];
-    w1[2] = pw_buf1[2];
-    w1[3] = pw_buf1[3];
-
-    u32 w2[4];
+    u32x w0[4] = { 0 };
+    u32x w1[4] = { 0 };
+    u32x w2[4] = { 0 };
+    u32x w3[4] = { 0 };
 
-    w2[0] = 0;
-    w2[1] = 0;
-    w2[2] = 0;
-    w2[3] = 0;
-
-    u32 w3[4];
-
-    w3[0] = 0;
-    w3[1] = 0;
-    w3[2] = 0;
-    w3[3] = 0;
-
-    const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
+    const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
     u32 w14 = out_len * 8;
 
-    u32 data[8];
+    u32x data[8];
 
     data[0] = w0[0];
     data[1] = w0[1];
@@ -1062,7 +1011,7 @@ __kernel void m06900_s04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     data[6] = w1[2];
     data[7] = w1[3];
 
-    u32 state[16];
+    u32x state[16];
 
     state[ 0] = 0;
     state[ 1] = 0;
@@ -1081,8 +1030,8 @@ __kernel void m06900_s04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     state[14] = data[6];
     state[15] = data[7];
 
-    u32 state_m[8];
-    u32 data_m[8];
+    u32x state_m[8];
+    u32x data_m[8];
 
     /* gost1 */
 
@@ -1104,9 +1053,9 @@ __kernel void m06900_s04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     data_m[6] = data[6];
     data_m[7] = data[7];
 
-    u32 tmp[8];
+    u32x tmp[8];
 
-    if (out_len > 0)
+    if (pw_len > 0)
     {
       PASS0 (state, tmp, state_m, data_m, s_tables);
       PASS2 (state, tmp, state_m, data_m, s_tables);
@@ -1196,12 +1145,7 @@ __kernel void m06900_s04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
 
     /* store */
 
-    const u32 r0 = state[0];
-    const u32 r1 = state[1];
-    const u32 r2 = state[2];
-    const u32 r3 = state[3];
-
-    #include COMPARE_S
+    COMPARE_S_SIMD (state[0], state[1], state[2], state[3]);
   }
 }