From bed7e8f466552bd55bc7b3e74d238a9f0f82f225 Mon Sep 17 00:00:00 2001
From: jsteube <jens.steube@gmail.com>
Date: Thu, 24 Aug 2017 20:07:43 +0200
Subject: [PATCH] Remove unused truncate_block_xxx_xx() functions and update
 kernels to use the _S function

---
 OpenCL/inc_common.cl          | 3554 ++++++++++++++++-----------------
 OpenCL/m00500-optimized.cl    |    2 +-
 OpenCL/m00500.cl              |    2 +-
 OpenCL/m01600-optimized.cl    |    2 +-
 OpenCL/m01600.cl              |    2 +-
 OpenCL/m01800.cl              |    6 +-
 OpenCL/m06300-optimized.cl    |    2 +-
 OpenCL/m06300.cl              |    2 +-
 OpenCL/m07400-optimized.cl    |    6 +-
 OpenCL/m07400.cl              |    6 +-
 OpenCL/m13100_a0-optimized.cl |    8 +-
 OpenCL/m13100_a0.cl           |    8 +-
 OpenCL/m13100_a1-optimized.cl |    8 +-
 OpenCL/m13100_a1.cl           |    8 +-
 OpenCL/m13100_a3-optimized.cl |    8 +-
 OpenCL/m13100_a3.cl           |    8 +-
 OpenCL/m13400.cl              |    4 +-
 17 files changed, 1763 insertions(+), 1873 deletions(-)

diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl
index 20c69899d..577445ccd 100644
--- a/OpenCL/inc_common.cl
+++ b/OpenCL/inc_common.cl
@@ -144,1838 +144,6 @@ static float get_entropy (const u32 *buf, const int elems)
  * vector functions
  */
 
