sha256_sse41.cpp raw

   1  // Copyright (c) 2018-2019 The Limenka developers
   2  // Distributed under the MIT software license, see the accompanying
   3  // file COPYING or http://www.opensource.org/licenses/mit-license.php.
   4  
   5  #ifdef ENABLE_SSE41
   6  
   7  #include <stdint.h>
   8  #include <immintrin.h>
   9  
  10  #include <attributes.h>
  11  #include <crypto/common.h>
  12  
  13  #if defined(__clang__)
  14  #pragma clang attribute push(__attribute__((__target__("sse4.1"))), apply_to = function)
  15  #elif defined(__GNUC__)
  16  #pragma GCC target ("sse4.1")
  17  #endif
  18  
  19  namespace sha256d64_sse41 {
  20  namespace {
  21  
  22  __m128i inline K(uint32_t x) { return _mm_set1_epi32(x); }
  23  
  24  __m128i inline Add(__m128i x, __m128i y) { return _mm_add_epi32(x, y); }
  25  __m128i inline Add(__m128i x, __m128i y, __m128i z) { return Add(Add(x, y), z); }
  26  __m128i inline Add(__m128i x, __m128i y, __m128i z, __m128i w) { return Add(Add(x, y), Add(z, w)); }
  27  __m128i inline Add(__m128i x, __m128i y, __m128i z, __m128i w, __m128i v) { return Add(Add(x, y, z), Add(w, v)); }
  28  __m128i inline Inc(__m128i& x, __m128i y) { x = Add(x, y); return x; }
  29  __m128i inline Inc(__m128i& x, __m128i y, __m128i z) { x = Add(x, y, z); return x; }
  30  __m128i inline Inc(__m128i& x, __m128i y, __m128i z, __m128i w) { x = Add(x, y, z, w); return x; }
  31  __m128i inline Xor(__m128i x, __m128i y) { return _mm_xor_si128(x, y); }
  32  __m128i inline Xor(__m128i x, __m128i y, __m128i z) { return Xor(Xor(x, y), z); }
  33  __m128i inline Or(__m128i x, __m128i y) { return _mm_or_si128(x, y); }
  34  __m128i inline And(__m128i x, __m128i y) { return _mm_and_si128(x, y); }
  35  __m128i inline ShR(__m128i x, int n) { return _mm_srli_epi32(x, n); }
  36  __m128i inline ShL(__m128i x, int n) { return _mm_slli_epi32(x, n); }
  37  
  38  __m128i inline Ch(__m128i x, __m128i y, __m128i z) { return Xor(z, And(x, Xor(y, z))); }
  39  __m128i inline Maj(__m128i x, __m128i y, __m128i z) { return Or(And(x, y), And(z, Or(x, y))); }
  40  __m128i inline Sigma0(__m128i x) { return Xor(Or(ShR(x, 2), ShL(x, 30)), Or(ShR(x, 13), ShL(x, 19)), Or(ShR(x, 22), ShL(x, 10))); }
  41  __m128i inline Sigma1(__m128i x) { return Xor(Or(ShR(x, 6), ShL(x, 26)), Or(ShR(x, 11), ShL(x, 21)), Or(ShR(x, 25), ShL(x, 7))); }
  42  __m128i inline sigma0(__m128i x) { return Xor(Or(ShR(x, 7), ShL(x, 25)), Or(ShR(x, 18), ShL(x, 14)), ShR(x, 3)); }
  43  __m128i inline sigma1(__m128i x) { return Xor(Or(ShR(x, 17), ShL(x, 15)), Or(ShR(x, 19), ShL(x, 13)), ShR(x, 10)); }
  44  
  45  /** One round of SHA-256. */
  46  void ALWAYS_INLINE Round(__m128i a, __m128i b, __m128i c, __m128i& d, __m128i e, __m128i f, __m128i g, __m128i& h, __m128i k)
  47  {
  48      __m128i t1 = Add(h, Sigma1(e), Ch(e, f, g), k);
  49      __m128i t2 = Add(Sigma0(a), Maj(a, b, c));
  50      d = Add(d, t1);
  51      h = Add(t1, t2);
  52  }
  53  
  54  __m128i inline Read4(const unsigned char* chunk, int offset) {
  55      __m128i ret = _mm_set_epi32(
  56          ReadLE32(chunk + 0 + offset),
  57          ReadLE32(chunk + 64 + offset),
  58          ReadLE32(chunk + 128 + offset),
  59          ReadLE32(chunk + 192 + offset)
  60      );
  61      return _mm_shuffle_epi8(ret, _mm_set_epi32(0x0C0D0E0FUL, 0x08090A0BUL, 0x04050607UL, 0x00010203UL));
  62  }
  63  
  64  void inline Write4(unsigned char* out, int offset, __m128i v) {
  65      v = _mm_shuffle_epi8(v, _mm_set_epi32(0x0C0D0E0FUL, 0x08090A0BUL, 0x04050607UL, 0x00010203UL));
  66      WriteLE32(out + 0 + offset, _mm_extract_epi32(v, 3));
  67      WriteLE32(out + 32 + offset, _mm_extract_epi32(v, 2));
  68      WriteLE32(out + 64 + offset, _mm_extract_epi32(v, 1));
  69      WriteLE32(out + 96 + offset, _mm_extract_epi32(v, 0));
  70  }
  71  
  72  }
  73  
  74  void Transform_4way(unsigned char* out, const unsigned char* in)
  75  {
  76      // Transform 1
  77      __m128i a = K(0x6a09e667ul);
  78      __m128i b = K(0xbb67ae85ul);
  79      __m128i c = K(0x3c6ef372ul);
  80      __m128i d = K(0xa54ff53aul);
  81      __m128i e = K(0x510e527ful);
  82      __m128i f = K(0x9b05688cul);
  83      __m128i g = K(0x1f83d9abul);
  84      __m128i h = K(0x5be0cd19ul);
  85  
  86      __m128i w0, w1, w2, w3, w4, w5, w6, w7, w8, w9, w10, w11, w12, w13, w14, w15;
  87  
  88      Round(a, b, c, d, e, f, g, h, Add(K(0x428a2f98ul), w0 = Read4(in, 0)));
  89      Round(h, a, b, c, d, e, f, g, Add(K(0x71374491ul), w1 = Read4(in, 4)));
  90      Round(g, h, a, b, c, d, e, f, Add(K(0xb5c0fbcful), w2 = Read4(in, 8)));
  91      Round(f, g, h, a, b, c, d, e, Add(K(0xe9b5dba5ul), w3 = Read4(in, 12)));
  92      Round(e, f, g, h, a, b, c, d, Add(K(0x3956c25bul), w4 = Read4(in, 16)));
  93      Round(d, e, f, g, h, a, b, c, Add(K(0x59f111f1ul), w5 = Read4(in, 20)));
  94      Round(c, d, e, f, g, h, a, b, Add(K(0x923f82a4ul), w6 = Read4(in, 24)));
  95      Round(b, c, d, e, f, g, h, a, Add(K(0xab1c5ed5ul), w7 = Read4(in, 28)));
  96      Round(a, b, c, d, e, f, g, h, Add(K(0xd807aa98ul), w8 = Read4(in, 32)));
  97      Round(h, a, b, c, d, e, f, g, Add(K(0x12835b01ul), w9 = Read4(in, 36)));
  98      Round(g, h, a, b, c, d, e, f, Add(K(0x243185beul), w10 = Read4(in, 40)));
  99      Round(f, g, h, a, b, c, d, e, Add(K(0x550c7dc3ul), w11 = Read4(in, 44)));
 100      Round(e, f, g, h, a, b, c, d, Add(K(0x72be5d74ul), w12 = Read4(in, 48)));
 101      Round(d, e, f, g, h, a, b, c, Add(K(0x80deb1feul), w13 = Read4(in, 52)));
 102      Round(c, d, e, f, g, h, a, b, Add(K(0x9bdc06a7ul), w14 = Read4(in, 56)));
 103      Round(b, c, d, e, f, g, h, a, Add(K(0xc19bf174ul), w15 = Read4(in, 60)));
 104      Round(a, b, c, d, e, f, g, h, Add(K(0xe49b69c1ul), Inc(w0, sigma1(w14), w9, sigma0(w1))));
 105      Round(h, a, b, c, d, e, f, g, Add(K(0xefbe4786ul), Inc(w1, sigma1(w15), w10, sigma0(w2))));
 106      Round(g, h, a, b, c, d, e, f, Add(K(0x0fc19dc6ul), Inc(w2, sigma1(w0), w11, sigma0(w3))));
 107      Round(f, g, h, a, b, c, d, e, Add(K(0x240ca1ccul), Inc(w3, sigma1(w1), w12, sigma0(w4))));
 108      Round(e, f, g, h, a, b, c, d, Add(K(0x2de92c6ful), Inc(w4, sigma1(w2), w13, sigma0(w5))));
 109      Round(d, e, f, g, h, a, b, c, Add(K(0x4a7484aaul), Inc(w5, sigma1(w3), w14, sigma0(w6))));
 110      Round(c, d, e, f, g, h, a, b, Add(K(0x5cb0a9dcul), Inc(w6, sigma1(w4), w15, sigma0(w7))));
 111      Round(b, c, d, e, f, g, h, a, Add(K(0x76f988daul), Inc(w7, sigma1(w5), w0, sigma0(w8))));
 112      Round(a, b, c, d, e, f, g, h, Add(K(0x983e5152ul), Inc(w8, sigma1(w6), w1, sigma0(w9))));
 113      Round(h, a, b, c, d, e, f, g, Add(K(0xa831c66dul), Inc(w9, sigma1(w7), w2, sigma0(w10))));
 114      Round(g, h, a, b, c, d, e, f, Add(K(0xb00327c8ul), Inc(w10, sigma1(w8), w3, sigma0(w11))));
 115      Round(f, g, h, a, b, c, d, e, Add(K(0xbf597fc7ul), Inc(w11, sigma1(w9), w4, sigma0(w12))));
 116      Round(e, f, g, h, a, b, c, d, Add(K(0xc6e00bf3ul), Inc(w12, sigma1(w10), w5, sigma0(w13))));
 117      Round(d, e, f, g, h, a, b, c, Add(K(0xd5a79147ul), Inc(w13, sigma1(w11), w6, sigma0(w14))));
 118      Round(c, d, e, f, g, h, a, b, Add(K(0x06ca6351ul), Inc(w14, sigma1(w12), w7, sigma0(w15))));
 119      Round(b, c, d, e, f, g, h, a, Add(K(0x14292967ul), Inc(w15, sigma1(w13), w8, sigma0(w0))));
 120      Round(a, b, c, d, e, f, g, h, Add(K(0x27b70a85ul), Inc(w0, sigma1(w14), w9, sigma0(w1))));
 121      Round(h, a, b, c, d, e, f, g, Add(K(0x2e1b2138ul), Inc(w1, sigma1(w15), w10, sigma0(w2))));
 122      Round(g, h, a, b, c, d, e, f, Add(K(0x4d2c6dfcul), Inc(w2, sigma1(w0), w11, sigma0(w3))));
 123      Round(f, g, h, a, b, c, d, e, Add(K(0x53380d13ul), Inc(w3, sigma1(w1), w12, sigma0(w4))));
 124      Round(e, f, g, h, a, b, c, d, Add(K(0x650a7354ul), Inc(w4, sigma1(w2), w13, sigma0(w5))));
 125      Round(d, e, f, g, h, a, b, c, Add(K(0x766a0abbul), Inc(w5, sigma1(w3), w14, sigma0(w6))));
 126      Round(c, d, e, f, g, h, a, b, Add(K(0x81c2c92eul), Inc(w6, sigma1(w4), w15, sigma0(w7))));
 127      Round(b, c, d, e, f, g, h, a, Add(K(0x92722c85ul), Inc(w7, sigma1(w5), w0, sigma0(w8))));
 128      Round(a, b, c, d, e, f, g, h, Add(K(0xa2bfe8a1ul), Inc(w8, sigma1(w6), w1, sigma0(w9))));
 129      Round(h, a, b, c, d, e, f, g, Add(K(0xa81a664bul), Inc(w9, sigma1(w7), w2, sigma0(w10))));
 130      Round(g, h, a, b, c, d, e, f, Add(K(0xc24b8b70ul), Inc(w10, sigma1(w8), w3, sigma0(w11))));
 131      Round(f, g, h, a, b, c, d, e, Add(K(0xc76c51a3ul), Inc(w11, sigma1(w9), w4, sigma0(w12))));
 132      Round(e, f, g, h, a, b, c, d, Add(K(0xd192e819ul), Inc(w12, sigma1(w10), w5, sigma0(w13))));
 133      Round(d, e, f, g, h, a, b, c, Add(K(0xd6990624ul), Inc(w13, sigma1(w11), w6, sigma0(w14))));
 134      Round(c, d, e, f, g, h, a, b, Add(K(0xf40e3585ul), Inc(w14, sigma1(w12), w7, sigma0(w15))));
 135      Round(b, c, d, e, f, g, h, a, Add(K(0x106aa070ul), Inc(w15, sigma1(w13), w8, sigma0(w0))));
 136      Round(a, b, c, d, e, f, g, h, Add(K(0x19a4c116ul), Inc(w0, sigma1(w14), w9, sigma0(w1))));
 137      Round(h, a, b, c, d, e, f, g, Add(K(0x1e376c08ul), Inc(w1, sigma1(w15), w10, sigma0(w2))));
 138      Round(g, h, a, b, c, d, e, f, Add(K(0x2748774cul), Inc(w2, sigma1(w0), w11, sigma0(w3))));
 139      Round(f, g, h, a, b, c, d, e, Add(K(0x34b0bcb5ul), Inc(w3, sigma1(w1), w12, sigma0(w4))));
 140      Round(e, f, g, h, a, b, c, d, Add(K(0x391c0cb3ul), Inc(w4, sigma1(w2), w13, sigma0(w5))));
 141      Round(d, e, f, g, h, a, b, c, Add(K(0x4ed8aa4aul), Inc(w5, sigma1(w3), w14, sigma0(w6))));
 142      Round(c, d, e, f, g, h, a, b, Add(K(0x5b9cca4ful), Inc(w6, sigma1(w4), w15, sigma0(w7))));
 143      Round(b, c, d, e, f, g, h, a, Add(K(0x682e6ff3ul), Inc(w7, sigma1(w5), w0, sigma0(w8))));
 144      Round(a, b, c, d, e, f, g, h, Add(K(0x748f82eeul), Inc(w8, sigma1(w6), w1, sigma0(w9))));
 145      Round(h, a, b, c, d, e, f, g, Add(K(0x78a5636ful), Inc(w9, sigma1(w7), w2, sigma0(w10))));
 146      Round(g, h, a, b, c, d, e, f, Add(K(0x84c87814ul), Inc(w10, sigma1(w8), w3, sigma0(w11))));
 147      Round(f, g, h, a, b, c, d, e, Add(K(0x8cc70208ul), Inc(w11, sigma1(w9), w4, sigma0(w12))));
 148      Round(e, f, g, h, a, b, c, d, Add(K(0x90befffaul), Inc(w12, sigma1(w10), w5, sigma0(w13))));
 149      Round(d, e, f, g, h, a, b, c, Add(K(0xa4506cebul), Inc(w13, sigma1(w11), w6, sigma0(w14))));
 150      Round(c, d, e, f, g, h, a, b, Add(K(0xbef9a3f7ul), Inc(w14, sigma1(w12), w7, sigma0(w15))));
 151      Round(b, c, d, e, f, g, h, a, Add(K(0xc67178f2ul), Inc(w15, sigma1(w13), w8, sigma0(w0))));
 152  
 153      a = Add(a, K(0x6a09e667ul));
 154      b = Add(b, K(0xbb67ae85ul));
 155      c = Add(c, K(0x3c6ef372ul));
 156      d = Add(d, K(0xa54ff53aul));
 157      e = Add(e, K(0x510e527ful));
 158      f = Add(f, K(0x9b05688cul));
 159      g = Add(g, K(0x1f83d9abul));
 160      h = Add(h, K(0x5be0cd19ul));
 161  
 162      __m128i t0 = a, t1 = b, t2 = c, t3 = d, t4 = e, t5 = f, t6 = g, t7 = h;
 163  
 164      // Transform 2
 165      Round(a, b, c, d, e, f, g, h, K(0xc28a2f98ul));
 166      Round(h, a, b, c, d, e, f, g, K(0x71374491ul));
 167      Round(g, h, a, b, c, d, e, f, K(0xb5c0fbcful));
 168      Round(f, g, h, a, b, c, d, e, K(0xe9b5dba5ul));
 169      Round(e, f, g, h, a, b, c, d, K(0x3956c25bul));
 170      Round(d, e, f, g, h, a, b, c, K(0x59f111f1ul));
 171      Round(c, d, e, f, g, h, a, b, K(0x923f82a4ul));
 172      Round(b, c, d, e, f, g, h, a, K(0xab1c5ed5ul));
 173      Round(a, b, c, d, e, f, g, h, K(0xd807aa98ul));
 174      Round(h, a, b, c, d, e, f, g, K(0x12835b01ul));
 175      Round(g, h, a, b, c, d, e, f, K(0x243185beul));
 176      Round(f, g, h, a, b, c, d, e, K(0x550c7dc3ul));
 177      Round(e, f, g, h, a, b, c, d, K(0x72be5d74ul));
 178      Round(d, e, f, g, h, a, b, c, K(0x80deb1feul));
 179      Round(c, d, e, f, g, h, a, b, K(0x9bdc06a7ul));
 180      Round(b, c, d, e, f, g, h, a, K(0xc19bf374ul));
 181      Round(a, b, c, d, e, f, g, h, K(0x649b69c1ul));
 182      Round(h, a, b, c, d, e, f, g, K(0xf0fe4786ul));
 183      Round(g, h, a, b, c, d, e, f, K(0x0fe1edc6ul));
 184      Round(f, g, h, a, b, c, d, e, K(0x240cf254ul));
 185      Round(e, f, g, h, a, b, c, d, K(0x4fe9346ful));
 186      Round(d, e, f, g, h, a, b, c, K(0x6cc984beul));
 187      Round(c, d, e, f, g, h, a, b, K(0x61b9411eul));
 188      Round(b, c, d, e, f, g, h, a, K(0x16f988faul));
 189      Round(a, b, c, d, e, f, g, h, K(0xf2c65152ul));
 190      Round(h, a, b, c, d, e, f, g, K(0xa88e5a6dul));
 191      Round(g, h, a, b, c, d, e, f, K(0xb019fc65ul));
 192      Round(f, g, h, a, b, c, d, e, K(0xb9d99ec7ul));
 193      Round(e, f, g, h, a, b, c, d, K(0x9a1231c3ul));
 194      Round(d, e, f, g, h, a, b, c, K(0xe70eeaa0ul));
 195      Round(c, d, e, f, g, h, a, b, K(0xfdb1232bul));
 196      Round(b, c, d, e, f, g, h, a, K(0xc7353eb0ul));
 197      Round(a, b, c, d, e, f, g, h, K(0x3069bad5ul));
 198      Round(h, a, b, c, d, e, f, g, K(0xcb976d5ful));
 199      Round(g, h, a, b, c, d, e, f, K(0x5a0f118ful));
 200      Round(f, g, h, a, b, c, d, e, K(0xdc1eeefdul));
 201      Round(e, f, g, h, a, b, c, d, K(0x0a35b689ul));
 202      Round(d, e, f, g, h, a, b, c, K(0xde0b7a04ul));
 203      Round(c, d, e, f, g, h, a, b, K(0x58f4ca9dul));
 204      Round(b, c, d, e, f, g, h, a, K(0xe15d5b16ul));
 205      Round(a, b, c, d, e, f, g, h, K(0x007f3e86ul));
 206      Round(h, a, b, c, d, e, f, g, K(0x37088980ul));
 207      Round(g, h, a, b, c, d, e, f, K(0xa507ea32ul));
 208      Round(f, g, h, a, b, c, d, e, K(0x6fab9537ul));
 209      Round(e, f, g, h, a, b, c, d, K(0x17406110ul));
 210      Round(d, e, f, g, h, a, b, c, K(0x0d8cd6f1ul));
 211      Round(c, d, e, f, g, h, a, b, K(0xcdaa3b6dul));
 212      Round(b, c, d, e, f, g, h, a, K(0xc0bbbe37ul));
 213      Round(a, b, c, d, e, f, g, h, K(0x83613bdaul));
 214      Round(h, a, b, c, d, e, f, g, K(0xdb48a363ul));
 215      Round(g, h, a, b, c, d, e, f, K(0x0b02e931ul));
 216      Round(f, g, h, a, b, c, d, e, K(0x6fd15ca7ul));
 217      Round(e, f, g, h, a, b, c, d, K(0x521afacaul));
 218      Round(d, e, f, g, h, a, b, c, K(0x31338431ul));
 219      Round(c, d, e, f, g, h, a, b, K(0x6ed41a95ul));
 220      Round(b, c, d, e, f, g, h, a, K(0x6d437890ul));
 221      Round(a, b, c, d, e, f, g, h, K(0xc39c91f2ul));
 222      Round(h, a, b, c, d, e, f, g, K(0x9eccabbdul));
 223      Round(g, h, a, b, c, d, e, f, K(0xb5c9a0e6ul));
 224      Round(f, g, h, a, b, c, d, e, K(0x532fb63cul));
 225      Round(e, f, g, h, a, b, c, d, K(0xd2c741c6ul));
 226      Round(d, e, f, g, h, a, b, c, K(0x07237ea3ul));
 227      Round(c, d, e, f, g, h, a, b, K(0xa4954b68ul));
 228      Round(b, c, d, e, f, g, h, a, K(0x4c191d76ul));
 229  
 230      w0 = Add(t0, a);
 231      w1 = Add(t1, b);
 232      w2 = Add(t2, c);
 233      w3 = Add(t3, d);
 234      w4 = Add(t4, e);
 235      w5 = Add(t5, f);
 236      w6 = Add(t6, g);
 237      w7 = Add(t7, h);
 238  
 239      // Transform 3
 240      a = K(0x6a09e667ul);
 241      b = K(0xbb67ae85ul);
 242      c = K(0x3c6ef372ul);
 243      d = K(0xa54ff53aul);
 244      e = K(0x510e527ful);
 245      f = K(0x9b05688cul);
 246      g = K(0x1f83d9abul);
 247      h = K(0x5be0cd19ul);
 248  
 249      Round(a, b, c, d, e, f, g, h, Add(K(0x428a2f98ul), w0));
 250      Round(h, a, b, c, d, e, f, g, Add(K(0x71374491ul), w1));
 251      Round(g, h, a, b, c, d, e, f, Add(K(0xb5c0fbcful), w2));
 252      Round(f, g, h, a, b, c, d, e, Add(K(0xe9b5dba5ul), w3));
 253      Round(e, f, g, h, a, b, c, d, Add(K(0x3956c25bul), w4));
 254      Round(d, e, f, g, h, a, b, c, Add(K(0x59f111f1ul), w5));
 255      Round(c, d, e, f, g, h, a, b, Add(K(0x923f82a4ul), w6));
 256      Round(b, c, d, e, f, g, h, a, Add(K(0xab1c5ed5ul), w7));
 257      Round(a, b, c, d, e, f, g, h, K(0x5807aa98ul));
 258      Round(h, a, b, c, d, e, f, g, K(0x12835b01ul));
 259      Round(g, h, a, b, c, d, e, f, K(0x243185beul));
 260      Round(f, g, h, a, b, c, d, e, K(0x550c7dc3ul));
 261      Round(e, f, g, h, a, b, c, d, K(0x72be5d74ul));
 262      Round(d, e, f, g, h, a, b, c, K(0x80deb1feul));
 263      Round(c, d, e, f, g, h, a, b, K(0x9bdc06a7ul));
 264      Round(b, c, d, e, f, g, h, a, K(0xc19bf274ul));
 265      Round(a, b, c, d, e, f, g, h, Add(K(0xe49b69c1ul), Inc(w0, sigma0(w1))));
 266      Round(h, a, b, c, d, e, f, g, Add(K(0xefbe4786ul), Inc(w1, K(0xa00000ul), sigma0(w2))));
 267      Round(g, h, a, b, c, d, e, f, Add(K(0x0fc19dc6ul), Inc(w2, sigma1(w0), sigma0(w3))));
 268      Round(f, g, h, a, b, c, d, e, Add(K(0x240ca1ccul), Inc(w3, sigma1(w1), sigma0(w4))));
 269      Round(e, f, g, h, a, b, c, d, Add(K(0x2de92c6ful), Inc(w4, sigma1(w2), sigma0(w5))));
 270      Round(d, e, f, g, h, a, b, c, Add(K(0x4a7484aaul), Inc(w5, sigma1(w3), sigma0(w6))));
 271      Round(c, d, e, f, g, h, a, b, Add(K(0x5cb0a9dcul), Inc(w6, sigma1(w4), K(0x100ul), sigma0(w7))));
 272      Round(b, c, d, e, f, g, h, a, Add(K(0x76f988daul), Inc(w7, sigma1(w5), w0, K(0x11002000ul))));
 273      Round(a, b, c, d, e, f, g, h, Add(K(0x983e5152ul), w8 = Add(K(0x80000000ul), sigma1(w6), w1)));
 274      Round(h, a, b, c, d, e, f, g, Add(K(0xa831c66dul), w9 = Add(sigma1(w7), w2)));
 275      Round(g, h, a, b, c, d, e, f, Add(K(0xb00327c8ul), w10 = Add(sigma1(w8), w3)));
 276      Round(f, g, h, a, b, c, d, e, Add(K(0xbf597fc7ul), w11 = Add(sigma1(w9), w4)));
 277      Round(e, f, g, h, a, b, c, d, Add(K(0xc6e00bf3ul), w12 = Add(sigma1(w10), w5)));
 278      Round(d, e, f, g, h, a, b, c, Add(K(0xd5a79147ul), w13 = Add(sigma1(w11), w6)));
 279      Round(c, d, e, f, g, h, a, b, Add(K(0x06ca6351ul), w14 = Add(sigma1(w12), w7, K(0x400022ul))));
 280      Round(b, c, d, e, f, g, h, a, Add(K(0x14292967ul), w15 = Add(K(0x100ul), sigma1(w13), w8, sigma0(w0))));
 281      Round(a, b, c, d, e, f, g, h, Add(K(0x27b70a85ul), Inc(w0, sigma1(w14), w9, sigma0(w1))));
 282      Round(h, a, b, c, d, e, f, g, Add(K(0x2e1b2138ul), Inc(w1, sigma1(w15), w10, sigma0(w2))));
 283      Round(g, h, a, b, c, d, e, f, Add(K(0x4d2c6dfcul), Inc(w2, sigma1(w0), w11, sigma0(w3))));
 284      Round(f, g, h, a, b, c, d, e, Add(K(0x53380d13ul), Inc(w3, sigma1(w1), w12, sigma0(w4))));
 285      Round(e, f, g, h, a, b, c, d, Add(K(0x650a7354ul), Inc(w4, sigma1(w2), w13, sigma0(w5))));
 286      Round(d, e, f, g, h, a, b, c, Add(K(0x766a0abbul), Inc(w5, sigma1(w3), w14, sigma0(w6))));
 287      Round(c, d, e, f, g, h, a, b, Add(K(0x81c2c92eul), Inc(w6, sigma1(w4), w15, sigma0(w7))));
 288      Round(b, c, d, e, f, g, h, a, Add(K(0x92722c85ul), Inc(w7, sigma1(w5), w0, sigma0(w8))));
 289      Round(a, b, c, d, e, f, g, h, Add(K(0xa2bfe8a1ul), Inc(w8, sigma1(w6), w1, sigma0(w9))));
 290      Round(h, a, b, c, d, e, f, g, Add(K(0xa81a664bul), Inc(w9, sigma1(w7), w2, sigma0(w10))));
 291      Round(g, h, a, b, c, d, e, f, Add(K(0xc24b8b70ul), Inc(w10, sigma1(w8), w3, sigma0(w11))));
 292      Round(f, g, h, a, b, c, d, e, Add(K(0xc76c51a3ul), Inc(w11, sigma1(w9), w4, sigma0(w12))));
 293      Round(e, f, g, h, a, b, c, d, Add(K(0xd192e819ul), Inc(w12, sigma1(w10), w5, sigma0(w13))));
 294      Round(d, e, f, g, h, a, b, c, Add(K(0xd6990624ul), Inc(w13, sigma1(w11), w6, sigma0(w14))));
 295      Round(c, d, e, f, g, h, a, b, Add(K(0xf40e3585ul), Inc(w14, sigma1(w12), w7, sigma0(w15))));
 296      Round(b, c, d, e, f, g, h, a, Add(K(0x106aa070ul), Inc(w15, sigma1(w13), w8, sigma0(w0))));
 297      Round(a, b, c, d, e, f, g, h, Add(K(0x19a4c116ul), Inc(w0, sigma1(w14), w9, sigma0(w1))));
 298      Round(h, a, b, c, d, e, f, g, Add(K(0x1e376c08ul), Inc(w1, sigma1(w15), w10, sigma0(w2))));
 299      Round(g, h, a, b, c, d, e, f, Add(K(0x2748774cul), Inc(w2, sigma1(w0), w11, sigma0(w3))));
 300      Round(f, g, h, a, b, c, d, e, Add(K(0x34b0bcb5ul), Inc(w3, sigma1(w1), w12, sigma0(w4))));
 301      Round(e, f, g, h, a, b, c, d, Add(K(0x391c0cb3ul), Inc(w4, sigma1(w2), w13, sigma0(w5))));
 302      Round(d, e, f, g, h, a, b, c, Add(K(0x4ed8aa4aul), Inc(w5, sigma1(w3), w14, sigma0(w6))));
 303      Round(c, d, e, f, g, h, a, b, Add(K(0x5b9cca4ful), Inc(w6, sigma1(w4), w15, sigma0(w7))));
 304      Round(b, c, d, e, f, g, h, a, Add(K(0x682e6ff3ul), Inc(w7, sigma1(w5), w0, sigma0(w8))));
 305      Round(a, b, c, d, e, f, g, h, Add(K(0x748f82eeul), Inc(w8, sigma1(w6), w1, sigma0(w9))));
 306      Round(h, a, b, c, d, e, f, g, Add(K(0x78a5636ful), Inc(w9, sigma1(w7), w2, sigma0(w10))));
 307      Round(g, h, a, b, c, d, e, f, Add(K(0x84c87814ul), Inc(w10, sigma1(w8), w3, sigma0(w11))));
 308      Round(f, g, h, a, b, c, d, e, Add(K(0x8cc70208ul), Inc(w11, sigma1(w9), w4, sigma0(w12))));
 309      Round(e, f, g, h, a, b, c, d, Add(K(0x90befffaul), Inc(w12, sigma1(w10), w5, sigma0(w13))));
 310      Round(d, e, f, g, h, a, b, c, Add(K(0xa4506cebul), Inc(w13, sigma1(w11), w6, sigma0(w14))));
 311      Round(c, d, e, f, g, h, a, b, Add(K(0xbef9a3f7ul), w14, sigma1(w12), w7, sigma0(w15)));
 312      Round(b, c, d, e, f, g, h, a, Add(K(0xc67178f2ul), w15, sigma1(w13), w8, sigma0(w0)));
 313  
 314      // Output
 315      Write4(out, 0, Add(a, K(0x6a09e667ul)));
 316      Write4(out, 4, Add(b, K(0xbb67ae85ul)));
 317      Write4(out, 8, Add(c, K(0x3c6ef372ul)));
 318      Write4(out, 12, Add(d, K(0xa54ff53aul)));
 319      Write4(out, 16, Add(e, K(0x510e527ful)));
 320      Write4(out, 20, Add(f, K(0x9b05688cul)));
 321      Write4(out, 24, Add(g, K(0x1f83d9abul)));
 322      Write4(out, 28, Add(h, K(0x5be0cd19ul)));
 323  }
 324  
 325  }
 326  
 327  #if defined(__clang__)
 328  #pragma clang attribute pop
 329  #endif
 330  
 331  #endif
 332