Add missing swaps to -m 6100

OpenCL/m06100_a0-pure.cl

@@ -69,7 +69,7 @@ __kernel void m06100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   for (int idx = 0; idx < pw_lenv; idx++)
   {
-    w[idx] = pws[gid].i[idx];
+    w[idx] = swap32_S (pws[gid].i[idx]);
 
     barrier (CLK_GLOBAL_MEM_FENCE);
   }
@@ -165,7 +165,7 @@ __kernel void m06100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   for (int idx = 0; idx < pw_lenv; idx++)
   {
-    w[idx] = pws[gid].i[idx];
+    w[idx] = swap32_S (pws[gid].i[idx]);
 
     barrier (CLK_GLOBAL_MEM_FENCE);
   }

OpenCL/m06100_a1-pure.cl

@@ -63,7 +63,7 @@ __kernel void m06100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   whirlpool_init (&ctx0, s_Ch, s_Cl);
 
-  whirlpool_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
+  whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
 
   /**
    * loop
@@ -73,7 +73,7 @@ __kernel void m06100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
   {
     whirlpool_ctx_t ctx = ctx0;
 
-    whirlpool_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
     whirlpool_final (&ctx);
 
@@ -148,7 +148,7 @@ __kernel void m06100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   whirlpool_init (&ctx0, s_Ch, s_Cl);
 
-  whirlpool_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
+  whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
 
   /**
    * loop
@@ -158,7 +158,7 @@ __kernel void m06100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
   {
     whirlpool_ctx_t ctx = ctx0;
 
-    whirlpool_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
     whirlpool_final (&ctx);