nv/m11700_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _GOST2012_256_
   7
   8 #include "include/constants.h"
   9 #include "include/kernel_vendor.h"
  10
  11 #ifdef  VLIW1
  12 #define VECT_SIZE1
  13 #endif
  14
  15 #ifdef  VLIW2
  16 #define VECT_SIZE1
  17 #endif
  18
  19 #define DGST_R0 0
  20 #define DGST_R1 1
  21 #define DGST_R2 2
  22 #define DGST_R3 3
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27
  28 #ifdef  VECT_SIZE1
  29 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  30 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  31 #endif
  32
  33 #ifdef  VECT_SIZE2
  34 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  36 #endif
  37
  38 #ifdef  VECT_SIZE4
  39 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  40 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  41 #endif
  42
  43 #define INITVAL 0x0101010101010101
  44
  45 #define SBOG_LPSti64                         \
  46   s_sbob_sl64[0][(t[0] >> (i * 8)) & 0xff] ^ \
  47   s_sbob_sl64[1][(t[1] >> (i * 8)) & 0xff] ^ \
  48   s_sbob_sl64[2][(t[2] >> (i * 8)) & 0xff] ^ \
  49   s_sbob_sl64[3][(t[3] >> (i * 8)) & 0xff] ^ \
  50   s_sbob_sl64[4][(t[4] >> (i * 8)) & 0xff] ^ \
  51   s_sbob_sl64[5][(t[5] >> (i * 8)) & 0xff] ^ \
  52   s_sbob_sl64[6][(t[6] >> (i * 8)) & 0xff] ^ \
  53   s_sbob_sl64[7][(t[7] >> (i * 8)) & 0xff]
  54
  55 // constants
  56
  57 __device__ __constant__ u64 sbob_sl64[8][256] =
  58 {
  59   {
  60     0xd031c397ce553fe6,
  61     0x16ba5b01b006b525,
  62     0xa89bade6296e70c8,
  63     0x6a1f525d77d3435b,
  64     0x6e103570573dfa0b,
  65     0x660efb2a17fc95ab,
  66     0x76327a9e97634bf6,
  67     0x4bad9d6462458bf5,
  68     0xf1830caedbc3f748,
  69     0xc5c8f542669131ff,
  70     0x95044a1cdc48b0cb,
  71     0x892962df3cf8b866,
  72     0xb0b9e208e930c135,
  73     0xa14fb3f0611a767c,
  74     0x8d2605f21c160136,
  75     0xd6b71922fecc549e,
  76     0x37089438a5907d8b,
  77     0x0b5da38e5803d49c,
  78     0x5a5bcc9cea6f3cbc,
  79     0xedae246d3b73ffe5,
  80     0xd2b87e0fde22edce,
  81     0x5e54abb1ca8185ec,
  82     0x1de7f88fe80561b9,
  83     0xad5e1a870135a08c,
  84     0x2f2adbd665cecc76,
  85     0x5780b5a782f58358,
  86     0x3edc8a2eede47b3f,
  87     0xc9d95c3506bee70f,
  88     0x83be111d6c4e05ee,
  89     0xa603b90959367410,
  90     0x103c81b4809fde5d,
  91     0x2c69b6027d0c774a,
  92     0x399080d7d5c87953,
  93     0x09d41e16487406b4,
  94     0xcdd63b1826505e5f,
  95     0xf99dc2f49b0298e8,
  96     0x9cd0540a943cb67f,
  97     0xbca84b7f891f17c5,
  98     0x723d1db3b78df2a6,
  99     0x78aa6e71e73b4f2e,
 100     0x1433e699a071670d,
 101     0x84f21be454620782,
 102     0x98df3327b4d20f2f,
 103     0xf049dce2d3769e5c,
 104     0xdb6c60199656eb7a,
 105     0x648746b2078b4783,
 106     0x32cd23598dcbadcf,
 107     0x1ea4955bf0c7da85,
 108     0xe9a143401b9d46b5,
 109     0xfd92a5d9bbec21b8,
 110     0xc8138c790e0b8e1b,
 111     0x2ee00b9a6d7ba562,
 112     0xf85712b893b7f1fc,
 113     0xeb28fed80bea949d,
 114     0x564a65eb8a40ea4c,
 115     0x6c9988e8474a2823,
 116     0x4535898b121d8f2d,
 117     0xabd8c03231accbf4,
 118     0xba2e91cab9867cbd,
 119     0x7960be3def8e263a,
 120     0x0c11a977602fd6f0,
 121     0xcb50e1ad16c93527,
 122     0xeae22e94035ffd89,
 123     0x2866d12f5de2ce1a,
 124     0xff1b1841ab9bf390,
 125     0x9f9339de8cfe0d43,
 126     0x964727c8c48a0bf7,
 127     0x524502c6aaae531c,
 128     0x9b9c5ef3ac10b413,
 129     0x4fa2fa4942ab32a5,
 130     0x3f165a62e551122b,
 131     0xc74148da76e6e3d7,
 132     0x924840e5e464b2a7,
 133     0xd372ae43d69784da,
 134     0x233b72a105e11a86,
 135     0xa48a04914941a638,
 136     0xb4b68525c9de7865,
 137     0xddeabaaca6cf8002,
 138     0x0a9773c250b6bd88,
 139     0xc284ffbb5ebd3393,
 140     0x8ba0df472c8f6a4e,
 141     0x2aef6cb74d951c32,
 142     0x427983722a318d41,
 143     0x73f7cdffbf389bb2,
 144     0x074c0af9382c026c,
 145     0x8a6a0f0b243a035a,
 146     0x6fdae53c5f88931f,
 147     0xc68b98967e538ac3,
 148     0x44ff59c71aa8e639,
 149     0xe2fce0ce439e9229,
 150     0xa20cde2479d8cd40,
 151     0x19e89fa2c8ebd8e9,
 152     0xf446bbcff398270c,
 153     0x43b3533e2284e455,
 154     0xd82f0dcd8e945046,
 155     0x51066f12b26ce820,
 156     0xe73957af6bc5426d,
 157     0x081ece5a40c16fa0,
 158     0x3b193d4fc5bfab7b,
 159     0x7fe66488df174d42,
 160     0x0e9814ef705804d8,
 161     0x8137ac857c39d7c6,
 162     0xb1733244e185a821,
 163     0x695c3f896f11f867,
 164     0xf6cf0657e3eff524,
 165     0x1aabf276d02963d5,
 166     0x2da3664e75b91e5e,
 167     0x0289bd981077d228,
 168     0x90c1fd7df413608f,
 169     0x3c5537b6fd93a917,
 170     0xaa12107e3919a2e0,
 171     0x0686dab530996b78,
 172     0xdaa6b0559ee3826e,
 173     0xc34e2ff756085a87,
 174     0x6d5358a44fff4137,
 175     0xfc587595b35948ac,
 176     0x7ca5095cc7d5f67e,
 177     0xfb147f6c8b754ac0,
 178     0xbfeb26ab91ddacf9,
 179     0x6896efc567a49173,
 180     0xca9a31e11e7c5c33,
 181     0xbbe44186b13315a9,
 182     0x0ddb793b689abfe4,
 183     0x70b4a02ba7fa208e,
 184     0xe47a3a7b7307f951,
 185     0x8cecd5be14a36822,
 186     0xeeed49b923b144d9,
 187     0x17708b4db8b3dc31,
 188     0x6088219f2765fed3,
 189     0xb3fa8fdcf1f27a09,
 190     0x910b2d31fca6099b,
 191     0x0f52c4a378ed6dcc,
 192     0x50ccbf5ebad98134,
 193     0x6bd582117f662a4f,
 194     0x94ce9a50d4fdd9df,
 195     0x2b25bcfb45207526,
 196     0x67c42b661f49fcbf,
 197     0x492420fc723259dd,
 198     0x03436dd418c2bb3c,
 199     0x1f6e4517f872b391,
 200     0xa08563bc69af1f68,
 201     0xd43ea4baeebb86b6,
 202     0x01cad04c08b56914,
 203     0xac94cacb0980c998,
 204     0x54c3d8739a373864,
 205     0x26fec5c02dbacac2,
 206     0xdea9d778be0d3b3e,
 207     0x040f672d20eeb950,
 208     0xe5b0ea377bb29045,
 209     0xf30ab136cbb42560,
 210     0x62019c0737122cfb,
 211     0xe86b930c13282fa1,
 212     0xcc1ceb542ee5374b,
 213     0x538fd28aa21b3a08,
 214     0x1b61223ad89c0ac1,
 215     0x36c24474ad25149f,
 216     0x7a23d3e9f74c9d06,
 217     0xbe21f6e79968c5ed,
 218     0xcf5f868036278c77,
 219     0xf705d61beb5a9c30,
 220     0x4d2b47d152dce08d,
 221     0x5f9e7bfdc234ecf8,
 222     0x247778583dcd18ea,
 223     0x867ba67c4415d5aa,
 224     0x4ce1979d5a698999,
 225     0x0000000000000000,
 226     0xec64f42133c696f1,
 227     0xb57c5569c16b1171,
 228     0xc1c7926f467f88af,
 229     0x654d96fe0f3e2e97,
 230     0x15f936d5a8c40e19,
 231     0xb8a72c52a9f1ae95,
 232     0xa9517daa21db19dc,
 233     0x58d27104fa18ee94,
 234     0x5918a148f2ad8780,
 235     0x5cdd1629daf657c4,
 236     0x8274c15164fb6cfa,
 237     0xd1fb13dbc6e056f2,
 238     0x7d6fd910cf609f6a,
 239     0xb63f38bdd9a9aa4d,
 240     0x3d9fe7faf526c003,
 241     0x74bbc706871499de,
 242     0xdf630734b6b8522a,
 243     0x3ad3ed03cd0ac26f,
 244     0xfadeaf2083c023d4,
 245     0xc00d42234ecae1bb,
 246     0x8538cba85cd76e96,
 247     0xc402250e6e2458eb,
 248     0x47bc3413026a5d05,
 249     0xafd7a71f114272a4,
 250     0x978df784cc3f62e3,
 251     0xb96dfc1ea144c781,
 252     0x21b2cf391596c8ae,
 253     0x318e4e8d950916f3,
 254     0xce9556cc3e92e563,
 255     0x385a509bdd7d1047,
 256     0x358129a0b5e7afa3,
 257     0xe6f387e363702b79,
 258     0xe0755d5653e94001,
 259     0x7be903a5fff9f412,
 260     0x12b53c2c90e80c75,
 261     0x3307f315857ec4db,
 262     0x8fafb86a0c61d31e,
 263     0xd9e5dd8186213952,
 264     0x77f8aad29fd622e2,
 265     0x25bda814357871fe,
 266     0x7571174a8fa1f0ca,
 267     0x137fec60985d6561,
 268     0x30449ec19dbc7fe7,
 269     0xa540d4dd41f4cf2c,
 270     0xdc206ae0ae7ae916,
 271     0x5b911cd0e2da55a8,
 272     0xb2305f90f947131d,
 273     0x344bf9ecbd52c6b7,
 274     0x5d17c665d2433ed0,
 275     0x18224feec05eb1fd,
 276     0x9e59e992844b6457,
 277     0x9a568ebfa4a5dd07,
 278     0xa3c60e68716da454,
 279     0x7e2cb4c4d7a22456,
 280     0x87b176304ca0bcbe,
 281     0x413aeea632f3367d,
 282     0x9915e36bbc67663b,
 283     0x40f03eea3a465f69,
 284     0x1c2d28c3e0b008ad,
 285     0x4e682a054a1e5bb1,
 286     0x05c5b761285bd044,
 287     0xe1bf8d1a5b5c2915,
 288     0xf2c0617ac3014c74,
 289     0xb7f5e8f1d11cc359,
 290     0x63cb4c4b3fa745ef,
 291     0x9d1a84469c89df6b,
 292     0xe33630824b2bfb3d,
 293     0xd5f474f6e60eefa2,
 294     0xf58c6b83fb2d4e18,
 295     0x4676e45f0adf3411,
 296     0x20781f751d23a1ba,
 297     0xbd629b3381aa7ed1,
 298     0xae1d775319f71bb0,
 299     0xfed1c80da32e9a84,
 300     0x5509083f92825170,
 301     0x29ac01635557a70e,
 302     0xa7c9694551831d04,
 303     0x8e65682604d4ba0a,
 304     0x11f651f8882ab749,
 305     0xd77dc96ef6793d8a,
 306     0xef2799f52b042dcd,
 307     0x48eef0b07a8730c9,
 308     0x22f1a2ed0d547392,
 309     0x6142f1d32fd097c7,
 310     0x4a674d286af0e2e1,
 311     0x80fd7cc9748cbed2,
 312     0x717e7067af4f499a,
 313     0x938290a9ecd1dbb3,
 314     0x88e3b293344dd172,
 315     0x2734158c250fa3d6,
 316   },
 317   {
 318     0x7e37e62dfc7d40c3,
 319     0x776f25a4ee939e5b,
 320     0xe045c850dd8fb5ad,
 321     0x86ed5ba711ff1952,
 322     0xe91d0bd9cf616b35,
 323     0x37e0ab256e408ffb,
 324     0x9607f6c031025a7a,
 325     0x0b02f5e116d23c9d,
 326     0xf3d8486bfb50650c,
 327     0x621cff27c40875f5,
 328     0x7d40cb71fa5fd34a,
 329     0x6daa6616daa29062,
 330     0x9f5f354923ec84e2,
 331     0xec847c3dc507c3b3,
 332     0x025a3668043ce205,
 333     0xa8bf9e6c4dac0b19,
 334     0xfa808be2e9bebb94,
 335     0xb5b99c5277c74fa3,
 336     0x78d9bc95f0397bcc,
 337     0xe332e50cdbad2624,
 338     0xc74fce129332797e,
 339     0x1729eceb2ea709ab,
 340     0xc2d6b9f69954d1f8,
 341     0x5d898cbfbab8551a,
 342     0x859a76fb17dd8adb,
 343     0x1be85886362f7fb5,
 344     0xf6413f8ff136cd8a,
 345     0xd3110fa5bbb7e35c,
 346     0x0a2feed514cc4d11,
 347     0xe83010edcd7f1ab9,
 348     0xa1e75de55f42d581,
 349     0xeede4a55c13b21b6,
 350     0xf2f5535ff94e1480,
 351     0x0cc1b46d1888761e,
 352     0xbce15fdb6529913b,
 353     0x2d25e8975a7181c2,
 354     0x71817f1ce2d7a554,
 355     0x2e52c5cb5c53124b,
 356     0xf9f7a6beef9c281d,
 357     0x9e722e7d21f2f56e,
 358     0xce170d9b81dca7e6,
 359     0x0e9b82051cb4941b,
 360     0x1e712f623c49d733,
 361     0x21e45cfa42f9f7dc,
 362     0xcb8e7a7f8bba0f60,
 363     0x8e98831a010fb646,
 364     0x474ccf0d8e895b23,
 365     0xa99285584fb27a95,
 366     0x8cc2b57205335443,
 367     0x42d5b8e984eff3a5,
 368     0x012d1b34021e718c,
 369     0x57a6626aae74180b,
 370     0xff19fc06e3d81312,
 371     0x35ba9d4d6a7c6dfe,
 372     0xc9d44c178f86ed65,
 373     0x506523e6a02e5288,
 374     0x03772d5c06229389,
 375     0x8b01f4fe0b691ec0,
 376     0xf8dabd8aed825991,
 377     0x4c4e3aec985b67be,
 378     0xb10df0827fbf96a9,
 379     0x6a69279ad4f8dae1,
 380     0xe78689dcd3d5ff2e,
 381     0x812e1a2b1fa553d1,
 382     0xfbad90d6eba0ca18,
 383     0x1ac543b234310e39,
 384     0x1604f7df2cb97827,
 385     0xa6241c6951189f02,
 386     0x753513cceaaf7c5e,
 387     0x64f2a59fc84c4efa,
 388     0x247d2b1e489f5f5a,
 389     0xdb64d718ab474c48,
 390     0x79f4a7a1f2270a40,
 391     0x1573da832a9bebae,
 392     0x3497867968621c72,
 393     0x514838d2a2302304,
 394     0xf0af6537fd72f685,
 395     0x1d06023e3a6b44ba,
 396     0x678588c3ce6edd73,
 397     0x66a893f7cc70acff,
 398     0xd4d24e29b5eda9df,
 399     0x3856321470ea6a6c,
 400     0x07c3418c0e5a4a83,
 401     0x2bcbb22f5635bacd,
 402     0x04b46cd00878d90a,
 403     0x06ee5ab80c443b0f,
 404     0x3b211f4876c8f9e5,
 405     0x0958c38912eede98,
 406     0xd14b39cdbf8b0159,
 407     0x397b292072f41be0,
 408     0x87c0409313e168de,
 409     0xad26e98847caa39f,
 410     0x4e140c849c6785bb,
 411     0xd5ff551db7f3d853,
 412     0xa0ca46d15d5ca40d,
 413     0xcd6020c787fe346f,
 414     0x84b76dcf15c3fb57,
 415     0xdefda0fca121e4ce,
 416     0x4b8d7b6096012d3d,
 417     0x9ac642ad298a2c64,
 418     0x0875d8bd10f0af14,
 419     0xb357c6ea7b8374ac,
 420     0x4d6321d89a451632,
 421     0xeda96709c719b23f,
 422     0xf76c24bbf328bc06,
 423     0xc662d526912c08f2,
 424     0x3ce25ec47892b366,
 425     0xb978283f6f4f39bd,
 426     0xc08c8f9e9d6833fd,
 427     0x4f3917b09e79f437,
 428     0x593de06fb2c08c10,
 429     0xd6887841b1d14bda,
 430     0x19b26eee32139db0,
 431     0xb494876675d93e2f,
 432     0x825937771987c058,
 433     0x90e9ac783d466175,
 434     0xf1827e03ff6c8709,
 435     0x945dc0a8353eb87f,
 436     0x4516f9658ab5b926,
 437     0x3f9573987eb020ef,
 438     0xb855330b6d514831,
 439     0x2ae6a91b542bcb41,
 440     0x6331e413c6160479,
 441     0x408f8e8180d311a0,
 442     0xeff35161c325503a,
 443     0xd06622f9bd9570d5,
 444     0x8876d9a20d4b8d49,
 445     0xa5533135573a0c8b,
 446     0xe168d364df91c421,
 447     0xf41b09e7f50a2f8f,
 448     0x12b09b0f24c1a12d,
 449     0xda49cc2ca9593dc4,
 450     0x1f5c34563e57a6bf,
 451     0x54d14f36a8568b82,
 452     0xaf7cdfe043f6419a,
 453     0xea6a2685c943f8bc,
 454     0xe5dcbfb4d7e91d2b,
 455     0xb27addde799d0520,
 456     0x6b443caed6e6ab6d,
 457     0x7bae91c9f61be845,
 458     0x3eb868ac7cae5163,
 459     0x11c7b65322e332a4,
 460     0xd23c1491b9a992d0,
 461     0x8fb5982e0311c7ca,
 462     0x70ac6428e0c9d4d8,
 463     0x895bc2960f55fcc5,
 464     0x76423e90ec8defd7,
 465     0x6ff0507ede9e7267,
 466     0x3dcf45f07a8cc2ea,
 467     0x4aa06054941f5cb1,
 468     0x5810fb5bb0defd9c,
 469     0x5efea1e3bc9ac693,
 470     0x6edd4b4adc8003eb,
 471     0x741808f8e8b10dd2,
 472     0x145ec1b728859a22,
 473     0x28bc9f7350172944,
 474     0x270a06424ebdccd3,
 475     0x972aedf4331c2bf6,
 476     0x059977e40a66a886,
 477     0x2550302a4a812ed6,
 478     0xdd8a8da0a7037747,
 479     0xc515f87a970e9b7b,
 480     0x3023eaa9601ac578,
 481     0xb7e3aa3a73fbada6,
 482     0x0fb699311eaae597,
 483     0x0000000000000000,
 484     0x310ef19d6204b4f4,
 485     0x229371a644db6455,
 486     0x0decaf591a960792,
 487     0x5ca4978bb8a62496,
 488     0x1c2b190a38753536,
 489     0x41a295b582cd602c,
 490     0x3279dcc16426277d,
 491     0xc1a194aa9f764271,
 492     0x139d803b26dfd0a1,
 493     0xae51c4d441e83016,
 494     0xd813fa44ad65dfc1,
 495     0xac0bf2bc45d4d213,
 496     0x23be6a9246c515d9,
 497     0x49d74d08923dcf38,
 498     0x9d05032127d066e7,
 499     0x2f7fdeff5e4d63c7,
 500     0xa47e2a0155247d07,
 501     0x99b16ff12fa8bfed,
 502     0x4661d4398c972aaf,
 503     0xdfd0bbc8a33f9542,
 504     0xdca79694a51d06cb,
 505     0xb020ebb67da1e725,
 506     0xba0f0563696daa34,
 507     0xe4f1a480d5f76ca7,
 508     0xc438e34e9510eaf7,
 509     0x939e81243b64f2fc,
 510     0x8defae46072d25cf,
 511     0x2c08f3a3586ff04e,
 512     0xd7a56375b3cf3a56,
 513     0x20c947ce40e78650,
 514     0x43f8a3dd86f18229,
 515     0x568b795eac6a6987,
 516     0x8003011f1dbb225d,
 517     0xf53612d3f7145e03,
 518     0x189f75da300dec3c,
 519     0x9570db9c3720c9f3,
 520     0xbb221e576b73dbb8,
 521     0x72f65240e4f536dd,
 522     0x443be25188abc8aa,
 523     0xe21ffe38d9b357a8,
 524     0xfd43ca6ee7e4f117,
 525     0xcaa3614b89a47eec,
 526     0xfe34e732e1c6629e,
 527     0x83742c431b99b1d4,
 528     0xcf3a16af83c2d66a,
 529     0xaae5a8044990e91c,
 530     0x26271d764ca3bd5f,
 531     0x91c4b74c3f5810f9,
 532     0x7c6dd045f841a2c6,
 533     0x7f1afd19fe63314f,
 534     0xc8f957238d989ce9,
 535     0xa709075d5306ee8e,
 536     0x55fc5402aa48fa0e,
 537     0x48fa563c9023beb4,
 538     0x65dfbeabca523f76,
 539     0x6c877d22d8bce1ee,
 540     0xcc4d3bf385e045e3,
 541     0xbebb69b36115733e,
 542     0x10eaad6720fd4328,
 543     0xb6ceb10e71e5dc2a,
 544     0xbdcc44ef6737e0b7,
 545     0x523f158ea412b08d,
 546     0x989c74c52db6ce61,
 547     0x9beb59992b945de8,
 548     0x8a2cefca09776f4c,
 549     0xa3bd6b8d5b7e3784,
 550     0xeb473db1cb5d8930,
 551     0xc3fba2c29b4aa074,
 552     0x9c28181525ce176b,
 553     0x683311f2d0c438e4,
 554     0x5fd3bad7be84b71f,
 555     0xfc6ed15ae5fa809b,
 556     0x36cdb0116c5efe77,
 557     0x29918447520958c8,
 558     0xa29070b959604608,
 559     0x53120ebaa60cc101,
 560     0x3a0c047c74d68869,
 561     0x691e0ac6d2da4968,
 562     0x73db4974e6eb4751,
 563     0x7a838afdf40599c9,
 564     0x5a4acd33b4e21f99,
 565     0x6046c94fc03497f0,
 566     0xe6ab92e8d1cb8ea2,
 567     0x3354c7f5663856f1,
 568     0xd93ee170af7bae4d,
 569     0x616bd27bc22ae67c,
 570     0x92b39a10397a8370,
 571     0xabc8b3304b8e9890,
 572     0xbf967287630b02b2,
 573     0x5b67d607b6fc6e15,
 574   },
 575   {
 576     0x8ab0a96846e06a6d,
 577     0x43c7e80b4bf0b33a,
 578     0x08c9b3546b161ee5,
 579     0x39f1c235eba990be,
 580     0xc1bef2376606c7b2,
 581     0x2c209233614569aa,
 582     0xeb01523b6fc3289a,
 583     0x946953ab935acedd,
 584     0x272838f63e13340e,
 585     0x8b0455eca12ba052,
 586     0x77a1b2c4978ff8a2,
 587     0xa55122ca13e54086,
 588     0x2276135862d3f1cd,
 589     0xdb8ddfde08b76cfe,
 590     0x5d1e12c89e4a178a,
 591     0x0e56816b03969867,
 592     0xee5f79953303ed59,
 593     0xafed748bab78d71d,
 594     0x6d929f2df93e53ee,
 595     0xf5d8a8f8ba798c2a,
 596     0xf619b1698e39cf6b,
 597     0x95ddaf2f749104e2,
 598     0xec2a9c80e0886427,
 599     0xce5c8fd8825b95ea,
 600     0xc4e0d9993ac60271,
 601     0x4699c3a5173076f9,
 602     0x3d1b151f50a29f42,
 603     0x9ed505ea2bc75946,
 604     0x34665acfdc7f4b98,
 605     0x61b1fb53292342f7,
 606     0xc721c0080e864130,
 607     0x8693cd1696fd7b74,
 608     0x872731927136b14b,
 609     0xd3446c8a63a1721b,
 610     0x669a35e8a6680e4a,
 611     0xcab658f239509a16,
 612     0xa4e5de4ef42e8ab9,
 613     0x37a7435ee83f08d9,
 614     0x134e6239e26c7f96,
 615     0x82791a3c2df67488,
 616     0x3f6ef00a8329163c,
 617     0x8e5a7e42fdeb6591,
 618     0x5caaee4c7981ddb5,
 619     0x19f234785af1e80d,
 620     0x255ddde3ed98bd70,
 621     0x50898a32a99cccac,
 622     0x28ca4519da4e6656,
 623     0xae59880f4cb31d22,
 624     0x0d9798fa37d6db26,
 625     0x32f968f0b4ffcd1a,
 626     0xa00f09644f258545,
 627     0xfa3ad5175e24de72,
 628     0xf46c547c5db24615,
 629     0x713e80fbff0f7e20,
 630     0x7843cf2b73d2aafa,
 631     0xbd17ea36aedf62b4,
 632     0xfd111bacd16f92cf,
 633     0x4abaa7dbc72d67e0,
 634     0xb3416b5dad49fad3,
 635     0xbca316b24914a88b,
 636     0x15d150068aecf914,
 637     0xe27c1debe31efc40,
 638     0x4fe48c759beda223,
 639     0x7edcfd141b522c78,
 640     0x4e5070f17c26681c,
 641     0xe696cac15815f3bc,
 642     0x35d2a64b3bb481a7,
 643     0x800cff29fe7dfdf6,
 644     0x1ed9fac3d5baa4b0,
 645     0x6c2663a91ef599d1,
 646     0x03c1199134404341,
 647     0xf7ad4ded69f20554,
 648     0xcd9d9649b61bd6ab,
 649     0xc8c3bde7eadb1368,
 650     0xd131899fb02afb65,
 651     0x1d18e352e1fae7f1,
 652     0xda39235aef7ca6c1,
 653     0xa1bbf5e0a8ee4f7a,
 654     0x91377805cf9a0b1e,
 655     0x3138716180bf8e5b,
 656     0xd9f83acbdb3ce580,
 657     0x0275e515d38b897e,
 658     0x472d3f21f0fbbcc6,
 659     0x2d946eb7868ea395,
 660     0xba3c248d21942e09,
 661     0xe7223645bfde3983,
 662     0xff64feb902e41bb1,
 663     0xc97741630d10d957,
 664     0xc3cb1722b58d4ecc,
 665     0xa27aec719cae0c3b,
 666     0x99fecb51a48c15fb,
 667     0x1465ac826d27332b,
 668     0xe1bd047ad75ebf01,
 669     0x79f733af941960c5,
 670     0x672ec96c41a3c475,
 671     0xc27feba6524684f3,
 672     0x64efd0fd75e38734,
 673     0xed9e60040743ae18,
 674     0xfb8e2993b9ef144d,
 675     0x38453eb10c625a81,
 676     0x6978480742355c12,
 677     0x48cf42ce14a6ee9e,
 678     0x1cac1fd606312dce,
 679     0x7b82d6ba4792e9bb,
 680     0x9d141c7b1f871a07,
 681     0x5616b80dc11c4a2e,
 682     0xb849c198f21fa777,
 683     0x7ca91801c8d9a506,
 684     0xb1348e487ec273ad,
 685     0x41b20d1e987b3a44,
 686     0x7460ab55a3cfbbe3,
 687     0x84e628034576f20a,
 688     0x1b87d16d897a6173,
 689     0x0fe27defe45d5258,
 690     0x83cde6b8ca3dbeb7,
 691     0x0c23647ed01d1119,
 692     0x7a362a3ea0592384,
 693     0xb61f40f3f1893f10,
 694     0x75d457d1440471dc,
 695     0x4558da34237035b8,
 696     0xdca6116587fc2043,
 697     0x8d9b67d3c9ab26d0,
 698     0x2b0b5c88ee0e2517,
 699     0x6fe77a382ab5da90,
 700     0x269cc472d9d8fe31,
 701     0x63c41e46faa8cb89,
 702     0xb7abbc771642f52f,
 703     0x7d1de4852f126f39,
 704     0xa8c6ba3024339ba0,
 705     0x600507d7cee888c8,
 706     0x8fee82c61a20afae,
 707     0x57a2448926d78011,
 708     0xfca5e72836a458f0,
 709     0x072bcebb8f4b4cbd,
 710     0x497bbe4af36d24a1,
 711     0x3cafe99bb769557d,
 712     0x12fa9ebd05a7b5a9,
 713     0xe8c04baa5b836bdb,
 714     0x4273148fac3b7905,
 715     0x908384812851c121,
 716     0xe557d3506c55b0fd,
 717     0x72ff996acb4f3d61,
 718     0x3eda0c8e64e2dc03,
 719     0xf0868356e6b949e9,
 720     0x04ead72abb0b0ffc,
 721     0x17a4b5135967706a,
 722     0xe3c8e16f04d5367f,
 723     0xf84f30028daf570c,
 724     0x1846c8fcbd3a2232,
 725     0x5b8120f7f6ca9108,
 726     0xd46fa231ecea3ea6,
 727     0x334d947453340725,
 728     0x58403966c28ad249,
 729     0xbed6f3a79a9f21f5,
 730     0x68ccb483a5fe962d,
 731     0xd085751b57e1315a,
 732     0xfed0023de52fd18e,
 733     0x4b0e5b5f20e6addf,
 734     0x1a332de96eb1ab4c,
 735     0xa3ce10f57b65c604,
 736     0x108f7ba8d62c3cd7,
 737     0xab07a3a11073d8e1,
 738     0x6b0dad1291bed56c,
 739     0xf2f366433532c097,
 740     0x2e557726b2cee0d4,
 741     0x0000000000000000,
 742     0xcb02a476de9b5029,
 743     0xe4e32fd48b9e7ac2,
 744     0x734b65ee2c84f75e,
 745     0x6e5386bccd7e10af,
 746     0x01b4fc84e7cbca3f,
 747     0xcfe8735c65905fd5,
 748     0x3613bfda0ff4c2e6,
 749     0x113b872c31e7f6e8,
 750     0x2fe18ba255052aeb,
 751     0xe974b72ebc48a1e4,
 752     0x0abc5641b89d979b,
 753     0xb46aa5e62202b66e,
 754     0x44ec26b0c4bbff87,
 755     0xa6903b5b27a503c7,
 756     0x7f680190fc99e647,
 757     0x97a84a3aa71a8d9c,
 758     0xdd12ede16037ea7c,
 759     0xc554251ddd0dc84e,
 760     0x88c54c7d956be313,
 761     0x4d91696048662b5d,
 762     0xb08072cc9909b992,
 763     0xb5de5962c5c97c51,
 764     0x81b803ad19b637c9,
 765     0xb2f597d94a8230ec,
 766     0x0b08aac55f565da4,
 767     0xf1327fd2017283d6,
 768     0xad98919e78f35e63,
 769     0x6ab9519676751f53,
 770     0x24e921670a53774f,
 771     0xb9fd3d1c15d46d48,
 772     0x92f66194fbda485f,
 773     0x5a35dc7311015b37,
 774     0xded3f4705477a93d,
 775     0xc00a0eb381cd0d8d,
 776     0xbb88d809c65fe436,
 777     0x16104997beacba55,
 778     0x21b70ac95693b28c,
 779     0x59f4c5e225411876,
 780     0xd5db5eb50b21f499,
 781     0x55d7a19cf55c096f,
 782     0xa97246b4c3f8519f,
 783     0x8552d487a2bd3835,
 784     0x54635d181297c350,
 785     0x23c2efdc85183bf2,
 786     0x9f61f96ecc0c9379,
 787     0x534893a39ddc8fed,
 788     0x5edf0b59aa0a54cb,
 789     0xac2c6d1a9f38945c,
 790     0xd7aebba0d8aa7de7,
 791     0x2abfa00c09c5ef28,
 792     0xd84cc64f3cf72fbf,
 793     0x2003f64db15878b3,
 794     0xa724c7dfc06ec9f8,
 795     0x069f323f68808682,
 796     0xcc296acd51d01c94,
 797     0x055e2bae5cc0c5c3,
 798     0x6270e2c21d6301b6,
 799     0x3b842720382219c0,
 800     0xd2f0900e846ab824,
 801     0x52fc6f277a1745d2,
 802     0xc6953c8ce94d8b0f,
 803     0xe009f8fe3095753e,
 804     0x655b2c7992284d0b,
 805     0x984a37d54347dfc4,
 806     0xeab5aebf8808e2a5,
 807     0x9a3fd2c090cc56ba,
 808     0x9ca0e0fff84cd038,
 809     0x4c2595e4afade162,
 810     0xdf6708f4b3bc6302,
 811     0xbf620f237d54ebca,
 812     0x93429d101c118260,
 813     0x097d4fd08cddd4da,
 814     0x8c2f9b572e60ecef,
 815     0x708a7c7f18c4b41f,
 816     0x3a30dba4dfe9d3ff,
 817     0x4006f19a7fb0f07b,
 818     0x5f6bf7dd4dc19ef4,
 819     0x1f6d064732716e8f,
 820     0xf9fbcc866a649d33,
 821     0x308c8de567744464,
 822     0x8971b0f972a0292c,
 823     0xd61a47243f61b7d8,
 824     0xefeb8511d4c82766,
 825     0x961cb6be40d147a3,
 826     0xaab35f25f7b812de,
 827     0x76154e407044329d,
 828     0x513d76b64e570693,
 829     0xf3479ac7d2f90aa8,
 830     0x9b8b2e4477079c85,
 831     0x297eb99d3d85ac69,
 832   },
 833   {
 834     0x3ef29d249b2c0a19,
 835     0xe9e16322b6f8622f,
 836     0x5536994047757f7a,
 837     0x9f4d56d5a47b0b33,
 838     0x822567466aa1174c,
 839     0xb8f5057deb082fb2,
 840     0xcc48c10bf4475f53,
 841     0x373088d4275dec3a,
 842     0x968f4325180aed10,
 843     0x173d232cf7016151,
 844     0xae4ed09f946fcc13,
 845     0xfd4b4741c4539873,
 846     0x1b5b3f0dd9933765,
 847     0x2ffcb0967b644052,
 848     0xe02376d20a89840c,
 849     0xa3ae3a70329b18d7,
 850     0x419cbd2335de8526,
 851     0xfafebf115b7c3199,
 852     0x0397074f85aa9b0d,
 853     0xc58ad4fb4836b970,
 854     0xbec60be3fc4104a8,
 855     0x1eff36dc4b708772,
 856     0x131fdc33ed8453b6,
 857     0x0844e33e341764d3,
 858     0x0ff11b6eab38cd39,
 859     0x64351f0a7761b85a,
 860     0x3b5694f509cfba0e,
 861     0x30857084b87245d0,
 862     0x47afb3bd2297ae3c,
 863     0xf2ba5c2f6f6b554a,
 864     0x74bdc4761f4f70e1,
 865     0xcfdfc64471edc45e,
 866     0xe610784c1dc0af16,
 867     0x7aca29d63c113f28,
 868     0x2ded411776a859af,
 869     0xac5f211e99a3d5ee,
 870     0xd484f949a87ef33b,
 871     0x3ce36ca596e013e4,
 872     0xd120f0983a9d432c,
 873     0x6bc40464dc597563,
 874     0x69d5f5e5d1956c9e,
 875     0x9ae95f043698bb24,
 876     0xc9ecc8da66a4ef44,
 877     0xd69508c8a5b2eac6,
 878     0xc40c2235c0503b80,
 879     0x38c193ba8c652103,
 880     0x1ceec75d46bc9e8f,
 881     0xd331011937515ad1,
 882     0xd8e2e56886eca50f,
 883     0xb137108d5779c991,
 884     0x709f3b6905ca4206,
 885     0x4feb50831680caef,
 886     0xec456af3241bd238,
 887     0x58d673afe181abbe,
 888     0x242f54e7cad9bf8c,
 889     0x0211f1810dcc19fd,
 890     0x90bc4dbb0f43c60a,
 891     0x9518446a9da0761d,
 892     0xa1bfcbf13f57012a,
 893     0x2bde4f8961e172b5,
 894     0x27b853a84f732481,
 895     0xb0b1e643df1f4b61,
 896     0x18cc38425c39ac68,
 897     0xd2b7f7d7bf37d821,
 898     0x3103864a3014c720,
 899     0x14aa246372abfa5c,
 900     0x6e600db54ebac574,
 901     0x394765740403a3f3,
 902     0x09c215f0bc71e623,
 903     0x2a58b947e987f045,
 904     0x7b4cdf18b477bdd8,
 905     0x9709b5eb906c6fe0,
 906     0x73083c268060d90b,
 907     0xfedc400e41f9037e,
 908     0x284948c6e44be9b8,
 909     0x728ecae808065bfb,
 910     0x06330e9e17492b1a,
 911     0x5950856169e7294e,
 912     0xbae4f4fce6c4364f,
 913     0xca7bcf95e30e7449,
 914     0x7d7fd186a33e96c2,
 915     0x52836110d85ad690,
 916     0x4dfaa1021b4cd312,
 917     0x913abb75872544fa,
 918     0xdd46ecb9140f1518,
 919     0x3d659a6b1e869114,
 920     0xc23f2cabd719109a,
 921     0xd713fe062dd46836,
 922     0xd0a60656b2fbc1dc,
 923     0x221c5a79dd909496,
 924     0xefd26dbca1b14935,
 925     0x0e77eda0235e4fc9,
 926     0xcbfd395b6b68f6b9,
 927     0x0de0eaefa6f4d4c4,
 928     0x0422ff1f1a8532e7,
 929     0xf969b85eded6aa94,
 930     0x7f6e2007aef28f3f,
 931     0x3ad0623b81a938fe,
 932     0x6624ee8b7aada1a7,
 933     0xb682e8ddc856607b,
 934     0xa78cc56f281e2a30,
 935     0xc79b257a45faa08d,
 936     0x5b4174e0642b30b3,
 937     0x5f638bff7eae0254,
 938     0x4bc9af9c0c05f808,
 939     0xce59308af98b46ae,
 940     0x8fc58da9cc55c388,
 941     0x803496c7676d0eb1,
 942     0xf33caae1e70dd7ba,
 943     0xbb6202326ea2b4bf,
 944     0xd5020f87201871cb,
 945     0x9d5ca754a9b712ce,
 946     0x841669d87de83c56,
 947     0x8a6184785eb6739f,
 948     0x420bba6cb0741e2b,
 949     0xf12d5b60eac1ce47,
 950     0x76ac35f71283691c,
 951     0x2c6bb7d9fecedb5f,
 952     0xfccdb18f4c351a83,
 953     0x1f79c012c3160582,
 954     0xf0abadae62a74cb7,
 955     0xe1a5801c82ef06fc,
 956     0x67a21845f2cb2357,
 957     0x5114665f5df04d9d,
 958     0xbf40fd2d74278658,
 959     0xa0393d3fb73183da,
 960     0x05a409d192e3b017,
 961     0xa9fb28cf0b4065f9,
 962     0x25a9a22942bf3d7c,
 963     0xdb75e22703463e02,
 964     0xb326e10c5ab5d06c,
 965     0xe7968e8295a62de6,
 966     0xb973f3b3636ead42,
 967     0xdf571d3819c30ce5,
 968     0xee549b7229d7cbc5,
 969     0x12992afd65e2d146,
 970     0xf8ef4e9056b02864,
 971     0xb7041e134030e28b,
 972     0xc02edd2adad50967,
 973     0x932b4af48ae95d07,
 974     0x6fe6fb7bc6dc4784,
 975     0x239aacb755f61666,
 976     0x401a4bedbdb807d6,
 977     0x485ea8d389af6305,
 978     0xa41bc220adb4b13d,
 979     0x753b32b89729f211,
 980     0x997e584bb3322029,
 981     0x1d683193ceda1c7f,
 982     0xff5ab6c0c99f818e,
 983     0x16bbd5e27f67e3a1,
 984     0xa59d34ee25d233cd,
 985     0x98f8ae853b54a2d9,
 986     0x6df70afacb105e79,
 987     0x795d2e99b9bba425,
 988     0x8e437b6744334178,
 989     0x0186f6ce886682f0,
 990     0xebf092a3bb347bd2,
 991     0xbcd7fa62f18d1d55,
 992     0xadd9d7d011c5571e,
 993     0x0bd3e471b1bdffde,
 994     0xaa6c2f808eeafef4,
 995     0x5ee57d31f6c880a4,
 996     0xf50fa47ff044fca0,
 997     0x1addc9c351f5b595,
 998     0xea76646d3352f922,
 999     0x0000000000000000,
1000     0x85909f16f58ebea6,
1001     0x46294573aaf12ccc,
1002     0x0a5512bf39db7d2e,
1003     0x78dbd85731dd26d5,
1004     0x29cfbe086c2d6b48,
1005     0x218b5d36583a0f9b,
1006     0x152cd2adfacd78ac,
1007     0x83a39188e2c795bc,
1008     0xc3b9da655f7f926a,
1009     0x9ecba01b2c1d89c3,
1010     0x07b5f8509f2fa9ea,
1011     0x7ee8d6c926940dcf,
1012     0x36b67e1aaf3b6eca,
1013     0x86079859702425ab,
1014     0xfb7849dfd31ab369,
1015     0x4c7c57cc932a51e2,
1016     0xd96413a60e8a27ff,
1017     0x263ea566c715a671,
1018     0x6c71fc344376dc89,
1019     0x4a4f595284637af8,
1020     0xdaf314e98b20bcf2,
1021     0x572768c14ab96687,
1022     0x1088db7c682ec8bb,
1023     0x887075f9537a6a62,
1024     0x2e7a4658f302c2a2,
1025     0x619116dbe582084d,
1026     0xa87dde018326e709,
1027     0xdcc01a779c6997e8,
1028     0xedc39c3dac7d50c8,
1029     0xa60a33a1a078a8c0,
1030     0xc1a82be452b38b97,
1031     0x3f746bea134a88e9,
1032     0xa228ccbebafd9a27,
1033     0xabead94e068c7c04,
1034     0xf48952b178227e50,
1035     0x5cf48cb0fb049959,
1036     0x6017e0156de48abd,
1037     0x4438b4f2a73d3531,
1038     0x8c528ae649ff5885,
1039     0xb515ef924dfcfb76,
1040     0x0c661c212e925634,
1041     0xb493195cc59a7986,
1042     0x9cda519a21d1903e,
1043     0x32948105b5be5c2d,
1044     0x194ace8cd45f2e98,
1045     0x438d4ca238129cdb,
1046     0x9b6fa9cabefe39d4,
1047     0x81b26009ef0b8c41,
1048     0xded1ebf691a58e15,
1049     0x4e6da64d9ee6481f,
1050     0x54b06f8ecf13fd8a,
1051     0x49d85e1d01c9e1f5,
1052     0xafc826511c094ee3,
1053     0xf698a33075ee67ad,
1054     0x5ac7822eec4db243,
1055     0x8dd47c28c199da75,
1056     0x89f68337db1ce892,
1057     0xcdce37c57c21dda3,
1058     0x530597de503c5460,
1059     0x6a42f2aa543ff793,
1060     0x5d727a7e73621ba9,
1061     0xe232875307459df1,
1062     0x56a19e0fc2dfe477,
1063     0xc61dd3b4cd9c227d,
1064     0xe5877f03986a341b,
1065     0x949eb2a415c6f4ed,
1066     0x6206119460289340,
1067     0x6380e75ae84e11b0,
1068     0x8be772b6d6d0f16f,
1069     0x50929091d596cf6d,
1070     0xe86795ec3e9ee0df,
1071     0x7cf927482b581432,
1072     0xc86a3e14eec26db4,
1073     0x7119cda78dacc0f6,
1074     0xe40189cd100cb6eb,
1075     0x92adbc3a028fdff7,
1076     0xb2a017c2d2d3529c,
1077     0x200dabf8d05c8d6b,
1078     0x34a78f9ba2f77737,
1079     0xe3b4719d8f231f01,
1080     0x45be423c2f5bb7c1,
1081     0xf71e55fefd88e55d,
1082     0x6853032b59f3ee6e,
1083     0x65b3e9c4ff073aaa,
1084     0x772ac3399ae5ebec,
1085     0x87816e97f842a75b,
1086     0x110e2db2e0484a4b,
1087     0x331277cb3dd8dedd,
1088     0xbd510cac79eb9fa5,
1089     0x352179552a91f5c7,
1090   },
1091   {
1092     0x05ba7bc82c9b3220,
1093     0x31a54665f8b65e4f,
1094     0xb1b651f77547f4d4,
1095     0x8bfa0d857ba46682,
1096     0x85a96c5aa16a98bb,
1097     0x990faef908eb79c9,
1098     0xa15e37a247f4a62d,
1099     0x76857dcd5d27741e,
1100     0xf8c50b800a1820bc,
1101     0xbe65dcb201f7a2b4,
1102     0x666d1b986f9426e7,
1103     0x4cc921bf53c4e648,
1104     0x95410a0f93d9ca42,
1105     0x20cdccaa647ba4ef,
1106     0x429a4060890a1871,
1107     0x0c4ea4f69b32b38b,
1108     0xccda362dde354cd3,
1109     0x96dc23bc7c5b2fa9,
1110     0xc309bb68aa851ab3,
1111     0xd26131a73648e013,
1112     0x021dc52941fc4db2,
1113     0xcd5adab7704be48a,
1114     0xa77965d984ed71e6,
1115     0x32386fd61734bba4,
1116     0xe82d6dd538ab7245,
1117     0x5c2147ea6177b4b1,
1118     0x5da1ab70cf091ce8,
1119     0xac907fce72b8bdff,
1120     0x57c85dfd972278a8,
1121     0xa4e44c6a6b6f940d,
1122     0x3851995b4f1fdfe4,
1123     0x62578ccaed71bc9e,
1124     0xd9882bb0c01d2c0a,
1125     0x917b9d5d113c503b,
1126     0xa2c31e11a87643c6,
1127     0xe463c923a399c1ce,
1128     0xf71686c57ea876dc,
1129     0x87b4a973e096d509,
1130     0xaf0d567d9d3a5814,
1131     0xb40c2a3f59dcc6f4,
1132     0x3602f88495d121dd,
1133     0xd3e1dd3d9836484a,
1134     0xf945e71aa46688e5,
1135     0x7518547eb2a591f5,
1136     0x9366587450c01d89,
1137     0x9ea81018658c065b,
1138     0x4f54080cbc4603a3,
1139     0x2d0384c65137bf3d,
1140     0xdc325078ec861e2a,
1141     0xea30a8fc79573ff7,
1142     0x214d2030ca050cb6,
1143     0x65f0322b8016c30c,
1144     0x69be96dd1b247087,
1145     0xdb95ee9981e161b8,
1146     0xd1fc1814d9ca05f8,
1147     0x820ed2bbcc0de729,
1148     0x63d76050430f14c7,
1149     0x3bccb0e8a09d3a0f,
1150     0x8e40764d573f54a2,
1151     0x39d175c1e16177bd,
1152     0x12f5a37c734f1f4b,
1153     0xab37c12f1fdfc26d,
1154     0x5648b167395cd0f1,
1155     0x6c04ed1537bf42a7,
1156     0xed97161d14304065,
1157     0x7d6c67daab72b807,
1158     0xec17fa87ba4ee83c,
1159     0xdfaf79cb0304fbc1,
1160     0x733f060571bc463e,
1161     0x78d61c1287e98a27,
1162     0xd07cf48e77b4ada1,
1163     0xb9c262536c90dd26,
1164     0xe2449b5860801605,
1165     0x8fc09ad7f941fcfb,
1166     0xfad8cea94be46d0e,
1167     0xa343f28b0608eb9f,
1168     0x9b126bd04917347b,
1169     0x9a92874ae7699c22,
1170     0x1b017c42c4e69ee0,
1171     0x3a4c5c720ee39256,
1172     0x4b6e9f5e3ea399da,
1173     0x6ba353f45ad83d35,
1174     0xe7fee0904c1b2425,
1175     0x22d009832587e95d,
1176     0x842980c00f1430e2,
1177     0xc6b3c0a0861e2893,
1178     0x087433a419d729f2,
1179     0x341f3dadd42d6c6f,
1180     0xee0a3faefbb2a58e,
1181     0x4aee73c490dd3183,
1182     0xaab72db5b1a16a34,
1183     0xa92a04065e238fdf,
1184     0x7b4b35a1686b6fcc,
1185     0x6a23bf6ef4a6956c,
1186     0x191cb96b851ad352,
1187     0x55d598d4d6de351a,
1188     0xc9604de5f2ae7ef3,
1189     0x1ca6c2a3a981e172,
1190     0xde2f9551ad7a5398,
1191     0x3025aaff56c8f616,
1192     0x15521d9d1e2860d9,
1193     0x506fe31cfa45073a,
1194     0x189c55f12b647b0b,
1195     0x0180ec9aae7ea859,
1196     0x7cec8b40050c105e,
1197     0x2350e5198bf94104,
1198     0xef8ad33455cc0dd7,
1199     0x07a7bee16d677f92,
1200     0xe5e325b90de76997,
1201     0x5a061591a26e637a,
1202     0xb611ef1618208b46,
1203     0x09f4df3eb7a981ab,
1204     0x1ebb078ae87dacc0,
1205     0xb791038cb65e231f,
1206     0x0fd38d4574b05660,
1207     0x67edf702c1ea8ebe,
1208     0xba5f4be0831238cd,
1209     0xe3c477c2cefebe5c,
1210     0x0dce486c354c1bd2,
1211     0x8c5db36416c31910,
1212     0x26ea9ed1a7627324,
1213     0x039d29b3ef82e5eb,
1214     0x9f28fc82cbf2ae02,
1215     0xa8aae89cf05d2786,
1216     0x431aacfa2774b028,
1217     0xcf471f9e31b7a938,
1218     0x581bd0b8e3922ec8,
1219     0xbc78199b400bef06,
1220     0x90fb71c7bf42f862,
1221     0x1f3beb1046030499,
1222     0x683e7a47b55ad8de,
1223     0x988f4263a695d190,
1224     0xd808c72a6e638453,
1225     0x0627527bc319d7cb,
1226     0xebb04466d72997ae,
1227     0xe67e0c0ae2658c7c,
1228     0x14d2f107b056c880,
1229     0x7122c32c30400b8c,
1230     0x8a7ae11fd5dacedb,
1231     0xa0dedb38e98a0e74,
1232     0xad109354dcc615a6,
1233     0x0be91a17f655cc19,
1234     0x8ddd5ffeb8bdb149,
1235     0xbfe53028af890aed,
1236     0xd65ba6f5b4ad7a6a,
1237     0x7956f0882997227e,
1238     0x10e8665532b352f9,
1239     0x0e5361dfdacefe39,
1240     0xcec7f3049fc90161,
1241     0xff62b561677f5f2e,
1242     0x975ccf26d22587f0,
1243     0x51ef0f86543baf63,
1244     0x2f1e41ef10cbf28f,
1245     0x52722635bbb94a88,
1246     0xae8dbae73344f04d,
1247     0x410769d36688fd9a,
1248     0xb3ab94de34bbb966,
1249     0x801317928df1aa9b,
1250     0xa564a0f0c5113c54,
1251     0xf131d4bebdb1a117,
1252     0x7f71a2f3ea8ef5b5,
1253     0x40878549c8f655c3,
1254     0x7ef14e6944f05dec,
1255     0xd44663dcf55137d8,
1256     0xf2acfd0d523344fc,
1257     0x0000000000000000,
1258     0x5fbc6e598ef5515a,
1259     0x16cf342ef1aa8532,
1260     0xb036bd6ddb395c8d,
1261     0x13754fe6dd31b712,
1262     0xbbdfa77a2d6c9094,
1263     0x89e7c8ac3a582b30,
1264     0x3c6b0e09cdfa459d,
1265     0xc4ae0589c7e26521,
1266     0x49735a777f5fd468,
1267     0xcafd64561d2c9b18,
1268     0xda1502032f9fc9e1,
1269     0x8867243694268369,
1270     0x3782141e3baf8984,
1271     0x9cb5d53124704be9,
1272     0xd7db4a6f1ad3d233,
1273     0xa6f989432a93d9bf,
1274     0x9d3539ab8a0ee3b0,
1275     0x53f2caaf15c7e2d1,
1276     0x6e19283c76430f15,
1277     0x3debe2936384edc4,
1278     0x5e3c82c3208bf903,
1279     0x33b8834cb94a13fd,
1280     0x6470deb12e686b55,
1281     0x359fd1377a53c436,
1282     0x61caa57902f35975,
1283     0x043a975282e59a79,
1284     0xfd7f70482683129c,
1285     0xc52ee913699ccd78,
1286     0x28b9ff0e7dac8d1d,
1287     0x5455744e78a09d43,
1288     0xcb7d88ccb3523341,
1289     0x44bd121b4a13cfba,
1290     0x4d49cd25fdba4e11,
1291     0x3e76cb208c06082f,
1292     0x3ff627ba2278a076,
1293     0xc28957f204fbb2ea,
1294     0x453dfe81e46d67e3,
1295     0x94c1e6953da7621b,
1296     0x2c83685cff491764,
1297     0xf32c1197fc4deca5,
1298     0x2b24d6bd922e68f6,
1299     0xb22b78449ac5113f,
1300     0x48f3b6edd1217c31,
1301     0x2e9ead75beb55ad6,
1302     0x174fd8b45fd42d6b,
1303     0x4ed4e4961238abfa,
1304     0x92e6b4eefebeb5d0,
1305     0x46a0d7320bef8208,
1306     0x47203ba8a5912a51,
1307     0x24f75bf8e69e3e96,
1308     0xf0b1382413cf094e,
1309     0xfee259fbc901f777,
1310     0x276a724b091cdb7d,
1311     0xbdf8f501ee75475f,
1312     0x599b3c224dec8691,
1313     0x6d84018f99c1eafe,
1314     0x7498b8e41cdb39ac,
1315     0xe0595e71217c5bb7,
1316     0x2aa43a273c50c0af,
1317     0xf50b43ec3f543b6e,
1318     0x838e3e2162734f70,
1319     0xc09492db4507ff58,
1320     0x72bfea9fdfc2ee67,
1321     0x11688acf9ccdfaa0,
1322     0x1a8190d86a9836b9,
1323     0x7acbd93bc615c795,
1324     0xc7332c3a286080ca,
1325     0x863445e94ee87d50,
1326     0xf6966a5fd0d6de85,
1327     0xe9ad814f96d5da1c,
1328     0x70a22fb69e3ea3d5,
1329     0x0a69f68d582b6440,
1330     0xb8428ec9c2ee757f,
1331     0x604a49e3ac8df12c,
1332     0x5b86f90b0c10cb23,
1333     0xe1d9b2eb8f02f3ee,
1334     0x29391394d3d22544,
1335     0xc8e0a17f5cd0d6aa,
1336     0xb58cc6a5f7a26ead,
1337     0x8193fb08238f02c2,
1338     0xd5c68f465b2f9f81,
1339     0xfcff9cd288fdbac5,
1340     0x77059157f359dc47,
1341     0x1d262e3907ff492b,
1342     0xfb582233e59ac557,
1343     0xddb2bce242f8b673,
1344     0x2577b76248e096cf,
1345     0x6f99c4a6d83da74c,
1346     0xc1147e41eb795701,
1347     0xf48baf76912a9337,
1348   },
1349   {
1350     0x45b268a93acde4cc,
1351     0xaf7f0be884549d08,
1352     0x048354b3c1468263,
1353     0x925435c2c80efed2,
1354     0xee4e37f27fdffba7,
1355     0x167a33920c60f14d,
1356     0xfb123b52ea03e584,
1357     0x4a0cab53fdbb9007,
1358     0x9deaf6380f788a19,
1359     0xcb48ec558f0cb32a,
1360     0xb59dc4b2d6fef7e0,
1361     0xdcdbca22f4f3ecb6,
1362     0x11df5813549a9c40,
1363     0xe33fdedf568aced3,
1364     0xa0c1c8124322e9c3,
1365     0x07a56b8158fa6d0d,
1366     0x77279579b1e1f3dd,
1367     0xd9b18b74422ac004,
1368     0xb8ec2d9fffabc294,
1369     0xf4acf8a82d75914f,
1370     0x7bbf69b1ef2b6878,
1371     0xc4f62faf487ac7e1,
1372     0x76ce809cc67e5d0c,
1373     0x6711d88f92e4c14c,
1374     0x627b99d9243dedfe,
1375     0x234aa5c3dfb68b51,
1376     0x909b1f15262dbf6d,
1377     0x4f66ea054b62bcb5,
1378     0x1ae2cf5a52aa6ae8,
1379     0xbea053fbd0ce0148,
1380     0xed6808c0e66314c9,
1381     0x43fe16cd15a82710,
1382     0xcd049231a06970f6,
1383     0xe7bc8a6c97cc4cb0,
1384     0x337ce835fcb3b9c0,
1385     0x65def2587cc780f3,
1386     0x52214ede4132bb50,
1387     0x95f15e4390f493df,
1388     0x870839625dd2e0f1,
1389     0x41313c1afb8b66af,
1390     0x91720af051b211bc,
1391     0x477d427ed4eea573,
1392     0x2e3b4ceef6e3be25,
1393     0x82627834eb0bcc43,
1394     0x9c03e3dd78e724c8,
1395     0x2877328ad9867df9,
1396     0x14b51945e243b0f2,
1397     0x574b0f88f7eb97e2,
1398     0x88b6fa989aa4943a,
1399     0x19c4f068cb168586,
1400     0x50ee6409af11faef,
1401     0x7df317d5c04eaba4,
1402     0x7a567c5498b4c6a9,
1403     0xb6bbfb804f42188e,
1404     0x3cc22bcf3bc5cd0b,
1405     0xd04336eaaa397713,
1406     0xf02fac1bec33132c,
1407     0x2506dba7f0d3488d,
1408     0xd7e65d6bf2c31a1e,
1409     0x5eb9b2161ff820f5,
1410     0x842e0650c46e0f9f,
1411     0x716beb1d9e843001,
1412     0xa933758cab315ed4,
1413     0x3fe414fda2792265,
1414     0x27c9f1701ef00932,
1415     0x73a4c1ca70a771be,
1416     0x94184ba6e76b3d0e,
1417     0x40d829ff8c14c87e,
1418     0x0fbec3fac77674cb,
1419     0x3616a9634a6a9572,
1420     0x8f139119c25ef937,
1421     0xf545ed4d5aea3f9e,
1422     0xe802499650ba387b,
1423     0x6437e7bd0b582e22,
1424     0xe6559f89e053e261,
1425     0x80ad52e305288dfc,
1426     0x6dc55a23e34b9935,
1427     0xde14e0f51ad0ad09,
1428     0xc6390578a659865e,
1429     0x96d7617109487cb1,
1430     0xe2d6cb3a21156002,
1431     0x01e915e5779faed1,
1432     0xadb0213f6a77dcb7,
1433     0x9880b76eb9a1a6ab,
1434     0x5d9f8d248644cf9b,
1435     0xfd5e4536c5662658,
1436     0xf1c6b9fe9bacbdfd,
1437     0xeacd6341be9979c4,
1438     0xefa7221708405576,
1439     0x510771ecd88e543e,
1440     0xc2ba51cb671f043d,
1441     0x0ad482ac71af5879,
1442     0xfe787a045cdac936,
1443     0xb238af338e049aed,
1444     0xbd866cc94972ee26,
1445     0x615da6ebbd810290,
1446     0x3295fdd08b2c1711,
1447     0xf834046073bf0aea,
1448     0xf3099329758ffc42,
1449     0x1caeb13e7dcfa934,
1450     0xba2307481188832b,
1451     0x24efce42874ce65c,
1452     0x0e57d61fb0e9da1a,
1453     0xb3d1bad6f99b343c,
1454     0xc0757b1c893c4582,
1455     0x2b510db8403a9297,
1456     0x5c7698c1f1db614a,
1457     0x3e0d0118d5e68cb4,
1458     0xd60f488e855cb4cf,
1459     0xae961e0df3cb33d9,
1460     0x3a8e55ab14a00ed7,
1461     0x42170328623789c1,
1462     0x838b6dd19c946292,
1463     0x895fef7ded3b3aeb,
1464     0xcfcbb8e64e4a3149,
1465     0x064c7e642f65c3dc,
1466     0x3d2b3e2a4c5a63da,
1467     0x5bd3f340a9210c47,
1468     0xb474d157a1615931,
1469     0xac5934da1de87266,
1470     0x6ee365117af7765b,
1471     0xc86ed36716b05c44,
1472     0x9ba6885c201d49c5,
1473     0xb905387a88346c45,
1474     0x131072c4bab9ddff,
1475     0xbf49461ea751af99,
1476     0xd52977bc1ce05ba1,
1477     0xb0f785e46027db52,
1478     0x546d30ba6e57788c,
1479     0x305ad707650f56ae,
1480     0xc987c682612ff295,
1481     0xa5ab8944f5fbc571,
1482     0x7ed528e759f244ca,
1483     0x8ddcbbce2c7db888,
1484     0xaa154abe328db1ba,
1485     0x1e619be993ece88b,
1486     0x09f2bd9ee813b717,
1487     0x7401aa4b285d1cb3,
1488     0x21858f143195caee,
1489     0x48c381841398d1b8,
1490     0xfcb750d3b2f98889,
1491     0x39a86a998d1ce1b9,
1492     0x1f888e0ce473465a,
1493     0x7899568376978716,
1494     0x02cf2ad7ee2341bf,
1495     0x85c713b5b3f1a14e,
1496     0xff916fe12b4567e7,
1497     0x7c1a0230b7d10575,
1498     0x0c98fcc85eca9ba5,
1499     0xa3e7f720da9e06ad,
1500     0x6a6031a2bbb1f438,
1501     0x973e74947ed7d260,
1502     0x2cf4663918c0ff9a,
1503     0x5f50a7f368678e24,
1504     0x34d983b4a449d4cd,
1505     0x68af1b755592b587,
1506     0x7f3c3d022e6dea1b,
1507     0xabfc5f5b45121f6b,
1508     0x0d71e92d29553574,
1509     0xdffdf5106d4f03d8,
1510     0x081ba87b9f8c19c6,
1511     0xdb7ea1a3ac0981bb,
1512     0xbbca12ad66172dfa,
1513     0x79704366010829c7,
1514     0x179326777bff5f9c,
1515     0x0000000000000000,
1516     0xeb2476a4c906d715,
1517     0x724dd42f0738df6f,
1518     0xb752ee6538ddb65f,
1519     0x37ffbc863df53ba3,
1520     0x8efa84fcb5c157e6,
1521     0xe9eb5c73272596aa,
1522     0x1b0bdabf2535c439,
1523     0x86e12c872a4d4e20,
1524     0x9969a28bce3e087a,
1525     0xfafb2eb79d9c4b55,
1526     0x056a4156b6d92cb2,
1527     0x5a3ae6a5debea296,
1528     0x22a3b026a8292580,
1529     0x53c85b3b36ad1581,
1530     0xb11e900117b87583,
1531     0xc51f3a4a3fe56930,
1532     0xe019e1edcf3621bd,
1533     0xec811d2591fcba18,
1534     0x445b7d4c4d524a1d,
1535     0xa8da6069dcaef005,
1536     0x58f5cc72309de329,
1537     0xd4c062596b7ff570,
1538     0xce22ad0339d59f98,
1539     0x591cd99747024df8,
1540     0x8b90c5aa03187b54,
1541     0xf663d27fc356d0f0,
1542     0xd8589e9135b56ed5,
1543     0x35309651d3d67a1c,
1544     0x12f96721cd26732e,
1545     0xd28c1c3d441a36ac,
1546     0x492a946164077f69,
1547     0x2d1d73dc6f5f514b,
1548     0x6f0a70f40d68d88a,
1549     0x60b4b30eca1eac41,
1550     0xd36509d83385987d,
1551     0x0b3d97490630f6a8,
1552     0x9eccc90a96c46577,
1553     0xa20ee2c5ad01a87c,
1554     0xe49ab55e0e70a3de,
1555     0xa4429ca182646ba0,
1556     0xda97b446db962f6a,
1557     0xcced87d4d7f6de27,
1558     0x2ab8185d37a53c46,
1559     0x9f25dcefe15bcba6,
1560     0xc19c6ef9fea3eb53,
1561     0xa764a3931bd884ce,
1562     0x2fd2590b817c10f4,
1563     0x56a21a6d80743933,
1564     0xe573a0bb79ef0d0f,
1565     0x155c0ca095dc1e23,
1566     0x6c2c4fc694d437e4,
1567     0x10364df623053291,
1568     0xdd32dfc7836c4267,
1569     0x03263f3299bcef6e,
1570     0x66f8cd6ae57b6f9d,
1571     0x8c35ae2b5be21659,
1572     0x31b3c2e21290f87f,
1573     0x93bd2027bf915003,
1574     0x69460e90220d1b56,
1575     0x299e276fae19d328,
1576     0x63928c3c53a2432f,
1577     0x7082fef8e91b9ed0,
1578     0xbc6f792c3eed40f7,
1579     0x4c40d537d2de53db,
1580     0x75e8bfae5fc2b262,
1581     0x4da9c0d2a541fd0a,
1582     0x4e8fffe03cfd1264,
1583     0x2620e495696fa7e3,
1584     0xe1f0f408b8a98f6c,
1585     0xd1aa230fdda6d9c2,
1586     0xc7d0109dd1c6288f,
1587     0x8a79d04f7487d585,
1588     0x4694579ba3710ba2,
1589     0x38417f7cfa834f68,
1590     0x1d47a4db0a5007e5,
1591     0x206c9af1460a643f,
1592     0xa128ddf734bd4712,
1593     0x8144470672b7232d,
1594     0xf2e086cc02105293,
1595     0x182de58dbc892b57,
1596     0xcaa1f9b0f8931dfb,
1597     0x6b892447cc2e5ae9,
1598     0xf9dd11850420a43b,
1599     0x4be5beb68a243ed6,
1600     0x5584255f19c8d65d,
1601     0x3b67404e633fa006,
1602     0xa68db6766c472a1f,
1603     0xf78ac79ab4c97e21,
1604     0xc353442e1080aaec,
1605     0x9a4f9db95782e714,
1606   },
1607   {
1608     0xc811a8058c3f55de,
1609     0x65f5b43196b50619,
1610     0xf74f96b1d6706e43,
1611     0x859d1e8bcb43d336,
1612     0x5aab8a85ccfa3d84,
1613     0xf9c7bf99c295fcfd,
1614     0xa21fd5a1de4b630f,
1615     0xcdb3ef763b8b456d,
1616     0x803f59f87cf7c385,
1617     0xb27c73be5f31913c,
1618     0x98e3ac6633b04821,
1619     0xbf61674c26b8f818,
1620     0x0ffbc995c4c130c8,
1621     0xaaa0862010761a98,
1622     0x6057f342210116aa,
1623     0xf63c760c0654cc35,
1624     0x2ddb45cc667d9042,
1625     0xbcf45a964bd40382,
1626     0x68e8a0c3ef3c6f3d,
1627     0xa7bd92d269ff73bc,
1628     0x290ae20201ed2287,
1629     0xb7de34cde885818f,
1630     0xd901eea7dd61059b,
1631     0xd6fa273219a03553,
1632     0xd56f1ae874cccec9,
1633     0xea31245c2e83f554,
1634     0x7034555da07be499,
1635     0xce26d2ac56e7bef7,
1636     0xfd161857a5054e38,
1637     0x6a0e7da4527436d1,
1638     0x5bd86a381cde9ff2,
1639     0xcaf7756231770c32,
1640     0xb09aaed9e279c8d0,
1641     0x5def1091c60674db,
1642     0x111046a2515e5045,
1643     0x23536ce4729802fc,
1644     0xc50cbcf7f5b63cfa,
1645     0x73a16887cd171f03,
1646     0x7d2941afd9f28dbd,
1647     0x3f5e3eb45a4f3b9d,
1648     0x84eefe361b677140,
1649     0x3db8e3d3e7076271,
1650     0x1a3a28f9f20fd248,
1651     0x7ebc7c75b49e7627,
1652     0x74e5f293c7eb565c,
1653     0x18dcf59e4f478ba4,
1654     0x0c6ef44fa9adcb52,
1655     0xc699812d98dac760,
1656     0x788b06dc6e469d0e,
1657     0xfc65f8ea7521ec4e,
1658     0x30a5f7219e8e0b55,
1659     0x2bec3f65bca57b6b,
1660     0xddd04969baf1b75e,
1661     0x99904cdbe394ea57,
1662     0x14b201d1e6ea40f6,
1663     0xbbb0c08241284add,
1664     0x50f20463bf8f1dff,
1665     0xe8d7f93b93cbacb8,
1666     0x4d8cb68e477c86e8,
1667     0xc1dd1b3992268e3f,
1668     0x7c5aa11209d62fcb,
1669     0x2f3d98abdb35c9ae,
1670     0x671369562bfd5ff5,
1671     0x15c1e16c36cee280,
1672     0x1d7eb2edf8f39b17,
1673     0xda94d37db00dfe01,
1674     0x877bc3ec760b8ada,
1675     0xcb8495dfe153ae44,
1676     0x05a24773b7b410b3,
1677     0x12857b783c32abdf,
1678     0x8eb770d06812513b,
1679     0x536739b9d2e3e665,
1680     0x584d57e271b26468,
1681     0xd789c78fc9849725,
1682     0xa935bbfa7d1ae102,
1683     0x8b1537a3dfa64188,
1684     0xd0cd5d9bc378de7a,
1685     0x4ac82c9a4d80cfb7,
1686     0x42777f1b83bdb620,
1687     0x72d2883a1d33bd75,
1688     0x5e7a2d4bab6a8f41,
1689     0xf4daab6bbb1c95d9,
1690     0x905cffe7fd8d31b6,
1691     0x83aa6422119b381f,
1692     0xc0aefb8442022c49,
1693     0xa0f908c663033ae3,
1694     0xa428af0804938826,
1695     0xade41c341a8a53c7,
1696     0xae7121ee77e6a85d,
1697     0xc47f5c4a25929e8c,
1698     0xb538e9aa55cdd863,
1699     0x06377aa9dad8eb29,
1700     0xa18ae87bb3279895,
1701     0x6edfda6a35e48414,
1702     0x6b7d9d19825094a7,
1703     0xd41cfa55a4e86cbf,
1704     0xe5caedc9ea42c59c,
1705     0xa36c351c0e6fc179,
1706     0x5181e4de6fabbf89,
1707     0xfff0c530184d17d4,
1708     0x9d41eb1584045892,
1709     0x1c0d525028d73961,
1710     0xf178ec180ca8856a,
1711     0x9a0571018ef811cd,
1712     0x4091a27c3ef5efcc,
1713     0x19af15239f6329d2,
1714     0x347450eff91eb990,
1715     0xe11b4a078dd27759,
1716     0xb9561de5fc601331,
1717     0x912f1f5a2da993c0,
1718     0x1654dcb65ba2191a,
1719     0x3e2dde098a6b99eb,
1720     0x8a66d71e0f82e3fe,
1721     0x8c51adb7d55a08d7,
1722     0x4533e50f8941ff7f,
1723     0x02e6dd67bd4859ec,
1724     0xe068aaba5df6d52f,
1725     0xc24826e3ff4a75a5,
1726     0x6c39070d88acddf8,
1727     0x6486548c4691a46f,
1728     0xd1bebd26135c7c0c,
1729     0xb30f93038f15334a,
1730     0x82d9849fc1bf9a69,
1731     0x9c320ba85420fae4,
1732     0xfa528243aff90767,
1733     0x9ed4d6cfe968a308,
1734     0xb825fd582c44b147,
1735     0x9b7691bc5edcb3bb,
1736     0xc7ea619048fe6516,
1737     0x1063a61f817af233,
1738     0x47d538683409a693,
1739     0x63c2ce984c6ded30,
1740     0x2a9fdfd86c81d91d,
1741     0x7b1e3b06032a6694,
1742     0x666089ebfbd9fd83,
1743     0x0a598ee67375207b,
1744     0x07449a140afc495f,
1745     0x2ca8a571b6593234,
1746     0x1f986f8a45bbc2fb,
1747     0x381aa4a050b372c2,
1748     0x5423a3add81faf3a,
1749     0x17273c0b8b86bb6c,
1750     0xfe83258dc869b5a2,
1751     0x287902bfd1c980f1,
1752     0xf5a94bd66b3837af,
1753     0x88800a79b2caba12,
1754     0x55504310083b0d4c,
1755     0xdf36940e07b9eeb2,
1756     0x04d1a7ce6790b2c5,
1757     0x612413fff125b4dc,
1758     0x26f12b97c52c124f,
1759     0x86082351a62f28ac,
1760     0xef93632f9937e5e7,
1761     0x3507b052293a1be6,
1762     0xe72c30ae570a9c70,
1763     0xd3586041ae1425e0,
1764     0xde4574b3d79d4cc4,
1765     0x92ba228040c5685a,
1766     0xf00b0ca5dc8c271c,
1767     0xbe1287f1f69c5a6e,
1768     0xf39e317fb1e0dc86,
1769     0x495d114020ec342d,
1770     0x699b407e3f18cd4b,
1771     0xdca3a9d46ad51528,
1772     0x0d1d14f279896924,
1773     0x0000000000000000,
1774     0x593eb75fa196c61e,
1775     0x2e4e78160b116bd8,
1776     0x6d4ae7b058887f8e,
1777     0xe65fd013872e3e06,
1778     0x7a6ddbbbd30ec4e2,
1779     0xac97fc89caaef1b1,
1780     0x09ccb33c1e19dbe1,
1781     0x89f3eac462ee1864,
1782     0x7770cf49aa87adc6,
1783     0x56c57eca6557f6d6,
1784     0x03953dda6d6cfb9a,
1785     0x36928d884456e07c,
1786     0x1eeb8f37959f608d,
1787     0x31d6179c4eaaa923,
1788     0x6fac3ad7e5c02662,
1789     0x43049fa653991456,
1790     0xabd3669dc052b8ee,
1791     0xaf02c153a7c20a2b,
1792     0x3ccb036e3723c007,
1793     0x93c9c23d90e1ca2c,
1794     0xc33bc65e2f6ed7d3,
1795     0x4cff56339758249e,
1796     0xb1e94e64325d6aa6,
1797     0x37e16d359472420a,
1798     0x79f8e661be623f78,
1799     0x5214d90402c74413,
1800     0x482ef1fdf0c8965b,
1801     0x13f69bc5ec1609a9,
1802     0x0e88292814e592be,
1803     0x4e198b542a107d72,
1804     0xccc00fcbebafe71b,
1805     0x1b49c844222b703e,
1806     0x2564164da840e9d5,
1807     0x20c6513e1ff4f966,
1808     0xbac3203f910ce8ab,
1809     0xf2edd1c261c47ef0,
1810     0x814cb945acd361f3,
1811     0x95feb8944a392105,
1812     0x5c9cf02c1622d6ad,
1813     0x971865f3f77178e9,
1814     0xbd87ba2b9bf0a1f4,
1815     0x444005b259655d09,
1816     0xed75be48247fbc0b,
1817     0x7596122e17cff42a,
1818     0xb44b091785e97a15,
1819     0x966b854e2755da9f,
1820     0xeee0839249134791,
1821     0x32432a4623c652b9,
1822     0xa8465b47ad3e4374,
1823     0xf8b45f2412b15e8b,
1824     0x2417f6f078644ba3,
1825     0xfb2162fe7fdda511,
1826     0x4bbbcc279da46dc1,
1827     0x0173e0bdd024a276,
1828     0x22208c59a2bca08a,
1829     0x8fc4906db836f34d,
1830     0xe4b90d743a6667ea,
1831     0x7147b5e0705f46ef,
1832     0x2782cb2a1508b039,
1833     0xec065ef5f45b1e7d,
1834     0x21b5b183cfd05b10,
1835     0xdbe733c060295c77,
1836     0x9fa73672394c017e,
1837     0xcf55321186c31c81,
1838     0xd8720e1a0d45a7ed,
1839     0x3b8f997a3ddf8958,
1840     0x3afc79c7edfb2b2e,
1841     0xe9a4198643ef0ece,
1842     0x5f09cdf67b4e2d37,
1843     0x4f6a6be9fa34df04,
1844     0xb6add47038a123f9,
1845     0x8d224d0a057eaaa1,
1846     0xc96248b85c1bf7a8,
1847     0xe3fd9760309a2eb5,
1848     0x0b2a6e5ba351820d,
1849     0xeb42c4e1fea75722,
1850     0x948d58299a1d8373,
1851     0x7fcf9cc864bad451,
1852     0xa55b4fb5d4b72a50,
1853     0x08bf5381ce3d7997,
1854     0x46a6d8d5e42d04e5,
1855     0xd22b80fc7e308796,
1856     0x57b69e77b57354a0,
1857     0x3969441d8097d0b4,
1858     0x3330cafbf3e2f0cf,
1859     0xe28e77dde0be8cc3,
1860     0x62b12e259c494f46,
1861     0xa6ce726fb9dbd1ca,
1862     0x41e242c1eed14dba,
1863     0x76032ff47aa30fb0,
1864   },
1865   {
1866     0xe6f87e5c5b711fd0,
1867     0x258377800924fa16,
1868     0xc849e07e852ea4a8,
1869     0x5b4686a18f06c16a,
1870     0x0b32e9a2d77b416e,
1871     0xabda37a467815c66,
1872     0xf61796a81a686676,
1873     0xf5dc0b706391954b,
1874     0x4862f38db7e64bf1,
1875     0xff5c629a68bd85c5,
1876     0xcb827da6fcd75795,
1877     0x66d36daf69b9f089,
1878     0x356c9f74483d83b0,
1879     0x7cbcecb1238c99a1,
1880     0x36a702ac31c4708d,
1881     0x9eb6a8d02fbcdfd6,
1882     0x8b19fa51e5b3ae37,
1883     0x9ccfb5408a127d0b,
1884     0xbc0c78b508208f5a,
1885     0xe533e3842288eced,
1886     0xcec2c7d377c15fd2,
1887     0xec7817b6505d0f5e,
1888     0xb94cc2c08336871d,
1889     0x8c205db4cb0b04ad,
1890     0x763c855b28a0892f,
1891     0x588d1b79f6ff3257,
1892     0x3fecf69e4311933e,
1893     0x0fc0d39f803a18c9,
1894     0xee010a26f5f3ad83,
1895     0x10efe8f4411979a6,
1896     0x5dcda10c7de93a10,
1897     0x4a1bee1d1248e92c,
1898     0x53bff2db21847339,
1899     0xb4f50ccfa6a23d09,
1900     0x5fb4bc9cd84798cd,
1901     0xe88a2d8b071c56f9,
1902     0x7f7771695a756a9c,
1903     0xc5f02e71a0ba1ebc,
1904     0xa663f9ab4215e672,
1905     0x2eb19e22de5fbb78,
1906     0x0db9ce0f2594ba14,
1907     0x82520e6397664d84,
1908     0x2f031e6a0208ea98,
1909     0x5c7f2144a1be6bf0,
1910     0x7a37cb1cd16362db,
1911     0x83e08e2b4b311c64,
1912     0xcf70479bab960e32,
1913     0x856ba986b9dee71e,
1914     0xb5478c877af56ce9,
1915     0xb8fe42885f61d6fd,
1916     0x1bdd0156966238c8,
1917     0x622157923ef8a92e,
1918     0xfc97ff42114476f8,
1919     0x9d7d350856452ceb,
1920     0x4c90c9b0e0a71256,
1921     0x2308502dfbcb016c,
1922     0x2d7a03faa7a64845,
1923     0xf46e8b38bfc6c4ab,
1924     0xbdbef8fdd477deba,
1925     0x3aac4cebc8079b79,
1926     0xf09cb105e8879d0c,
1927     0x27fa6a10ac8a58cb,
1928     0x8960e7c1401d0cea,
1929     0x1a6f811e4a356928,
1930     0x90c4fb0773d196ff,
1931     0x43501a2f609d0a9f,
1932     0xf7a516e0c63f3796,
1933     0x1ce4a6b3b8da9252,
1934     0x1324752c38e08a9b,
1935     0xa5a864733bec154f,
1936     0x2bf124575549b33f,
1937     0xd766db15440dc5c7,
1938     0xa7d179e39e42b792,
1939     0xdadf151a61997fd3,
1940     0x86a0345ec0271423,
1941     0x38d5517b6da939a4,
1942     0x6518f077104003b4,
1943     0x02791d90a5aea2dd,
1944     0x88d267899c4a5d0a,
1945     0x930f66df0a2865c2,
1946     0x4ee9d4204509b08b,
1947     0x325538916685292a,
1948     0x412907bfc533a842,
1949     0xb27e2b62544dc673,
1950     0x6c5304456295e007,
1951     0x5af406e95351908a,
1952     0x1f2f3b6bc123616f,
1953     0xc37b09dc5255e5c6,
1954     0x3967d133b1fe6844,
1955     0x298839c7f0e711e2,
1956     0x409b87f71964f9a2,
1957     0xe938adc3db4b0719,
1958     0x0c0b4e47f9c3ebf4,
1959     0x5534d576d36b8843,
1960     0x4610a05aeb8b02d8,
1961     0x20c3cdf58232f251,
1962     0x6de1840dbec2b1e7,
1963     0xa0e8de06b0fa1d08,
1964     0x7b854b540d34333b,
1965     0x42e29a67bcca5b7f,
1966     0xd8a6088ac437dd0e,
1967     0xc63bb3a9d943ed81,
1968     0x21714dbd5e65a3b1,
1969     0x6761ede7b5eea169,
1970     0x2431f7c8d573abf6,
1971     0xd51fc685e1a3671a,
1972     0x5e063cd40410c92d,
1973     0x283ab98f2cb04002,
1974     0x8febc06cb2f2f790,
1975     0x17d64f116fa1d33c,
1976     0xe07359f1a99ee4aa,
1977     0x784ed68c74cdc006,
1978     0x6e2a19d5c73b42da,
1979     0x8712b4161c7045c3,
1980     0x371582e4ed93216d,
1981     0xace390414939f6fc,
1982     0x7ec5f12186223b7c,
1983     0xc0b094042bac16fb,
1984     0xf9d745379a527ebf,
1985     0x737c3f2ea3b68168,
1986     0x33e7b8d9bad278ca,
1987     0xa9a32a34c22ffebb,
1988     0xe48163ccfedfbd0d,
1989     0x8e5940246ea5a670,
1990     0x51c6ef4b842ad1e4,
1991     0x22bad065279c508c,
1992     0xd91488c218608cee,
1993     0x319ea5491f7cda17,
1994     0xd394e128134c9c60,
1995     0x094bf43272d5e3b3,
1996     0x9bf612a5a4aad791,
1997     0xccbbda43d26ffd0f,
1998     0x34de1f3c946ad250,
1999     0x4f5b5468995ee16b,
2000     0xdf9faf6fea8f7794,
2001     0x2648ea5870dd092b,
2002     0xbfc7e56d71d97c67,
2003     0xdde6b2ff4f21d549,
2004     0x3c276b463ae86003,
2005     0x91767b4faf86c71f,
2006     0x68a13e7835d4b9a0,
2007     0xb68c115f030c9fd4,
2008     0x141dd2c916582001,
2009     0x983d8f7ddd5324ac,
2010     0x64aa703fcc175254,
2011     0xc2c989948e02b426,
2012     0x3e5e76d69f46c2de,
2013     0x50746f03587d8004,
2014     0x45db3d829272f1e5,
2015     0x60584a029b560bf3,
2016     0xfbae58a73ffcdc62,
2017     0xa15a5e4e6cad4ce8,
2018     0x4ba96e55ce1fb8cc,
2019     0x08f9747aae82b253,
2020     0xc102144cf7fb471b,
2021     0x9f042898f3eb8e36,
2022     0x068b27adf2effb7a,
2023     0xedca97fe8c0a5ebe,
2024     0x778e0513f4f7d8cf,
2025     0x302c2501c32b8bf7,
2026     0x8d92ddfc175c554d,
2027     0xf865c57f46052f5f,
2028     0xeaf3301ba2b2f424,
2029     0xaa68b7ecbbd60d86,
2030     0x998f0f350104754c,
2031     0x0000000000000000,
2032     0xf12e314d34d0ccec,
2033     0x710522be061823b5,
2034     0xaf280d9930c005c1,
2035     0x97fd5ce25d693c65,
2036     0x19a41cc633cc9a15,
2037     0x95844172f8c79eb8,
2038     0xdc5432b7937684a9,
2039     0x9436c13a2490cf58,
2040     0x802b13f332c8ef59,
2041     0xc442ae397ced4f5c,
2042     0xfa1cd8efe3ab8d82,
2043     0xf2e5ac954d293fd1,
2044     0x6ad823e8907a1b7d,
2045     0x4d2249f83cf043b6,
2046     0x03cb9dd879f9f33d,
2047     0xde2d2f2736d82674,
2048     0x2a43a41f891ee2df,
2049     0x6f98999d1b6c133a,
2050     0xd4ad46cd3df436fa,
2051     0xbb35df50269825c0,
2052     0x964fdcaa813e6d85,
2053     0xeb41b0537ee5a5c4,
2054     0x0540ba758b160847,
2055     0xa41ae43be7bb44af,
2056     0xe3b8c429d0671797,
2057     0x819993bbee9fbeb9,
2058     0xae9a8dd1ec975421,
2059     0xf3572cdd917e6e31,
2060     0x6393d7dae2aff8ce,
2061     0x47a2201237dc5338,
2062     0xa32343dec903ee35,
2063     0x79fc56c4a89a91e6,
2064     0x01b28048dc5751e0,
2065     0x1296f564e4b7db7b,
2066     0x75f7188351597a12,
2067     0xdb6d9552bdce2e33,
2068     0x1e9dbb231d74308f,
2069     0x520d7293fdd322d9,
2070     0xe20a44610c304677,
2071     0xfeeee2d2b4ead425,
2072     0xca30fdee20800675,
2073     0x61eaca4a47015a13,
2074     0xe74afe1487264e30,
2075     0x2cc883b27bf119a5,
2076     0x1664cf59b3f682dc,
2077     0xa811aa7c1e78af5b,
2078     0x1d5626fb648dc3b2,
2079     0xb73e9117df5bce34,
2080     0xd05f7cf06ab56f5d,
2081     0xfd257f0acd132718,
2082     0x574dc8e676c52a9e,
2083     0x0739a7e52eb8aa9a,
2084     0x5486553e0f3cd9a3,
2085     0x56ff48aeaa927b7e,
2086     0xbe756525ad8e2d87,
2087     0x7d0e6cf9ffdbc841,
2088     0x3b1ecca31450ca99,
2089     0x6913be30e983e840,
2090     0xad511009956ea71c,
2091     0xb1b5b6ba2db4354e,
2092     0x4469bdca4e25a005,
2093     0x15af5281ca0f71e1,
2094     0x744598cb8d0e2bf2,
2095     0x593f9b312aa863b7,
2096     0xefb38a6e29a4fc63,
2097     0x6b6aa3a04c2d4a9d,
2098     0x3d95eb0ee6bf31e3,
2099     0xa291c3961554bfd5,
2100     0x18169c8eef9bcbf5,
2101     0x115d68bc9d4e2846,
2102     0xba875f18facf7420,
2103     0xd1edfcb8b6e23ebd,
2104     0xb00736f2f1e364ae,
2105     0x84d929ce6589b6fe,
2106     0x70b7a2f6da4f7255,
2107     0x0e7253d75c6d4929,
2108     0x04f23a3d574159a7,
2109     0x0a8069ea0b2c108e,
2110     0x49d073c56bb11a11,
2111     0x8aab7a1939e4ffd7,
2112     0xcd095a0b0e38acef,
2113     0xc9fb60365979f548,
2114     0x92bde697d67f3422,
2115     0xc78933e10514bc61,
2116     0xe1c1d9b975c9b54a,
2117     0xd2266160cf1bcd80,
2118     0x9a4492ed78fd8671,
2119     0xb3ccab2a881a9793,
2120     0x72cebf667fe1d088,
2121     0xd6d45b5d985a9427,
2122   },
2123 };
2124
2125 __device__ __constant__ u64 sbob_rc64[12][8] =
2126 {
2127   {
2128     0xe9daca1eda5b08b1,
2129     0x1f7c65c0812fcbeb,
2130     0x16d0452e43766a2f,
2131     0xfcc485758db84e71,
2132     0x0169679291e07c4b,
2133     0x15d360a4082a42a2,
2134     0x234d74cc36747605,
2135     0x0745a6f2596580dd,
2136   },
2137   {
2138     0x1a2f9da98ab5a36f,
2139     0xd7b5700f469de34f,
2140     0x982b230a72eafef3,
2141     0x3101b5160f5ed561,
2142     0x5899d6126b17b59a,
2143     0xcaa70adbc261b55c,
2144     0x56cdcbd71ba2dd55,
2145     0xb79bb121700479e6,
2146   },
2147   {
2148     0xc72fce2bacdc74f5,
2149     0x35843d6a28fc390a,
2150     0x8b1f9c525f5ef106,
2151     0x7b7b29b11475eaf2,
2152     0xb19e3590e40fe2d3,
2153     0x09db6260373ac9c1,
2154     0x31db7a8643f4b6c2,
2155     0xb20aba0af5961e99,
2156   },
2157   {
2158     0xd26615e8b3df1fef,
2159     0xdde4715da0e148f9,
2160     0x7d3c5c337e858e48,
2161     0x3f355e68ad1c729d,
2162     0x75d603ed822cd7a9,
2163     0xbe0352933313b7d8,
2164     0xf137e893a1ea5334,
2165     0x2ed1e384bcbe0c22,
2166   },
2167   {
2168     0x994747adac6bea4b,
2169     0x6323a96c0c413f9a,
2170     0x4a1086161f1c157f,
2171     0xbdff0f80d7359e35,
2172     0xa3f53a254717cdbf,
2173     0x161a2723b700ffdf,
2174     0xf563eaa97ea2567a,
2175     0x57fe6c7cfd581760,
2176   },
2177   {
2178     0xd9d33a1daeae4fae,
2179     0xc039307a3bc3a46f,
2180     0x6ca44251f9c4662d,
2181     0xc68ef09ab49a7f18,
2182     0xb4b79a1cb7a6facf,
2183     0xb6c6bec2661ff20a,
2184     0x354f903672c571bf,
2185     0x6e7d64467a4068fa,
2186   },
2187   {
2188     0xecc5aaee160ec7f4,
2189     0x540924bffe86ac51,
2190     0xc987bfe6c7c69e39,
2191     0xc9937a19333e47d3,
2192     0x372c822dc5ab9209,
2193     0x04054a2883694706,
2194     0xf34a3ca24c451735,
2195     0x93d4143a4d568688,
2196   },
2197   {
2198     0xa7c9934d425b1f9b,
2199     0x41416e0c02aae703,
2200     0x1ede369c71f8b74e,
2201     0x9ac4db4d3b44b489,
2202     0x90069b92cb2b89f4,
2203     0x2fc4a5d12b8dd169,
2204     0xd9a8515935c2ac36,
2205     0x1ee702bfd40d7fa4,
2206   },
2207   {
2208     0x9b223116545a8f37,
2209     0xde5f16ecd89a4c94,
2210     0x244289251b3a7d3a,
2211     0x84090de0b755d93c,
2212     0xb1ceb2db0b440a80,
2213     0x549c07a69a8a2b7b,
2214     0x602a1fcb92dc380e,
2215     0xdb5a238351446172,
2216   },
2217   {
2218     0x526f0580a6debeab,
2219     0xf3f3e4b248e52a38,
2220     0xdb788aff1ce74189,
2221     0x0361331b8ae1ff1f,
2222     0x4b3369af0267e79f,
2223     0xf452763b306c1e7a,
2224     0xc3b63b15d1fa9836,
2225     0xed9c4598fbc7b474,
2226   },
2227   {
2228     0xfb89c8efd09ecd7b,
2229     0x94fe5a63cdc60230,
2230     0x6107abebbb6bfad8,
2231     0x7966841421800120,
2232     0xcab948eaef711d8a,
2233     0x986e477d1dcdbaef,
2234     0x5dd86fc04a59a2de,
2235     0x1b2df381cda4ca6b,
2236   },
2237   {
2238     0xba3116f167e78e37,
2239     0x7ab14904b08013d2,
2240     0x771ddfbc323ca4cd,
2241     0x9b9f2130d41220f8,
2242     0x86cc91189def805d,
2243     0x5228e188aaa41de7,
2244     0x991bb2d9d517f4fa,
2245     0x20d71bf14a92bc48,
2246   },
2247 };
2248
2249 __device__ static void streebog_g (u64 h[8], const u64 m[8], u64 s_sbob_sl64[8][256])
2250 {
2251   u64 k[8];
2252   u64 s[8];
2253   u64 t[8];
2254
2255   #pragma unroll
2256   for (int i = 0; i < 8; i++)
2257   {
2258     t[i] = h[i];
2259   }
2260
2261   for (int i = 0; i < 8; i++)
2262   {
2263     k[i] = SBOG_LPSti64;
2264   }
2265
2266   #pragma unroll
2267   for (int i = 0; i < 8; i++)
2268   {
2269     s[i] = m[i];
2270   }
2271
2272   for (int r = 0; r < 12; r++)
2273   {
2274     #pragma unroll
2275     for (int i = 0; i < 8; i++)
2276     {
2277       t[i] = s[i] ^ k[i];
2278     }
2279
2280     #pragma unroll
2281     for (int i = 0; i < 8; i++)
2282     {
2283       s[i] = SBOG_LPSti64;
2284     }
2285
2286     for (int i = 0; i < 8; i++)
2287     {
2288       t[i] = k[i] ^ sbob_rc64[r][i];
2289     }
2290
2291     #pragma unroll
2292     for (int i = 0; i < 8; i++)
2293     {
2294       k[i] = SBOG_LPSti64;
2295     }
2296   }
2297
2298   #pragma unroll
2299   for (int i = 0; i < 8; i++)
2300   {
2301     h[i] ^= s[i] ^ k[i] ^ m[i];
2302   }
2303 }
2304
2305 __device__ __constant__ bf_t c_bfs[1024];
2306
2307 __device__ static void m11700m (u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
2308 {
2309   /**
2310    * modifier
2311    */
2312
2313   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2314   const u32 lid = threadIdx.x;
2315
2316   /**
2317    * loop
2318    */
2319
2320   u32x w0l = w[0];
2321
2322   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
2323   {
2324     const u32 w0r = c_bfs[il_pos].i;
2325
2326     w[0] = w0l | w0r;
2327
2328     /**
2329      * reverse message block
2330      */
2331
2332     u64 m[8];
2333
2334     m[0] = hl32_to_64 (w[15], w[14]);
2335     m[1] = hl32_to_64 (w[13], w[12]);
2336     m[2] = hl32_to_64 (w[11], w[10]);
2337     m[3] = hl32_to_64 (w[ 9], w[ 8]);
2338     m[4] = hl32_to_64 (w[ 7], w[ 6]);
2339     m[5] = hl32_to_64 (w[ 5], w[ 4]);
2340     m[6] = hl32_to_64 (w[ 3], w[ 2]);
2341     m[7] = hl32_to_64 (w[ 1], w[ 0]);
2342
2343     m[0] = swap_workaround (m[0]);
2344     m[1] = swap_workaround (m[1]);
2345     m[2] = swap_workaround (m[2]);
2346     m[3] = swap_workaround (m[3]);
2347     m[4] = swap_workaround (m[4]);
2348     m[5] = swap_workaround (m[5]);
2349     m[6] = swap_workaround (m[6]);
2350     m[7] = swap_workaround (m[7]);
2351
2352     // state buffer (hash)
2353
2354     u64 h[8];
2355
2356     h[0] = INITVAL;
2357     h[1] = INITVAL;
2358     h[2] = INITVAL;
2359     h[3] = INITVAL;
2360     h[4] = INITVAL;
2361     h[5] = INITVAL;
2362     h[6] = INITVAL;
2363     h[7] = INITVAL;
2364
2365     streebog_g (h, m, s_sbob_sl64);
2366
2367     u64 z[8];
2368
2369     z[0] = 0;
2370     z[1] = 0;
2371     z[2] = 0;
2372     z[3] = 0;
2373     z[4] = 0;
2374     z[5] = 0;
2375     z[6] = 0;
2376     z[7] = swap_workaround ((u64) (pw_len * 8));
2377
2378     streebog_g (h, z, s_sbob_sl64);
2379     streebog_g (h, m, s_sbob_sl64);
2380
2381     const u32 r0 = l32_from_64 (h[0]);
2382     const u32 r1 = h32_from_64 (h[0]);
2383     const u32 r2 = l32_from_64 (h[1]);
2384     const u32 r3 = h32_from_64 (h[1]);
2385
2386     #include VECT_COMPARE_M
2387   }
2388 }
2389
2390 __device__ static void m11700s (u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
2391 {
2392   /**
2393    * modifier
2394    */
2395
2396   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2397   const u32 lid = threadIdx.x;
2398
2399   /**
2400    * digest
2401    */
2402
2403   const u32 search[4] =
2404   {
2405     digests_buf[digests_offset].digest_buf[DGST_R0],
2406     digests_buf[digests_offset].digest_buf[DGST_R1],
2407     digests_buf[digests_offset].digest_buf[DGST_R2],
2408     digests_buf[digests_offset].digest_buf[DGST_R3]
2409   };
2410
2411   /**
2412    * loop
2413    */
2414
2415   u32x w0l = w[0];
2416
2417   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
2418   {
2419     const u32 w0r = c_bfs[il_pos].i;
2420
2421     w[0] = w0l | w0r;
2422
2423     /**
2424      * reverse message block
2425      */
2426
2427     u64 m[8];
2428
2429     m[0] = hl32_to_64 (w[15], w[14]);
2430     m[1] = hl32_to_64 (w[13], w[12]);
2431     m[2] = hl32_to_64 (w[11], w[10]);
2432     m[3] = hl32_to_64 (w[ 9], w[ 8]);
2433     m[4] = hl32_to_64 (w[ 7], w[ 6]);
2434     m[5] = hl32_to_64 (w[ 5], w[ 4]);
2435     m[6] = hl32_to_64 (w[ 3], w[ 2]);
2436     m[7] = hl32_to_64 (w[ 1], w[ 0]);
2437
2438     m[0] = swap_workaround (m[0]);
2439     m[1] = swap_workaround (m[1]);
2440     m[2] = swap_workaround (m[2]);
2441     m[3] = swap_workaround (m[3]);
2442     m[4] = swap_workaround (m[4]);
2443     m[5] = swap_workaround (m[5]);
2444     m[6] = swap_workaround (m[6]);
2445     m[7] = swap_workaround (m[7]);
2446
2447     // state buffer (hash)
2448
2449     u64 h[8];
2450
2451     h[0] = INITVAL;
2452     h[1] = INITVAL;
2453     h[2] = INITVAL;
2454     h[3] = INITVAL;
2455     h[4] = INITVAL;
2456     h[5] = INITVAL;
2457     h[6] = INITVAL;
2458     h[7] = INITVAL;
2459
2460     streebog_g (h, m, s_sbob_sl64);
2461
2462     u64 z[8];
2463
2464     z[0] = 0;
2465     z[1] = 0;
2466     z[2] = 0;
2467     z[3] = 0;
2468     z[4] = 0;
2469     z[5] = 0;
2470     z[6] = 0;
2471     z[7] = swap_workaround ((u64) (pw_len * 8));
2472
2473     streebog_g (h, z, s_sbob_sl64);
2474     streebog_g (h, m, s_sbob_sl64);
2475
2476     const u32 r0 = l32_from_64 (h[0]);
2477     const u32 r1 = h32_from_64 (h[0]);
2478     const u32 r2 = l32_from_64 (h[1]);
2479     const u32 r3 = h32_from_64 (h[1]);
2480
2481     #include VECT_COMPARE_S
2482   }
2483 }
2484
2485 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2486 {
2487   /**
2488    * base
2489    */
2490
2491   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2492   const u32 lid = threadIdx.x;
2493
2494   u32x w[16];
2495
2496   w[ 0] = pws[gid].i[ 0];
2497   w[ 1] = pws[gid].i[ 1];
2498   w[ 2] = pws[gid].i[ 2];
2499   w[ 3] = pws[gid].i[ 3];
2500   w[ 4] = 0;
2501   w[ 5] = 0;
2502   w[ 6] = 0;
2503   w[ 7] = 0;
2504   w[ 8] = 0;
2505   w[ 9] = 0;
2506   w[10] = 0;
2507   w[11] = 0;
2508   w[12] = 0;
2509   w[13] = 0;
2510   w[14] = 0;
2511   w[15] = 0;
2512
2513   const u32 pw_len = pws[gid].pw_len;
2514
2515   /**
2516    * shared lookup table
2517    */
2518
2519   __shared__ u64 s_sbob_sl64[8][256];
2520
2521   if (lid < 64)
2522   {
2523     const u32 lid4 = lid * 4;
2524
2525     s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2526     s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2527     s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2528     s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2529     s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2530     s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2531     s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2532     s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2533     s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2534     s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2535     s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2536     s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2537     s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2538     s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2539     s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2540     s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2541     s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2542     s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2543     s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2544     s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2545     s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2546     s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2547     s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2548     s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2549     s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2550     s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2551     s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2552     s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2553     s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2554     s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2555     s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2556     s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2557   }
2558
2559   __syncthreads ();
2560
2561   if (gid >= gid_max) return;
2562
2563   /**
2564    * main
2565    */
2566
2567   m11700m (s_sbob_sl64, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
2568 }
2569
2570 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2571 {
2572   /**
2573    * base
2574    */
2575
2576   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2577   const u32 lid = threadIdx.x;
2578
2579   u32x w[16];
2580
2581   w[ 0] = pws[gid].i[ 0];
2582   w[ 1] = pws[gid].i[ 1];
2583   w[ 2] = pws[gid].i[ 2];
2584   w[ 3] = pws[gid].i[ 3];
2585   w[ 4] = pws[gid].i[ 4];
2586   w[ 5] = pws[gid].i[ 5];
2587   w[ 6] = pws[gid].i[ 6];
2588   w[ 7] = pws[gid].i[ 7];
2589   w[ 8] = 0;
2590   w[ 9] = 0;
2591   w[10] = 0;
2592   w[11] = 0;
2593   w[12] = 0;
2594   w[13] = 0;
2595   w[14] = 0;
2596   w[15] = 0;
2597
2598   const u32 pw_len = pws[gid].pw_len;
2599
2600   /**
2601    * shared lookup table
2602    */
2603
2604   __shared__ u64 s_sbob_sl64[8][256];
2605
2606   if (lid < 64)
2607   {
2608     const u32 lid4 = lid * 4;
2609
2610     s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2611     s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2612     s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2613     s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2614     s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2615     s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2616     s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2617     s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2618     s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2619     s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2620     s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2621     s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2622     s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2623     s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2624     s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2625     s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2626     s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2627     s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2628     s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2629     s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2630     s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2631     s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2632     s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2633     s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2634     s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2635     s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2636     s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2637     s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2638     s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2639     s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2640     s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2641     s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2642   }
2643
2644   __syncthreads ();
2645
2646   if (gid >= gid_max) return;
2647
2648   /**
2649    * main
2650    */
2651
2652   m11700m (s_sbob_sl64, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
2653 }
2654
2655 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2656 {
2657   /**
2658    * base
2659    */
2660
2661   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2662   const u32 lid = threadIdx.x;
2663
2664   u32x w[16];
2665
2666   w[ 0] = pws[gid].i[ 0];
2667   w[ 1] = pws[gid].i[ 1];
2668   w[ 2] = pws[gid].i[ 2];
2669   w[ 3] = pws[gid].i[ 3];
2670   w[ 4] = pws[gid].i[ 4];
2671   w[ 5] = pws[gid].i[ 5];
2672   w[ 6] = pws[gid].i[ 6];
2673   w[ 7] = pws[gid].i[ 7];
2674   w[ 8] = pws[gid].i[ 8];
2675   w[ 9] = pws[gid].i[ 9];
2676   w[10] = pws[gid].i[10];
2677   w[11] = pws[gid].i[11];
2678   w[12] = pws[gid].i[12];
2679   w[13] = pws[gid].i[13];
2680   w[14] = pws[gid].i[14];
2681   w[15] = pws[gid].i[15];
2682
2683   const u32 pw_len = pws[gid].pw_len;
2684
2685   /**
2686    * shared lookup table
2687    */
2688
2689   __shared__ u64 s_sbob_sl64[8][256];
2690
2691   if (lid < 64)
2692   {
2693     const u32 lid4 = lid * 4;
2694
2695     s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2696     s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2697     s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2698     s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2699     s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2700     s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2701     s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2702     s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2703     s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2704     s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2705     s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2706     s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2707     s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2708     s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2709     s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2710     s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2711     s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2712     s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2713     s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2714     s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2715     s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2716     s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2717     s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2718     s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2719     s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2720     s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2721     s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2722     s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2723     s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2724     s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2725     s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2726     s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2727   }
2728
2729   __syncthreads ();
2730
2731   if (gid >= gid_max) return;
2732
2733   /**
2734    * main
2735    */
2736
2737   m11700m (s_sbob_sl64, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
2738 }
2739
2740 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2741 {
2742   /**
2743    * base
2744    */
2745
2746   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2747   const u32 lid = threadIdx.x;
2748
2749   u32x w[16];
2750
2751   w[ 0] = pws[gid].i[ 0];
2752   w[ 1] = pws[gid].i[ 1];
2753   w[ 2] = pws[gid].i[ 2];
2754   w[ 3] = pws[gid].i[ 3];
2755   w[ 4] = 0;
2756   w[ 5] = 0;
2757   w[ 6] = 0;
2758   w[ 7] = 0;
2759   w[ 8] = 0;
2760   w[ 9] = 0;
2761   w[10] = 0;
2762   w[11] = 0;
2763   w[12] = 0;
2764   w[13] = 0;
2765   w[14] = 0;
2766   w[15] = 0;
2767
2768   const u32 pw_len = pws[gid].pw_len;
2769
2770   /**
2771    * shared lookup table
2772    */
2773
2774   __shared__ u64 s_sbob_sl64[8][256];
2775
2776   if (lid < 64)
2777   {
2778     const u32 lid4 = lid * 4;
2779
2780     s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2781     s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2782     s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2783     s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2784     s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2785     s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2786     s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2787     s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2788     s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2789     s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2790     s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2791     s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2792     s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2793     s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2794     s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2795     s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2796     s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2797     s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2798     s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2799     s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2800     s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2801     s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2802     s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2803     s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2804     s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2805     s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2806     s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2807     s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2808     s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2809     s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2810     s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2811     s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2812   }
2813
2814   __syncthreads ();
2815
2816   if (gid >= gid_max) return;
2817
2818   /**
2819    * main
2820    */
2821
2822   m11700s (s_sbob_sl64, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
2823 }
2824
2825 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2826 {
2827   /**
2828    * base
2829    */
2830
2831   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2832   const u32 lid = threadIdx.x;
2833
2834   u32x w[16];
2835
2836   w[ 0] = pws[gid].i[ 0];
2837   w[ 1] = pws[gid].i[ 1];
2838   w[ 2] = pws[gid].i[ 2];
2839   w[ 3] = pws[gid].i[ 3];
2840   w[ 4] = pws[gid].i[ 4];
2841   w[ 5] = pws[gid].i[ 5];
2842   w[ 6] = pws[gid].i[ 6];
2843   w[ 7] = pws[gid].i[ 7];
2844   w[ 8] = 0;
2845   w[ 9] = 0;
2846   w[10] = 0;
2847   w[11] = 0;
2848   w[12] = 0;
2849   w[13] = 0;
2850   w[14] = 0;
2851   w[15] = 0;
2852
2853   const u32 pw_len = pws[gid].pw_len;
2854
2855   /**
2856    * shared lookup table
2857    */
2858
2859   __shared__ u64 s_sbob_sl64[8][256];
2860
2861   if (lid < 64)
2862   {
2863     const u32 lid4 = lid * 4;
2864
2865     s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2866     s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2867     s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2868     s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2869     s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2870     s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2871     s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2872     s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2873     s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2874     s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2875     s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2876     s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2877     s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2878     s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2879     s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2880     s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2881     s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2882     s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2883     s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2884     s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2885     s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2886     s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2887     s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2888     s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2889     s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2890     s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2891     s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2892     s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2893     s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2894     s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2895     s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2896     s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2897   }
2898
2899   __syncthreads ();
2900
2901   if (gid >= gid_max) return;
2902
2903   /**
2904    * main
2905    */
2906
2907   m11700s (s_sbob_sl64, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
2908 }
2909
2910 extern "C" __global__ void __launch_bounds__ (256, 1) m11700_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
2911 {
2912   /**
2913    * base
2914    */
2915
2916   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
2917   const u32 lid = threadIdx.x;
2918
2919   u32x w[16];
2920
2921   w[ 0] = pws[gid].i[ 0];
2922   w[ 1] = pws[gid].i[ 1];
2923   w[ 2] = pws[gid].i[ 2];
2924   w[ 3] = pws[gid].i[ 3];
2925   w[ 4] = pws[gid].i[ 4];
2926   w[ 5] = pws[gid].i[ 5];
2927   w[ 6] = pws[gid].i[ 6];
2928   w[ 7] = pws[gid].i[ 7];
2929   w[ 8] = pws[gid].i[ 8];
2930   w[ 9] = pws[gid].i[ 9];
2931   w[10] = pws[gid].i[10];
2932   w[11] = pws[gid].i[11];
2933   w[12] = pws[gid].i[12];
2934   w[13] = pws[gid].i[13];
2935   w[14] = pws[gid].i[14];
2936   w[15] = pws[gid].i[15];
2937
2938   const u32 pw_len = pws[gid].pw_len;
2939
2940   /**
2941    * shared lookup table
2942    */
2943
2944   __shared__ u64 s_sbob_sl64[8][256];
2945
2946   if (lid < 64)
2947   {
2948     const u32 lid4 = lid * 4;
2949
2950     s_sbob_sl64[0][lid4 + 0] = sbob_sl64[0][lid4 + 0];
2951     s_sbob_sl64[0][lid4 + 1] = sbob_sl64[0][lid4 + 1];
2952     s_sbob_sl64[0][lid4 + 2] = sbob_sl64[0][lid4 + 2];
2953     s_sbob_sl64[0][lid4 + 3] = sbob_sl64[0][lid4 + 3];
2954     s_sbob_sl64[1][lid4 + 0] = sbob_sl64[1][lid4 + 0];
2955     s_sbob_sl64[1][lid4 + 1] = sbob_sl64[1][lid4 + 1];
2956     s_sbob_sl64[1][lid4 + 2] = sbob_sl64[1][lid4 + 2];
2957     s_sbob_sl64[1][lid4 + 3] = sbob_sl64[1][lid4 + 3];
2958     s_sbob_sl64[2][lid4 + 0] = sbob_sl64[2][lid4 + 0];
2959     s_sbob_sl64[2][lid4 + 1] = sbob_sl64[2][lid4 + 1];
2960     s_sbob_sl64[2][lid4 + 2] = sbob_sl64[2][lid4 + 2];
2961     s_sbob_sl64[2][lid4 + 3] = sbob_sl64[2][lid4 + 3];
2962     s_sbob_sl64[3][lid4 + 0] = sbob_sl64[3][lid4 + 0];
2963     s_sbob_sl64[3][lid4 + 1] = sbob_sl64[3][lid4 + 1];
2964     s_sbob_sl64[3][lid4 + 2] = sbob_sl64[3][lid4 + 2];
2965     s_sbob_sl64[3][lid4 + 3] = sbob_sl64[3][lid4 + 3];
2966     s_sbob_sl64[4][lid4 + 0] = sbob_sl64[4][lid4 + 0];
2967     s_sbob_sl64[4][lid4 + 1] = sbob_sl64[4][lid4 + 1];
2968     s_sbob_sl64[4][lid4 + 2] = sbob_sl64[4][lid4 + 2];
2969     s_sbob_sl64[4][lid4 + 3] = sbob_sl64[4][lid4 + 3];
2970     s_sbob_sl64[5][lid4 + 0] = sbob_sl64[5][lid4 + 0];
2971     s_sbob_sl64[5][lid4 + 1] = sbob_sl64[5][lid4 + 1];
2972     s_sbob_sl64[5][lid4 + 2] = sbob_sl64[5][lid4 + 2];
2973     s_sbob_sl64[5][lid4 + 3] = sbob_sl64[5][lid4 + 3];
2974     s_sbob_sl64[6][lid4 + 0] = sbob_sl64[6][lid4 + 0];
2975     s_sbob_sl64[6][lid4 + 1] = sbob_sl64[6][lid4 + 1];
2976     s_sbob_sl64[6][lid4 + 2] = sbob_sl64[6][lid4 + 2];
2977     s_sbob_sl64[6][lid4 + 3] = sbob_sl64[6][lid4 + 3];
2978     s_sbob_sl64[7][lid4 + 0] = sbob_sl64[7][lid4 + 0];
2979     s_sbob_sl64[7][lid4 + 1] = sbob_sl64[7][lid4 + 1];
2980     s_sbob_sl64[7][lid4 + 2] = sbob_sl64[7][lid4 + 2];
2981     s_sbob_sl64[7][lid4 + 3] = sbob_sl64[7][lid4 + 3];
2982   }
2983
2984   __syncthreads ();
2985
2986   if (gid >= gid_max) return;
2987
2988   /**
2989    * main
2990    */
2991
2992   m11700s (s_sbob_sl64, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
2993 }