-static void truncate_block_4x4_le (u32x w0[4], const u32 len)
-{
-  switch (len)
-  {
-    case  0:
-      w0[0]  = 0;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  1:
-      w0[0] &= 0x000000ff;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  2:
-      w0[0] &= 0x0000ffff;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  3:
-      w0[0] &= 0x00ffffff;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  4:
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  5:
-      w0[1] &= 0x000000ff;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  6:
-      w0[1] &= 0x0000ffff;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  7:
-      w0[1] &= 0x00ffffff;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  8:
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  9:
-      w0[2] &= 0x000000ff;
-      w0[3]  = 0;
-
-      break;
-
-    case 10:
-      w0[2] &= 0x0000ffff;
-      w0[3]  = 0;
-
-      break;
-
-    case 11:
-      w0[2] &= 0x00ffffff;
-      w0[3]  = 0;
-
-      break;
-
-    case 12:
-      w0[3]  = 0;
-
-      break;
-
-    case 13:
-      w0[3] &= 0x000000ff;
-
-      break;
-
-    case 14:
-      w0[3] &= 0x0000ffff;
-
-      break;
-
-    case 15:
-      w0[3] &= 0x00ffffff;
-
-      break;
-  }
-}
-
-static void truncate_block_16x4_le (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 len)
-{
-  switch (len)
-  {
-    case  0:
-      w0[0]  = 0;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  1:
-      w0[0] &= 0x000000ff;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  2:
-      w0[0] &= 0x0000ffff;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  3:
-      w0[0] &= 0x00ffffff;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  4:
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  5:
-      w0[1] &= 0x000000ff;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  6:
-      w0[1] &= 0x0000ffff;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  7:
-      w0[1] &= 0x00ffffff;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  8:
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  9:
-      w0[2] &= 0x000000ff;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 10:
-      w0[2] &= 0x0000ffff;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 11:
-      w0[2] &= 0x00ffffff;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 12:
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 13:
-      w0[3] &= 0x000000ff;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 14:
-      w0[3] &= 0x0000ffff;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 15:
-      w0[3] &= 0x00ffffff;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 16:
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 17:
-      w1[0] &= 0x000000ff;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 18:
-      w1[0] &= 0x0000ffff;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 19:
-      w1[0] &= 0x00ffffff;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 20:
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 21:
-      w1[1] &= 0x000000ff;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 22:
-      w1[1] &= 0x0000ffff;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 23:
-      w1[1] &= 0x00ffffff;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 24:
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 25:
-      w1[2] &= 0x000000ff;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 26:
-      w1[2] &= 0x0000ffff;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 27:
-      w1[2] &= 0x00ffffff;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 28:
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 29:
-      w1[3] &= 0x000000ff;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 30:
-      w1[3] &= 0x0000ffff;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 31:
-      w1[3] &= 0x00ffffff;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 32:
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 33:
-      w2[0] &= 0x000000ff;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 34:
-      w2[0] &= 0x0000ffff;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 35:
-      w2[0] &= 0x00ffffff;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 36:
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 37:
-      w2[1] &= 0x000000ff;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 38:
-      w2[1] &= 0x0000ffff;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 39:
-      w2[1] &= 0x00ffffff;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 40:
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 41:
-      w2[2] &= 0x000000ff;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 42:
-      w2[2] &= 0x0000ffff;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 43:
-      w2[2] &= 0x00ffffff;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 44:
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 45:
-      w2[3] &= 0x000000ff;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 46:
-      w2[3] &= 0x0000ffff;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 47:
-      w2[3] &= 0x00ffffff;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 48:
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 49:
-      w3[0] &= 0x000000ff;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 50:
-      w3[0] &= 0x0000ffff;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 51:
-      w3[0] &= 0x00ffffff;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 52:
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 53:
-      w3[1] &= 0x000000ff;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 54:
-      w3[1] &= 0x0000ffff;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 55:
-      w3[1] &= 0x00ffffff;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 56:
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 57:
-      w3[2] &= 0x000000ff;
-      w3[3]  = 0;
-
-      break;
-
-    case 58:
-      w3[2] &= 0x0000ffff;
-      w3[3]  = 0;
-
-      break;
-
-    case 59:
-      w3[2] &= 0x00ffffff;
-      w3[3]  = 0;
-
-      break;
-
-    case 60:
-      w3[3]  = 0;
-
-      break;
-
-    case 61:
-      w3[3] &= 0x000000ff;
-
-      break;
-
-    case 62:
-      w3[3] &= 0x0000ffff;
-
-      break;
-
-    case 63:
-      w3[3] &= 0x00ffffff;
-
-      break;
-  }
-}
-
-static void truncate_block_4x4_be (u32x w0[4], const u32 len)
-{
-  switch (len)
-  {
-    case  0:
-      w0[0]  = 0;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  1:
-      w0[0] &= 0xff000000;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  2:
-      w0[0] &= 0xffff0000;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  3:
-      w0[0] &= 0xffffff00;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  4:
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  5:
-      w0[1] &= 0xff000000;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  6:
-      w0[1] &= 0xffff0000;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  7:
-      w0[1] &= 0xffffff00;
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  8:
-      w0[2]  = 0;
-      w0[3]  = 0;
-
-      break;
-
-    case  9:
-      w0[2] &= 0xff000000;
-      w0[3]  = 0;
-
-      break;
-
-    case 10:
-      w0[2] &= 0xffff0000;
-      w0[3]  = 0;
-
-      break;
-
-    case 11:
-      w0[2] &= 0xffffff00;
-      w0[3]  = 0;
-
-      break;
-
-    case 12:
-      w0[3]  = 0;
-
-      break;
-
-    case 13:
-      w0[3] &= 0xff000000;
-
-      break;
-
-    case 14:
-      w0[3] &= 0xffff0000;
-
-      break;
-
-    case 15:
-      w0[3] &= 0xffffff00;
-
-      break;
-  }
-}
-
-static void truncate_block_16x4_be (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 len)
-{
-  switch (len)
-  {
-    case  0:
-      w0[0]  = 0;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  1:
-      w0[0] &= 0xff000000;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  2:
-      w0[0] &= 0xffff0000;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  3:
-      w0[0] &= 0xffffff00;
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  4:
-      w0[1]  = 0;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  5:
-      w0[1] &= 0xff000000;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  6:
-      w0[1] &= 0xffff0000;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  7:
-      w0[1] &= 0xffffff00;
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  8:
-      w0[2]  = 0;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case  9:
-      w0[2] &= 0xff000000;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 10:
-      w0[2] &= 0xffff0000;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 11:
-      w0[2] &= 0xffffff00;
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 12:
-      w0[3]  = 0;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 13:
-      w0[3] &= 0xff000000;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 14:
-      w0[3] &= 0xffff0000;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 15:
-      w0[3] &= 0xffffff00;
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 16:
-      w1[0]  = 0;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 17:
-      w1[0] &= 0xff000000;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 18:
-      w1[0] &= 0xffff0000;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 19:
-      w1[0] &= 0xffffff00;
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 20:
-      w1[1]  = 0;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 21:
-      w1[1] &= 0xff000000;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 22:
-      w1[1] &= 0xffff0000;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 23:
-      w1[1] &= 0xffffff00;
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 24:
-      w1[2]  = 0;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 25:
-      w1[2] &= 0xff000000;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 26:
-      w1[2] &= 0xffff0000;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 27:
-      w1[2] &= 0xffffff00;
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 28:
-      w1[3]  = 0;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 29:
-      w1[3] &= 0xff000000;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 30:
-      w1[3] &= 0xffff0000;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 31:
-      w1[3] &= 0xffffff00;
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 32:
-      w2[0]  = 0;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 33:
-      w2[0] &= 0xff000000;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 34:
-      w2[0] &= 0xffff0000;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 35:
-      w2[0] &= 0xffffff00;
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 36:
-      w2[1]  = 0;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 37:
-      w2[1] &= 0xff000000;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 38:
-      w2[1] &= 0xffff0000;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 39:
-      w2[1] &= 0xffffff00;
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 40:
-      w2[2]  = 0;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 41:
-      w2[2] &= 0xff000000;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 42:
-      w2[2] &= 0xffff0000;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 43:
-      w2[2] &= 0xffffff00;
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 44:
-      w2[3]  = 0;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 45:
-      w2[3] &= 0xff000000;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 46:
-      w2[3] &= 0xffff0000;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 47:
-      w2[3] &= 0xffffff00;
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 48:
-      w3[0]  = 0;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 49:
-      w3[0] &= 0xff000000;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 50:
-      w3[0] &= 0xffff0000;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 51:
-      w3[0] &= 0xffffff00;
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 52:
-      w3[1]  = 0;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 53:
-      w3[1] &= 0xff000000;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 54:
-      w3[1] &= 0xffff0000;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 55:
-      w3[1] &= 0xffffff00;
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 56:
-      w3[2]  = 0;
-      w3[3]  = 0;
-
-      break;
-
-    case 57:
-      w3[2] &= 0xff000000;
-      w3[3]  = 0;
-
-      break;
-
-    case 58:
-      w3[2] &= 0xffff0000;
-      w3[3]  = 0;
-
-      break;
-
-    case 59:
-      w3[2] &= 0xffffff00;
-      w3[3]  = 0;
-
-      break;
-
-    case 60:
-      w3[3]  = 0;
-
-      break;
-
-    case 61:
-      w3[3] &= 0xff000000;
-
-      break;
-
-    case 62:
-      w3[3] &= 0xffff0000;
-
-      break;
-
-    case 63:
-      w3[3] &= 0xffffff00;
-
-      break;
-  }
-}
-
 static void make_utf16be (const u32x in[4], u32x out1[4], u32x out2[4])
 {
   #ifdef IS_NV
@@ -31616,6 +29784,1728 @@ static void truncate_block_4x4_le_S (u32 w0[4], const u32 len)
   }
 }
 
