sha256_avx2.cpp raw

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