bitcracker/src_CUDA/sha256_header.h

473 lines
20 KiB
C
Executable File

/*
* BitCracker: BitLocker password cracking tool, CUDA version.
* Copyright (C) 2013-2017 Elena Ago <elena dot ago at gmail dot com>
* Massimo Bernaschi <massimo dot bernaschi at gmail dot com>
*
* This file is part of the BitCracker project: https://github.com/e-ago/bitcracker
*
* BitCracker is free software: you can redistribute it and/or modify
* it under the terms of the GNU General Public License as published by
* the Free Software Foundation, either version 2 of the License, or
* (at your option) any later version.
*
* BitCracker is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* along with BitCracker. If not, see <http://www.gnu.org/licenses/>.
*/
#define ROR(x, i) (((x) << (32 - (i))) | ((x) >> (i)))
#define BLOCK_SIZE 64
#define ROR7(x) (((x) << 25) | ((x) >> 7))
#define ROR18(x) (((x) << 14) | ((x) >> 18))
#define ROR17(x) (((x) << 15) | ((x) >> 17))
#define ROR19(x) (((x) << 13) | ((x) >> 19))
#if (__CUDA_ARCH__ >= 500)
static __device__ __forceinline__ uint32_t IADD3(uint32_t a, uint32_t b, uint32_t c) {
uint32_t d;
asm("iadd3 %0, %1, %2, %3;" : "=r"(d) : "r"(a), "r"(b), "r"(c));
return d;
}
static __device__ __forceinline__ uint32_t LOP3LUT_XOR(uint32_t a, uint32_t b, uint32_t c) {
uint32_t d;
asm("lop3.b32 %0, %1, %2, %3, 0x96;" : "=r"(d) : "r"(a), "r"(b), "r"(c));
return d;
}
static __device__ __forceinline__ uint32_t LOP3LUT_XORAND(uint32_t a, uint32_t b, uint32_t c) {
uint32_t d;
asm("lop3.b32 %0, %1, %2, %3, 0xb8;" : "=r"(d) : "r"(a), "r"(b), "r"(c));
return d;
}
static __device__ __forceinline__ uint32_t LOP3LUT_ANDOR(uint32_t a, uint32_t b, uint32_t c) {
uint32_t d;
asm("lop3.b32 %0, %1, %2, %3, 0xe8;" : "=r"(d) : "r"(a), "r"(b), "r"(c));
return d;
}
#else
static __device__ __forceinline__ uint32_t IADD3(uint32_t a, uint32_t b, uint32_t c) {
return a+b+c;
}
static __device__ __forceinline__ uint32_t LOP3LUT_XOR(uint32_t a, uint32_t b, uint32_t c) {
return a^b^c;
}
static __device__ __forceinline__ uint32_t LOP3LUT_XORAND(uint32_t g, uint32_t e, uint32_t f) {
return (g ^ (e & (f ^ g)));
}
static __device__ __forceinline__ uint32_t LOP3LUT_ANDOR(uint32_t a, uint32_t b, uint32_t c) {
return ((a & (b | c)) | (b & c));
}
#endif
#define SCHEDULE0() \
schedule0 = schedule16 + schedule25 \
+ LOP3LUT_XOR(ROR7(schedule17) , ROR18(schedule17) , (schedule17 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule30) , ROR19(schedule30) , (schedule30 >> 10));
#define SCHEDULE1() \
schedule1 = schedule17 + schedule26 \
+ LOP3LUT_XOR(ROR7(schedule18) , ROR18(schedule18) , (schedule18 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule31) , ROR19(schedule31) , (schedule31 >> 10));
#define SCHEDULE2() \
schedule2 = schedule18 + schedule27 \
+ LOP3LUT_XOR(ROR7(schedule19) , ROR18(schedule19) , (schedule19 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule0) , ROR19(schedule0) , (schedule0 >> 10));
#define SCHEDULE3() \
schedule3 = schedule19 + schedule28 \
+ LOP3LUT_XOR(ROR7(schedule20) , ROR18(schedule20) , (schedule20 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule1) , ROR19(schedule1) , (schedule1 >> 10));
#define SCHEDULE4() \
schedule4 = schedule20 + schedule29 \
+ LOP3LUT_XOR(ROR7(schedule21) , ROR18(schedule21) , (schedule21 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule2) , ROR19(schedule2) , (schedule2 >> 10));
#define SCHEDULE5() \
schedule5 = schedule21 + schedule30 \
+ LOP3LUT_XOR(ROR7(schedule22) , ROR18(schedule22) , (schedule22 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule3) , ROR19(schedule3) , (schedule3 >> 10));
#define SCHEDULE6() \
schedule6 = schedule22 + schedule31 \
+ LOP3LUT_XOR(ROR7(schedule23) , ROR18(schedule23) , (schedule23 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule4) , ROR19(schedule4) , (schedule4 >> 10));
#define SCHEDULE7() \
schedule7 = schedule23 + schedule0 \
+ LOP3LUT_XOR(ROR7(schedule24) , ROR18(schedule24) , (schedule24 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule5) , ROR19(schedule5) , (schedule5 >> 10));
#define SCHEDULE8() \
schedule8 = schedule24 + schedule1 \
+ LOP3LUT_XOR(ROR7(schedule25) , ROR18(schedule25) , (schedule25 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule6) , ROR19(schedule6) , (schedule6 >> 10));
#define SCHEDULE9() \
schedule9 = schedule25 + schedule2 \
+ LOP3LUT_XOR(ROR7(schedule26) , ROR18(schedule26) , (schedule26 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule7) , ROR19(schedule7) , (schedule7 >> 10));
#define SCHEDULE10() \
schedule10 = schedule26 + schedule3 \
+ LOP3LUT_XOR(ROR7(schedule27) , ROR18(schedule27) , (schedule27 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule8) , ROR19(schedule8) , (schedule8 >> 10));
#define SCHEDULE11() \
schedule11 = schedule27 + schedule4 \
+ LOP3LUT_XOR(ROR7(schedule28) , ROR18(schedule28) , (schedule28 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule9) , ROR19(schedule9) , (schedule9 >> 10));
#define SCHEDULE12() \
schedule12 = schedule28 + schedule5 \
+ LOP3LUT_XOR(ROR7(schedule29) , ROR18(schedule29) , (schedule29 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule10) , ROR19(schedule10) , (schedule10 >> 10));
#define SCHEDULE13() \
schedule13 = schedule29 + schedule6 \
+ LOP3LUT_XOR(ROR7(schedule30) , ROR18(schedule30) , (schedule30 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule11) , ROR19(schedule11) , (schedule11 >> 10));
#define SCHEDULE14() \
schedule14 = schedule30 + schedule7 \
+ LOP3LUT_XOR(ROR7(schedule31) , ROR18(schedule31) , (schedule31 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule12) , ROR19(schedule12) , (schedule12 >> 10));
#define SCHEDULE15() \
schedule15 = schedule31 + schedule8 \
+ LOP3LUT_XOR(ROR7(schedule0) , ROR18(schedule0) , (schedule0 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule13) , ROR19(schedule13) , (schedule13 >> 10));
#define SCHEDULE16() \
schedule16 = schedule0 + schedule9 \
+ LOP3LUT_XOR( ROR7(schedule1), ROR18(schedule1), (schedule1 >> 3)) \
+ LOP3LUT_XOR( ROR17(schedule14), ROR19(schedule14), (schedule14 >> 10));
#define SCHEDULE17() \
schedule17 = schedule1 + schedule10 \
+ LOP3LUT_XOR(ROR7(schedule2) , ROR18(schedule2) , (schedule2 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule15) , ROR19(schedule15) , (schedule15 >> 10));
#define SCHEDULE18() \
schedule18 = schedule2 + schedule11 \
+ LOP3LUT_XOR(ROR7(schedule3) ,ROR18(schedule3) ,(schedule3 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule16), ROR19(schedule16), (schedule16 >> 10));
#define SCHEDULE19() \
schedule19 = schedule3 + schedule12 \
+ LOP3LUT_XOR(ROR7(schedule4) , ROR18(schedule4) , (schedule4 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule17) , ROR19(schedule17) , (schedule17 >> 10));
#define SCHEDULE20() \
schedule20 = schedule4 + schedule13 \
+ LOP3LUT_XOR(ROR7(schedule5) , ROR18(schedule5) , (schedule5 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule18) , ROR19(schedule18) , (schedule18 >> 10));
#define SCHEDULE21() \
schedule21 = schedule5 + schedule14 \
+ LOP3LUT_XOR(ROR7(schedule6) , ROR18(schedule6) , (schedule6 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule19) , ROR19(schedule19) , (schedule19 >> 10));
#define SCHEDULE22() \
schedule22 = schedule6 + schedule15 \
+ LOP3LUT_XOR(ROR7(schedule7) , ROR18(schedule7) , (schedule7 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule20) , ROR19(schedule20) , (schedule20 >> 10));
#define SCHEDULE23() \
schedule23 = schedule7 + schedule16 \
+ LOP3LUT_XOR(ROR7(schedule8) , ROR18(schedule8) , (schedule8 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule21) , ROR19(schedule21) , (schedule21 >> 10));
#define SCHEDULE24() \
schedule24 = schedule8 + schedule17 \
+ LOP3LUT_XOR(ROR7(schedule9) , ROR18(schedule9) , (schedule9 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule22) , ROR19(schedule22) , (schedule22 >> 10));
#define SCHEDULE25() \
schedule25 = schedule9 + schedule18 \
+ LOP3LUT_XOR(ROR7(schedule10) , ROR18(schedule10) , (schedule10 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule23) , ROR19(schedule23) , (schedule23 >> 10));
#define SCHEDULE26() \
schedule26 = schedule10 + schedule19 \
+ LOP3LUT_XOR(ROR7(schedule11) , ROR18(schedule11) , (schedule11 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule24) , ROR19(schedule24) , (schedule24 >> 10));
#define SCHEDULE27() \
schedule27 = schedule11 + schedule20 \
+ LOP3LUT_XOR(ROR7(schedule12) , ROR18(schedule12) , (schedule12 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule25) , ROR19(schedule25) , (schedule25 >> 10));
#define SCHEDULE28() \
schedule28 = schedule12 + schedule21 \
+ LOP3LUT_XOR(ROR7(schedule13) , ROR18(schedule13) , (schedule13 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule26) , ROR19(schedule26) , (schedule26 >> 10));
#define SCHEDULE29() \
schedule29 = schedule13 + schedule22 \
+ LOP3LUT_XOR(ROR7(schedule14) , ROR18(schedule14) , (schedule14 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule27) , ROR19(schedule27) , (schedule27 >> 10));
#define SCHEDULE30() \
schedule30 = schedule14 + schedule23 \
+ LOP3LUT_XOR(ROR7(schedule15) , ROR18(schedule15) , (schedule15 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule28) , ROR19(schedule28) , (schedule28 >> 10));
#define SCHEDULE31() \
schedule31 = schedule15 + schedule24 \
+ LOP3LUT_XOR(ROR7(schedule16) , ROR18(schedule16) , (schedule16 >> 3)) \
+ LOP3LUT_XOR(ROR17(schedule29) , ROR19(schedule29) , (schedule29 >> 10));
#define ALL_SCHEDULE32() \
SCHEDULE0() \
SCHEDULE1() \
SCHEDULE2() \
SCHEDULE3() \
SCHEDULE4() \
SCHEDULE5() \
SCHEDULE6() \
SCHEDULE7() \
SCHEDULE8() \
SCHEDULE9() \
SCHEDULE10() \
SCHEDULE11() \
SCHEDULE12() \
SCHEDULE13() \
SCHEDULE14() \
SCHEDULE15() \
SCHEDULE16() \
SCHEDULE17() \
SCHEDULE18() \
SCHEDULE19() \
SCHEDULE20() \
SCHEDULE21() \
SCHEDULE22() \
SCHEDULE23() \
SCHEDULE24() \
SCHEDULE25() \
SCHEDULE26() \
SCHEDULE27() \
SCHEDULE28() \
SCHEDULE29() \
SCHEDULE30() \
SCHEDULE31()
#define ALL_SCHEDULE_LAST16() \
SCHEDULE16() \
SCHEDULE17() \
SCHEDULE18() \
SCHEDULE19() \
SCHEDULE20() \
SCHEDULE21() \
SCHEDULE22() \
SCHEDULE23() \
SCHEDULE24() \
SCHEDULE25() \
SCHEDULE26() \
SCHEDULE27() \
SCHEDULE28() \
SCHEDULE29() \
SCHEDULE30() \
SCHEDULE31()
#define ROR6(x) (((x) << 26) | ((x) >> 6))
#define ROR11(x) (((x) << 21) | ((x) >> 11))
#define ROR25(x) (((x) << 7) | ((x) >> 25))
#define ROR2(x) (((x) << 30) | ((x) >> 2))
#define ROR13(x) (((x) << 19) | ((x) >> 13))
#define ROR22(x) (((x) << 10) | ((x) >> 22))
#define ROUND(a, b, c, d, e, f, g, h, W, k) \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + W; \
d += h; \
h += LOP3LUT_XOR(ROR2(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
#define ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, i, k, indexW) \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + tex1Dfetch(w_texture, (indexW+i)); \
d += h; \
h += LOP3LUT_XOR(ROR2(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
#define ROUND_SECOND_BLOCK_CONST(a, b, c, d, e, f, g, h, i, k, w) \
h += LOP3LUT_XOR(ROR6(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + w; \
d += h; \
h += LOP3LUT_XOR(ROR2(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
#define LOADSCHEDULE_WPRE(i, j) \
w_blocks_d[j] = \
(uint32_t)block[i * 4 + 0] << 24 \
| (uint32_t)block[i * 4 + 1] << 16 \
| (uint32_t)block[i * 4 + 2] << 8 \
| (uint32_t)block[i * 4 + 3];
#define SCHEDULE_WPRE(i) \
w_blocks_d[i] = w_blocks_d[i - 16] + w_blocks_d[i - 7] \
+ (ROR(w_blocks_d[i - 15], 7) ^ ROR(w_blocks_d[i - 15], 18) ^ (w_blocks_d[i - 15] >> 3)) \
+ (ROR(w_blocks_d[i - 2], 17) ^ ROR(w_blocks_d[i - 2], 19) ^ (w_blocks_d[i - 2] >> 10));
#define ALL_ROUND_B1_1() \
ROUND(a, b, c, d, e, f, g, h, schedule0, 0x428A2F98) \
ROUND(h, a, b, c, d, e, f, g, schedule1, 0x71374491) \
ROUND(g, h, a, b, c, d, e, f, schedule2, 0xB5C0FBCF) \
ROUND(f, g, h, a, b, c, d, e, schedule3, 0xE9B5DBA5) \
ROUND(e, f, g, h, a, b, c, d, schedule4, 0x3956C25B) \
ROUND(d, e, f, g, h, a, b, c, schedule5, 0x59F111F1) \
ROUND(c, d, e, f, g, h, a, b, schedule6, 0x923F82A4) \
ROUND(b, c, d, e, f, g, h, a, schedule7, 0xAB1C5ED5) \
ROUND(a, b, c, d, e, f, g, h, schedule8, 0xD807AA98) \
ROUND(h, a, b, c, d, e, f, g, schedule9, 0x12835B01) \
ROUND(g, h, a, b, c, d, e, f, schedule10, 0x243185BE) \
ROUND(f, g, h, a, b, c, d, e, schedule11, 0x550C7DC3) \
ROUND(e, f, g, h, a, b, c, d, schedule12, 0x72BE5D74) \
ROUND(d, e, f, g, h, a, b, c, schedule13, 0x80DEB1FE) \
ROUND(c, d, e, f, g, h, a, b, schedule14, 0x9BDC06A7) \
ROUND(b, c, d, e, f, g, h, a, schedule15, 0xC19BF174) \
ROUND(a, b, c, d, e, f, g, h, schedule16, 0xE49B69C1) \
ROUND(h, a, b, c, d, e, f, g, schedule17, 0xEFBE4786) \
ROUND(g, h, a, b, c, d, e, f, schedule18, 0x0FC19DC6) \
ROUND(f, g, h, a, b, c, d, e, schedule19, 0x240CA1CC) \
ROUND(e, f, g, h, a, b, c, d, schedule20, 0x2DE92C6F) \
ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4A7484AA) \
ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5CB0A9DC) \
ROUND(b, c, d, e, f, g, h, a, schedule23, 0x76F988DA) \
ROUND(a, b, c, d, e, f, g, h, schedule24, 0x983E5152) \
ROUND(h, a, b, c, d, e, f, g, schedule25, 0xA831C66D) \
ROUND(g, h, a, b, c, d, e, f, schedule26, 0xB00327C8) \
ROUND(f, g, h, a, b, c, d, e, schedule27, 0xBF597FC7) \
ROUND(e, f, g, h, a, b, c, d, schedule28, 0xC6E00BF3) \
ROUND(d, e, f, g, h, a, b, c, schedule29, 0xD5A79147) \
ROUND(c, d, e, f, g, h, a, b, schedule30, 0x06CA6351) \
ROUND(b, c, d, e, f, g, h, a, schedule31, 0x14292967)
#define ALL_ROUND_B1_2() \
ROUND(a, b, c, d, e, f, g, h, schedule0, 0x27B70A85) \
ROUND(h, a, b, c, d, e, f, g, schedule1, 0x2E1B2138) \
ROUND(g, h, a, b, c, d, e, f, schedule2, 0x4D2C6DFC) \
ROUND(f, g, h, a, b, c, d, e, schedule3, 0x53380D13) \
ROUND(e, f, g, h, a, b, c, d, schedule4, 0x650A7354) \
ROUND(d, e, f, g, h, a, b, c, schedule5, 0x766A0ABB) \
ROUND(c, d, e, f, g, h, a, b, schedule6, 0x81C2C92E) \
ROUND(b, c, d, e, f, g, h, a, schedule7, 0x92722C85) \
ROUND(a, b, c, d, e, f, g, h, schedule8, 0xA2BFE8A1) \
ROUND(h, a, b, c, d, e, f, g, schedule9, 0xA81A664B) \
ROUND(g, h, a, b, c, d, e, f, schedule10, 0xC24B8B70) \
ROUND(f, g, h, a, b, c, d, e, schedule11, 0xC76C51A3) \
ROUND(e, f, g, h, a, b, c, d, schedule12, 0xD192E819) \
ROUND(d, e, f, g, h, a, b, c, schedule13, 0xD6990624) \
ROUND(c, d, e, f, g, h, a, b, schedule14, 0xF40E3585) \
ROUND(b, c, d, e, f, g, h, a, schedule15, 0x106AA070) \
ROUND(a, b, c, d, e, f, g, h, schedule16, 0x19A4C116) \
ROUND(h, a, b, c, d, e, f, g, schedule17, 0x1E376C08) \
ROUND(g, h, a, b, c, d, e, f, schedule18, 0x2748774C) \
ROUND(f, g, h, a, b, c, d, e, schedule19, 0x34B0BCB5) \
ROUND(e, f, g, h, a, b, c, d, schedule20, 0x391C0CB3) \
ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4ED8AA4A) \
ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5B9CCA4F) \
ROUND(b, c, d, e, f, g, h, a, schedule23, 0x682E6FF3) \
ROUND(a, b, c, d, e, f, g, h, schedule24, 0x748F82EE) \
ROUND(h, a, b, c, d, e, f, g, schedule25, 0x78A5636F) \
ROUND(g, h, a, b, c, d, e, f, schedule26, 0x84C87814) \
ROUND(f, g, h, a, b, c, d, e, schedule27, 0x8CC70208) \
ROUND(e, f, g, h, a, b, c, d, schedule28, 0x90BEFFFA) \
ROUND(d, e, f, g, h, a, b, c, schedule29, 0xA4506CEB) \
ROUND(c, d, e, f, g, h, a, b, schedule30, 0xBEF9A3F7) \
ROUND(b, c, d, e, f, g, h, a, schedule31, 0xC67178F2)
#define ALL_ROUND_B2() \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 0, 0x428A2F98, 0) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 1, 0x71374491, 0) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 2, 0xB5C0FBCF, 0) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 3, 0xE9B5DBA5, 0) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 4, 0x3956C25B, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 5, 0x59F111F1, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 6, 0x923F82A4, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 7, 0xAB1C5ED5, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 8, 0xD807AA98, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 9, 0x12835B01, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 10, 0x243185BE, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 11, 0x550C7DC3, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 12, 0x72BE5D74, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 13, 0x80DEB1FE, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 14, 0x9BDC06A7, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 15, 0xC19BF174, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 16, 0xE49B69C1, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 17, 0xEFBE4786, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 18, 0x0FC19DC6, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 19, 0x240CA1CC, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 20, 0x2DE92C6F, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 21, 0x4A7484AA, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 22, 0x5CB0A9DC, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 23, 0x76F988DA, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 24, 0x983E5152, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 25, 0xA831C66D, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 26, 0xB00327C8, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 27, 0xBF597FC7, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 28, 0xC6E00BF3, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 29, 0xD5A79147, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 30, 0x06CA6351, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 31, 0x14292967, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 32, 0x27B70A85, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 33, 0x2E1B2138, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 34, 0x4D2C6DFC, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 35, 0x53380D13, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 36, 0x650A7354, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 37, 0x766A0ABB, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 38, 0x81C2C92E, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 39, 0x92722C85, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 40, 0xA2BFE8A1, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 41, 0xA81A664B, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 42, 0xC24B8B70, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 43, 0xC76C51A3, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 44, 0xD192E819, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 45, 0xD6990624, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 46, 0xF40E3585, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 47, 0x106AA070, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 48, 0x19A4C116, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 49, 0x1E376C08, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 50, 0x2748774C, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 51, 0x34B0BCB5, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 52, 0x391C0CB3, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 53, 0x4ED8AA4A, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 54, 0x5B9CCA4F, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 55, 0x682E6FF3, indexW) \
ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, 56, 0x748F82EE, indexW) \
ROUND_SECOND_BLOCK(h, a, b, c, d, e, f, g, 57, 0x78A5636F, indexW) \
ROUND_SECOND_BLOCK(g, h, a, b, c, d, e, f, 58, 0x84C87814, indexW) \
ROUND_SECOND_BLOCK(f, g, h, a, b, c, d, e, 59, 0x8CC70208, indexW) \
ROUND_SECOND_BLOCK(e, f, g, h, a, b, c, d, 60, 0x90BEFFFA, indexW) \
ROUND_SECOND_BLOCK(d, e, f, g, h, a, b, c, 61, 0xA4506CEB, indexW) \
ROUND_SECOND_BLOCK(c, d, e, f, g, h, a, b, 62, 0xBEF9A3F7, indexW) \
ROUND_SECOND_BLOCK(b, c, d, e, f, g, h, a, 63, 0xC67178F2, indexW)