+static void truncate_block_4x4_be_S (u32 w0[4], const u32 len)
+{
+  switch (len)
+  {
+    case  0:
+      w0[0]  = 0;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  1:
+      w0[0] &= 0xff000000;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  2:
+      w0[0] &= 0xffff0000;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  3:
+      w0[0] &= 0xffffff00;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  4:
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  5:
+      w0[1] &= 0xff000000;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  6:
+      w0[1] &= 0xffff0000;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  7:
+      w0[1] &= 0xffffff00;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  8:
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  9:
+      w0[2] &= 0xff000000;
+      w0[3]  = 0;
+
+      break;
+
+    case 10:
+      w0[2] &= 0xffff0000;
+      w0[3]  = 0;
+
+      break;
+
+    case 11:
+      w0[2] &= 0xffffff00;
+      w0[3]  = 0;
+
+      break;
+
+    case 12:
+      w0[3]  = 0;
+
+      break;
+
+    case 13:
+      w0[3] &= 0xff000000;
+
+      break;
+
+    case 14:
+      w0[3] &= 0xffff0000;
+
+      break;
+
+    case 15:
+      w0[3] &= 0xffffff00;
+
+      break;
+  }
+}
+
+static void truncate_block_16x4_le_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 len)
+{
+  switch (len)
+  {
+    case  0:
+      w0[0]  = 0;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  1:
+      w0[0] &= 0x000000ff;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  2:
+      w0[0] &= 0x0000ffff;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  3:
+      w0[0] &= 0x00ffffff;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  4:
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  5:
+      w0[1] &= 0x000000ff;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  6:
+      w0[1] &= 0x0000ffff;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  7:
+      w0[1] &= 0x00ffffff;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  8:
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  9:
+      w0[2] &= 0x000000ff;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 10:
+      w0[2] &= 0x0000ffff;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 11:
+      w0[2] &= 0x00ffffff;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 12:
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 13:
+      w0[3] &= 0x000000ff;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 14:
+      w0[3] &= 0x0000ffff;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 15:
+      w0[3] &= 0x00ffffff;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 16:
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 17:
+      w1[0] &= 0x000000ff;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 18:
+      w1[0] &= 0x0000ffff;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 19:
+      w1[0] &= 0x00ffffff;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 20:
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 21:
+      w1[1] &= 0x000000ff;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 22:
+      w1[1] &= 0x0000ffff;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 23:
+      w1[1] &= 0x00ffffff;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 24:
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 25:
+      w1[2] &= 0x000000ff;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 26:
+      w1[2] &= 0x0000ffff;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 27:
+      w1[2] &= 0x00ffffff;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 28:
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 29:
+      w1[3] &= 0x000000ff;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 30:
+      w1[3] &= 0x0000ffff;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 31:
+      w1[3] &= 0x00ffffff;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 32:
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 33:
+      w2[0] &= 0x000000ff;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 34:
+      w2[0] &= 0x0000ffff;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 35:
+      w2[0] &= 0x00ffffff;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 36:
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 37:
+      w2[1] &= 0x000000ff;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 38:
+      w2[1] &= 0x0000ffff;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 39:
+      w2[1] &= 0x00ffffff;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 40:
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 41:
+      w2[2] &= 0x000000ff;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 42:
+      w2[2] &= 0x0000ffff;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 43:
+      w2[2] &= 0x00ffffff;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 44:
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 45:
+      w2[3] &= 0x000000ff;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 46:
+      w2[3] &= 0x0000ffff;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 47:
+      w2[3] &= 0x00ffffff;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 48:
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 49:
+      w3[0] &= 0x000000ff;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 50:
+      w3[0] &= 0x0000ffff;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 51:
+      w3[0] &= 0x00ffffff;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 52:
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 53:
+      w3[1] &= 0x000000ff;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 54:
+      w3[1] &= 0x0000ffff;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 55:
+      w3[1] &= 0x00ffffff;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 56:
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 57:
+      w3[2] &= 0x000000ff;
+      w3[3]  = 0;
+
+      break;
+
+    case 58:
+      w3[2] &= 0x0000ffff;
+      w3[3]  = 0;
+
+      break;
+
+    case 59:
+      w3[2] &= 0x00ffffff;
+      w3[3]  = 0;
+
+      break;
+
+    case 60:
+      w3[3]  = 0;
+
+      break;
+
+    case 61:
+      w3[3] &= 0x000000ff;
+
+      break;
+
+    case 62:
+      w3[3] &= 0x0000ffff;
+
+      break;
+
+    case 63:
+      w3[3] &= 0x00ffffff;
+
+      break;
+  }
+}
+
+static void truncate_block_16x4_be_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 len)
+{
+  switch (len)
+  {
+    case  0:
+      w0[0]  = 0;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  1:
+      w0[0] &= 0xff000000;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  2:
+      w0[0] &= 0xffff0000;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  3:
+      w0[0] &= 0xffffff00;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  4:
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  5:
+      w0[1] &= 0xff000000;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  6:
+      w0[1] &= 0xffff0000;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  7:
+      w0[1] &= 0xffffff00;
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  8:
+      w0[2]  = 0;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case  9:
+      w0[2] &= 0xff000000;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 10:
+      w0[2] &= 0xffff0000;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 11:
+      w0[2] &= 0xffffff00;
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 12:
+      w0[3]  = 0;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 13:
+      w0[3] &= 0xff000000;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 14:
+      w0[3] &= 0xffff0000;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 15:
+      w0[3] &= 0xffffff00;
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 16:
+      w1[0]  = 0;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 17:
+      w1[0] &= 0xff000000;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 18:
+      w1[0] &= 0xffff0000;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 19:
+      w1[0] &= 0xffffff00;
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 20:
+      w1[1]  = 0;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 21:
+      w1[1] &= 0xff000000;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 22:
+      w1[1] &= 0xffff0000;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 23:
+      w1[1] &= 0xffffff00;
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 24:
+      w1[2]  = 0;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 25:
+      w1[2] &= 0xff000000;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 26:
+      w1[2] &= 0xffff0000;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 27:
+      w1[2] &= 0xffffff00;
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 28:
+      w1[3]  = 0;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 29:
+      w1[3] &= 0xff000000;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 30:
+      w1[3] &= 0xffff0000;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 31:
+      w1[3] &= 0xffffff00;
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 32:
+      w2[0]  = 0;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 33:
+      w2[0] &= 0xff000000;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 34:
+      w2[0] &= 0xffff0000;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 35:
+      w2[0] &= 0xffffff00;
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 36:
+      w2[1]  = 0;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 37:
+      w2[1] &= 0xff000000;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 38:
+      w2[1] &= 0xffff0000;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 39:
+      w2[1] &= 0xffffff00;
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 40:
+      w2[2]  = 0;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 41:
+      w2[2] &= 0xff000000;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 42:
+      w2[2] &= 0xffff0000;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 43:
+      w2[2] &= 0xffffff00;
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 44:
+      w2[3]  = 0;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 45:
+      w2[3] &= 0xff000000;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 46:
+      w2[3] &= 0xffff0000;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 47:
+      w2[3] &= 0xffffff00;
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 48:
+      w3[0]  = 0;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 49:
+      w3[0] &= 0xff000000;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 50:
+      w3[0] &= 0xffff0000;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 51:
+      w3[0] &= 0xffffff00;
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 52:
+      w3[1]  = 0;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 53:
+      w3[1] &= 0xff000000;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 54:
+      w3[1] &= 0xffff0000;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 55:
+      w3[1] &= 0xffffff00;
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 56:
+      w3[2]  = 0;
+      w3[3]  = 0;
+
+      break;
+
+    case 57:
+      w3[2] &= 0xff000000;
+      w3[3]  = 0;
+
+      break;
+
+    case 58:
+      w3[2] &= 0xffff0000;
+      w3[3]  = 0;
+
+      break;
+
+    case 59:
+      w3[2] &= 0xffffff00;
+      w3[3]  = 0;
+
+      break;
+
+    case 60:
+      w3[3]  = 0;
+
+      break;
+
+    case 61:
+      w3[3] &= 0xff000000;
+
+      break;
+
+    case 62:
+      w3[3] &= 0xffff0000;
+
+      break;
+
+    case 63:
+      w3[3] &= 0xffffff00;
+
+      break;
+  }
+}
+
 static void append_0x01_2x4_S (u32 w0[4], u32 w1[4], const u32 offset)
 {
   const u32 tmp = 0x01 << ((offset & 3) * 8);
diff --git a/OpenCL/m00500-optimized.cl b/OpenCL/m00500-optimized.cl
index d2996f814..dda6c3923 100644
--- a/OpenCL/m00500-optimized.cl
+++ b/OpenCL/m00500-optimized.cl
@@ -776,7 +776,7 @@ __kernel void m00500_init (__global pw_t *pws, __constant const kernel_rule_t *r
 
   block_len += salt_len;
 
-  truncate_block_4x4_le (digest, pw_len);
+  truncate_block_4x4_le_S (digest, pw_len);
 
   memcat16 (block0, block1, block2, block3, block_len, digest);
 
diff --git a/OpenCL/m00500.cl b/OpenCL/m00500.cl
index 17972c21b..b071ca6ca 100644
--- a/OpenCL/m00500.cl
+++ b/OpenCL/m00500.cl
@@ -93,7 +93,7 @@ __kernel void m00500_init (__global pw_t *pws, __constant const kernel_rule_t *r
     md5_update (&md5_ctx, final, 16);
   }
 
-  truncate_block_4x4_le (final, pl);
+  truncate_block_4x4_le_S (final, pl);
 
   md5_update (&md5_ctx, final, pl);
 
diff --git a/OpenCL/m01600-optimized.cl b/OpenCL/m01600-optimized.cl
index 5501af7ba..ab2d438ab 100644
--- a/OpenCL/m01600-optimized.cl
+++ b/OpenCL/m01600-optimized.cl
@@ -791,7 +791,7 @@ __kernel void m01600_init (__global pw_t *pws, __constant const kernel_rule_t *r
 
   block_len += salt_len;
 
-  truncate_block_4x4_le (digest, pw_len);
+  truncate_block_4x4_le_S (digest, pw_len);
 
   memcat16 (block0, block1, block2, block3, block_len, digest);
 
diff --git a/OpenCL/m01600.cl b/OpenCL/m01600.cl
index 1ebedc79d..2c69ce3e0 100644
--- a/OpenCL/m01600.cl
+++ b/OpenCL/m01600.cl
@@ -95,7 +95,7 @@ __kernel void m01600_init (__global pw_t *pws, __constant const kernel_rule_t *r
     md5_update (&md5_ctx, final, 16);
   }
 
-  truncate_block_4x4_le (final, pl);
+  truncate_block_4x4_le_S (final, pl);
 
   md5_update (&md5_ctx, final, pl);
 
diff --git a/OpenCL/m01800.cl b/OpenCL/m01800.cl
index 22539bed0..d70dbd419 100644
--- a/OpenCL/m01800.cl
+++ b/OpenCL/m01800.cl
@@ -112,7 +112,7 @@ __kernel void m01800_init (__global pw_t *pws, __constant const kernel_rule_t *r
   #endif
   for (int i = 0; i < 16; i++) t_final[i] = final[i];
 
-  truncate_block_16x4_be (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
+  truncate_block_16x4_be_S (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
 
   sha512_update (&ctx, t_final, pl);
 
@@ -199,7 +199,7 @@ __kernel void m01800_init (__global pw_t *pws, __constant const kernel_rule_t *r
     p_final[idx + 15] = final[15];
   }
 
-  truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
+  truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
 
   p_final[idx +  0] = final[ 0];
   p_final[idx +  1] = final[ 1];
@@ -273,7 +273,7 @@ __kernel void m01800_init (__global pw_t *pws, __constant const kernel_rule_t *r
     s_final[idx + 15] = final[15];
   }
 
-  truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
+  truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
 
   s_final[idx +  0] = final[ 0];
   s_final[idx +  1] = final[ 1];
diff --git a/OpenCL/m06300-optimized.cl b/OpenCL/m06300-optimized.cl
index 5feeedd15..8d5a6bafc 100644
--- a/OpenCL/m06300-optimized.cl
+++ b/OpenCL/m06300-optimized.cl
@@ -685,7 +685,7 @@ __kernel void m06300_init (__global pw_t *pws, __constant const kernel_rule_t *r
 
   block_len += salt_len;
 
-  truncate_block_4x4_le (digest, pw_len);
+  truncate_block_4x4_le_S (digest, pw_len);
 
   memcat16 (block0, block1, block2, block3, block_len, digest);
 
diff --git a/OpenCL/m06300.cl b/OpenCL/m06300.cl
index 79614fd55..7f228d073 100644
--- a/OpenCL/m06300.cl
+++ b/OpenCL/m06300.cl
@@ -85,7 +85,7 @@ __kernel void m06300_init (__global pw_t *pws, __constant const kernel_rule_t *r
     md5_update (&md5_ctx, final, 16);
   }
 
-  truncate_block_4x4_le (final, pl);
+  truncate_block_4x4_le_S (final, pl);
 
   md5_update (&md5_ctx, final, pl);
 
diff --git a/OpenCL/m07400-optimized.cl b/OpenCL/m07400-optimized.cl
index 97d831798..a6173ece3 100644
--- a/OpenCL/m07400-optimized.cl
+++ b/OpenCL/m07400-optimized.cl
@@ -780,7 +780,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
   alt_result_tmp[6] = 0;
   alt_result_tmp[7] = 0;
 
-  truncate_block_4x4_le (alt_result_tmp, pw_len);
+  truncate_block_4x4_le_S (alt_result_tmp, pw_len);
 
   /* Add the key string.  */
 
@@ -890,7 +890,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
 
   bswap8 (p_bytes);
 
-  truncate_block_4x4_le (p_bytes, pw_len);
+  truncate_block_4x4_le_S (p_bytes, pw_len);
 
   tmps[gid].p_bytes[0] = p_bytes[0];
   tmps[gid].p_bytes[1] = p_bytes[1];
@@ -933,7 +933,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
 
   bswap8 (s_bytes);
 
-  truncate_block_4x4_le (s_bytes, salt_len);
+  truncate_block_4x4_le_S (s_bytes, salt_len);
 
   tmps[gid].s_bytes[0] = s_bytes[0];
   tmps[gid].s_bytes[1] = s_bytes[1];
diff --git a/OpenCL/m07400.cl b/OpenCL/m07400.cl
index 50def1f3d..0eac4292c 100644
--- a/OpenCL/m07400.cl
+++ b/OpenCL/m07400.cl
@@ -104,7 +104,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
   #endif
   for (int i = 0; i < 8; i++) t_final[i] = final[i];
 
-  truncate_block_16x4_be (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
+  truncate_block_16x4_be_S (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
 
   sha256_update (&ctx, t_final, pl);
 
@@ -175,7 +175,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
     p_final[idx + 7] = final[7];
   }
 
-  truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
+  truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
 
   p_final[idx + 0] = final[0];
   p_final[idx + 1] = final[1];
@@ -233,7 +233,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
     s_final[idx + 7] = final[7];
   }
 
-  truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
+  truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
 
   s_final[idx + 0] = final[0];
   s_final[idx + 1] = final[1];
diff --git a/OpenCL/m13100_a0-optimized.cl b/OpenCL/m13100_a0-optimized.cl
index 8f1149623..f1a566a8e 100644
--- a/OpenCL/m13100_a0-optimized.cl
+++ b/OpenCL/m13100_a0-optimized.cl
@@ -321,7 +321,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w0, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w0, edata2_left & 0xf);
 
     append_0x80_1x4 (w0, edata2_left & 0xf);
 
@@ -335,7 +335,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w1, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w1, edata2_left & 0xf);
 
     append_0x80_1x4 (w1, edata2_left & 0xf);
 
@@ -350,7 +350,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w2, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w2, edata2_left & 0xf);
 
     append_0x80_1x4 (w2, edata2_left & 0xf);
 
@@ -366,7 +366,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w3, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w3, edata2_left & 0xf);
 
     append_0x80_1x4 (w3, edata2_left & 0xf);
 
diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl
index 2fe1b8399..270cef3b5 100644
--- a/OpenCL/m13100_a0.cl
+++ b/OpenCL/m13100_a0.cl
@@ -234,14 +234,14 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w0, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w0, edata2_left & 0xf);
   }
   else if (edata2_left < 32)
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w1, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w1, edata2_left & 0xf);
   }
   else if (edata2_left < 48)
   {
@@ -249,7 +249,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w2, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w2, edata2_left & 0xf);
   }
   else
   {
@@ -258,7 +258,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w3, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w3, edata2_left & 0xf);
   }
 
   md5_hmac_update_64 (&ctx, w0, w1, w2, w3, edata2_left);
diff --git a/OpenCL/m13100_a1-optimized.cl b/OpenCL/m13100_a1-optimized.cl
index f697c2f5f..5ed9ae989 100644
--- a/OpenCL/m13100_a1-optimized.cl
+++ b/OpenCL/m13100_a1-optimized.cl
@@ -319,7 +319,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w0, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w0, edata2_left & 0xf);
 
     append_0x80_1x4 (w0, edata2_left & 0xf);
 
@@ -333,7 +333,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w1, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w1, edata2_left & 0xf);
 
     append_0x80_1x4 (w1, edata2_left & 0xf);
 
@@ -348,7 +348,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w2, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w2, edata2_left & 0xf);
 
     append_0x80_1x4 (w2, edata2_left & 0xf);
 
@@ -364,7 +364,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w3, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w3, edata2_left & 0xf);
 
     append_0x80_1x4 (w3, edata2_left & 0xf);
 
diff --git a/OpenCL/m13100_a1.cl b/OpenCL/m13100_a1.cl
index 5d985491a..aadc57e5c 100644
--- a/OpenCL/m13100_a1.cl
+++ b/OpenCL/m13100_a1.cl
@@ -232,14 +232,14 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w0, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w0, edata2_left & 0xf);
   }
   else if (edata2_left < 32)
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w1, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w1, edata2_left & 0xf);
   }
   else if (edata2_left < 48)
   {
@@ -247,7 +247,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w2, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w2, edata2_left & 0xf);
   }
   else
   {
@@ -256,7 +256,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w3, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w3, edata2_left & 0xf);
   }
 
   md5_hmac_update_64 (&ctx, w0, w1, w2, w3, edata2_left);
diff --git a/OpenCL/m13100_a3-optimized.cl b/OpenCL/m13100_a3-optimized.cl
index ce6e3863b..11ca38a79 100644
--- a/OpenCL/m13100_a3-optimized.cl
+++ b/OpenCL/m13100_a3-optimized.cl
@@ -319,7 +319,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w0, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w0, edata2_left & 0xf);
 
     append_0x80_1x4 (w0, edata2_left & 0xf);
 
@@ -333,7 +333,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w1, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w1, edata2_left & 0xf);
 
     append_0x80_1x4 (w1, edata2_left & 0xf);
 
@@ -348,7 +348,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w2, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w2, edata2_left & 0xf);
 
     append_0x80_1x4 (w2, edata2_left & 0xf);
 
@@ -364,7 +364,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w3, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w3, edata2_left & 0xf);
 
     append_0x80_1x4 (w3, edata2_left & 0xf);
 
diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3.cl
index 78c52ea02..a13dcc8ea 100644
--- a/OpenCL/m13100_a3.cl
+++ b/OpenCL/m13100_a3.cl
@@ -232,14 +232,14 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w0, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w0, edata2_left & 0xf);
   }
   else if (edata2_left < 32)
   {
     j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w1, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w1, edata2_left & 0xf);
   }
   else if (edata2_left < 48)
   {
@@ -247,7 +247,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
     j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w2, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w2, edata2_left & 0xf);
   }
   else
   {
@@ -256,7 +256,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
     j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
     j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
 
-    truncate_block_4x4_le (w3, edata2_left & 0xf);
+    truncate_block_4x4_le_S (w3, edata2_left & 0xf);
   }
 
   md5_hmac_update_64 (&ctx, w0, w1, w2, w3, edata2_left);
diff --git a/OpenCL/m13400.cl b/OpenCL/m13400.cl
index 9fe6f7002..25a2db8b7 100644
--- a/OpenCL/m13400.cl
+++ b/OpenCL/m13400.cl
@@ -524,7 +524,7 @@ __kernel void m13400_comp (__global pw_t *pws, __constant const kernel_rule_t *r
 
       // we need to clear the buffer of the padding data
 
-      truncate_block_4x4_be (out, 16 - pad_byte);
+      truncate_block_4x4_be_S (out, 16 - pad_byte);
 
       u32 w0[4] = { 0 };
       u32 w1[4] = { 0 };
@@ -613,7 +613,7 @@ __kernel void m13400_comp (__global pw_t *pws, __constant const kernel_rule_t *r
 
       // we need to clear the buffer of the padding data
 
-      truncate_block_4x4_be (out, 16 - pad_byte);
+      truncate_block_4x4_be_S (out, 16 - pad_byte);
 
       u32 w0[4] = { 0 };
       u32 w1[4] = { 0 };