nv/m05800.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA1_
   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_M "check_multi_vect1_comp4.c"
  30 #endif
  31
  32 #ifdef  VECT_SIZE2
  33 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  34 #endif
  35
  36 typedef struct
  37 {
  38   u32 dec;
  39   u32 len;
  40
  41 } entry_t;
  42
  43 __device__ __constant__ entry_t pc[1024] =
  44 {
  45         0x00000030, 1, 0x00000031, 1, 0x00000032, 1, 0x00000033, 1, 0x00000034, 1, 0x00000035, 1, 0x00000036, 1, 0x00000037, 1,
  46         0x00000038, 1, 0x00000039, 1, 0x00003031, 2, 0x00003131, 2, 0x00003231, 2, 0x00003331, 2, 0x00003431, 2, 0x00003531, 2,
  47         0x00003631, 2, 0x00003731, 2, 0x00003831, 2, 0x00003931, 2, 0x00003032, 2, 0x00003132, 2, 0x00003232, 2, 0x00003332, 2,
  48         0x00003432, 2, 0x00003532, 2, 0x00003632, 2, 0x00003732, 2, 0x00003832, 2, 0x00003932, 2, 0x00003033, 2, 0x00003133, 2,
  49         0x00003233, 2, 0x00003333, 2, 0x00003433, 2, 0x00003533, 2, 0x00003633, 2, 0x00003733, 2, 0x00003833, 2, 0x00003933, 2,
  50         0x00003034, 2, 0x00003134, 2, 0x00003234, 2, 0x00003334, 2, 0x00003434, 2, 0x00003534, 2, 0x00003634, 2, 0x00003734, 2,
  51         0x00003834, 2, 0x00003934, 2, 0x00003035, 2, 0x00003135, 2, 0x00003235, 2, 0x00003335, 2, 0x00003435, 2, 0x00003535, 2,
  52         0x00003635, 2, 0x00003735, 2, 0x00003835, 2, 0x00003935, 2, 0x00003036, 2, 0x00003136, 2, 0x00003236, 2, 0x00003336, 2,
  53         0x00003436, 2, 0x00003536, 2, 0x00003636, 2, 0x00003736, 2, 0x00003836, 2, 0x00003936, 2, 0x00003037, 2, 0x00003137, 2,
  54         0x00003237, 2, 0x00003337, 2, 0x00003437, 2, 0x00003537, 2, 0x00003637, 2, 0x00003737, 2, 0x00003837, 2, 0x00003937, 2,
  55         0x00003038, 2, 0x00003138, 2, 0x00003238, 2, 0x00003338, 2, 0x00003438, 2, 0x00003538, 2, 0x00003638, 2, 0x00003738, 2,
  56         0x00003838, 2, 0x00003938, 2, 0x00003039, 2, 0x00003139, 2, 0x00003239, 2, 0x00003339, 2, 0x00003439, 2, 0x00003539, 2,
  57         0x00003639, 2, 0x00003739, 2, 0x00003839, 2, 0x00003939, 2, 0x00303031, 3, 0x00313031, 3, 0x00323031, 3, 0x00333031, 3,
  58         0x00343031, 3, 0x00353031, 3, 0x00363031, 3, 0x00373031, 3, 0x00383031, 3, 0x00393031, 3, 0x00303131, 3, 0x00313131, 3,
  59         0x00323131, 3, 0x00333131, 3, 0x00343131, 3, 0x00353131, 3, 0x00363131, 3, 0x00373131, 3, 0x00383131, 3, 0x00393131, 3,
  60         0x00303231, 3, 0x00313231, 3, 0x00323231, 3, 0x00333231, 3, 0x00343231, 3, 0x00353231, 3, 0x00363231, 3, 0x00373231, 3,
  61         0x00383231, 3, 0x00393231, 3, 0x00303331, 3, 0x00313331, 3, 0x00323331, 3, 0x00333331, 3, 0x00343331, 3, 0x00353331, 3,
  62         0x00363331, 3, 0x00373331, 3, 0x00383331, 3, 0x00393331, 3, 0x00303431, 3, 0x00313431, 3, 0x00323431, 3, 0x00333431, 3,
  63         0x00343431, 3, 0x00353431, 3, 0x00363431, 3, 0x00373431, 3, 0x00383431, 3, 0x00393431, 3, 0x00303531, 3, 0x00313531, 3,
  64         0x00323531, 3, 0x00333531, 3, 0x00343531, 3, 0x00353531, 3, 0x00363531, 3, 0x00373531, 3, 0x00383531, 3, 0x00393531, 3,
  65         0x00303631, 3, 0x00313631, 3, 0x00323631, 3, 0x00333631, 3, 0x00343631, 3, 0x00353631, 3, 0x00363631, 3, 0x00373631, 3,
  66         0x00383631, 3, 0x00393631, 3, 0x00303731, 3, 0x00313731, 3, 0x00323731, 3, 0x00333731, 3, 0x00343731, 3, 0x00353731, 3,
  67         0x00363731, 3, 0x00373731, 3, 0x00383731, 3, 0x00393731, 3, 0x00303831, 3, 0x00313831, 3, 0x00323831, 3, 0x00333831, 3,
  68         0x00343831, 3, 0x00353831, 3, 0x00363831, 3, 0x00373831, 3, 0x00383831, 3, 0x00393831, 3, 0x00303931, 3, 0x00313931, 3,
  69         0x00323931, 3, 0x00333931, 3, 0x00343931, 3, 0x00353931, 3, 0x00363931, 3, 0x00373931, 3, 0x00383931, 3, 0x00393931, 3,
  70         0x00303032, 3, 0x00313032, 3, 0x00323032, 3, 0x00333032, 3, 0x00343032, 3, 0x00353032, 3, 0x00363032, 3, 0x00373032, 3,
  71         0x00383032, 3, 0x00393032, 3, 0x00303132, 3, 0x00313132, 3, 0x00323132, 3, 0x00333132, 3, 0x00343132, 3, 0x00353132, 3,
  72         0x00363132, 3, 0x00373132, 3, 0x00383132, 3, 0x00393132, 3, 0x00303232, 3, 0x00313232, 3, 0x00323232, 3, 0x00333232, 3,
  73         0x00343232, 3, 0x00353232, 3, 0x00363232, 3, 0x00373232, 3, 0x00383232, 3, 0x00393232, 3, 0x00303332, 3, 0x00313332, 3,
  74         0x00323332, 3, 0x00333332, 3, 0x00343332, 3, 0x00353332, 3, 0x00363332, 3, 0x00373332, 3, 0x00383332, 3, 0x00393332, 3,
  75         0x00303432, 3, 0x00313432, 3, 0x00323432, 3, 0x00333432, 3, 0x00343432, 3, 0x00353432, 3, 0x00363432, 3, 0x00373432, 3,
  76         0x00383432, 3, 0x00393432, 3, 0x00303532, 3, 0x00313532, 3, 0x00323532, 3, 0x00333532, 3, 0x00343532, 3, 0x00353532, 3,
  77         0x00363532, 3, 0x00373532, 3, 0x00383532, 3, 0x00393532, 3, 0x00303632, 3, 0x00313632, 3, 0x00323632, 3, 0x00333632, 3,
  78         0x00343632, 3, 0x00353632, 3, 0x00363632, 3, 0x00373632, 3, 0x00383632, 3, 0x00393632, 3, 0x00303732, 3, 0x00313732, 3,
  79         0x00323732, 3, 0x00333732, 3, 0x00343732, 3, 0x00353732, 3, 0x00363732, 3, 0x00373732, 3, 0x00383732, 3, 0x00393732, 3,
  80         0x00303832, 3, 0x00313832, 3, 0x00323832, 3, 0x00333832, 3, 0x00343832, 3, 0x00353832, 3, 0x00363832, 3, 0x00373832, 3,
  81         0x00383832, 3, 0x00393832, 3, 0x00303932, 3, 0x00313932, 3, 0x00323932, 3, 0x00333932, 3, 0x00343932, 3, 0x00353932, 3,
  82         0x00363932, 3, 0x00373932, 3, 0x00383932, 3, 0x00393932, 3, 0x00303033, 3, 0x00313033, 3, 0x00323033, 3, 0x00333033, 3,
  83         0x00343033, 3, 0x00353033, 3, 0x00363033, 3, 0x00373033, 3, 0x00383033, 3, 0x00393033, 3, 0x00303133, 3, 0x00313133, 3,
  84         0x00323133, 3, 0x00333133, 3, 0x00343133, 3, 0x00353133, 3, 0x00363133, 3, 0x00373133, 3, 0x00383133, 3, 0x00393133, 3,
  85         0x00303233, 3, 0x00313233, 3, 0x00323233, 3, 0x00333233, 3, 0x00343233, 3, 0x00353233, 3, 0x00363233, 3, 0x00373233, 3,
  86         0x00383233, 3, 0x00393233, 3, 0x00303333, 3, 0x00313333, 3, 0x00323333, 3, 0x00333333, 3, 0x00343333, 3, 0x00353333, 3,
  87         0x00363333, 3, 0x00373333, 3, 0x00383333, 3, 0x00393333, 3, 0x00303433, 3, 0x00313433, 3, 0x00323433, 3, 0x00333433, 3,
  88         0x00343433, 3, 0x00353433, 3, 0x00363433, 3, 0x00373433, 3, 0x00383433, 3, 0x00393433, 3, 0x00303533, 3, 0x00313533, 3,
  89         0x00323533, 3, 0x00333533, 3, 0x00343533, 3, 0x00353533, 3, 0x00363533, 3, 0x00373533, 3, 0x00383533, 3, 0x00393533, 3,
  90         0x00303633, 3, 0x00313633, 3, 0x00323633, 3, 0x00333633, 3, 0x00343633, 3, 0x00353633, 3, 0x00363633, 3, 0x00373633, 3,
  91         0x00383633, 3, 0x00393633, 3, 0x00303733, 3, 0x00313733, 3, 0x00323733, 3, 0x00333733, 3, 0x00343733, 3, 0x00353733, 3,
  92         0x00363733, 3, 0x00373733, 3, 0x00383733, 3, 0x00393733, 3, 0x00303833, 3, 0x00313833, 3, 0x00323833, 3, 0x00333833, 3,
  93         0x00343833, 3, 0x00353833, 3, 0x00363833, 3, 0x00373833, 3, 0x00383833, 3, 0x00393833, 3, 0x00303933, 3, 0x00313933, 3,
  94         0x00323933, 3, 0x00333933, 3, 0x00343933, 3, 0x00353933, 3, 0x00363933, 3, 0x00373933, 3, 0x00383933, 3, 0x00393933, 3,
  95         0x00303034, 3, 0x00313034, 3, 0x00323034, 3, 0x00333034, 3, 0x00343034, 3, 0x00353034, 3, 0x00363034, 3, 0x00373034, 3,
  96         0x00383034, 3, 0x00393034, 3, 0x00303134, 3, 0x00313134, 3, 0x00323134, 3, 0x00333134, 3, 0x00343134, 3, 0x00353134, 3,
  97         0x00363134, 3, 0x00373134, 3, 0x00383134, 3, 0x00393134, 3, 0x00303234, 3, 0x00313234, 3, 0x00323234, 3, 0x00333234, 3,
  98         0x00343234, 3, 0x00353234, 3, 0x00363234, 3, 0x00373234, 3, 0x00383234, 3, 0x00393234, 3, 0x00303334, 3, 0x00313334, 3,
  99         0x00323334, 3, 0x00333334, 3, 0x00343334, 3, 0x00353334, 3, 0x00363334, 3, 0x00373334, 3, 0x00383334, 3, 0x00393334, 3,
 100         0x00303434, 3, 0x00313434, 3, 0x00323434, 3, 0x00333434, 3, 0x00343434, 3, 0x00353434, 3, 0x00363434, 3, 0x00373434, 3,
 101         0x00383434, 3, 0x00393434, 3, 0x00303534, 3, 0x00313534, 3, 0x00323534, 3, 0x00333534, 3, 0x00343534, 3, 0x00353534, 3,
 102         0x00363534, 3, 0x00373534, 3, 0x00383534, 3, 0x00393534, 3, 0x00303634, 3, 0x00313634, 3, 0x00323634, 3, 0x00333634, 3,
 103         0x00343634, 3, 0x00353634, 3, 0x00363634, 3, 0x00373634, 3, 0x00383634, 3, 0x00393634, 3, 0x00303734, 3, 0x00313734, 3,
 104         0x00323734, 3, 0x00333734, 3, 0x00343734, 3, 0x00353734, 3, 0x00363734, 3, 0x00373734, 3, 0x00383734, 3, 0x00393734, 3,
 105         0x00303834, 3, 0x00313834, 3, 0x00323834, 3, 0x00333834, 3, 0x00343834, 3, 0x00353834, 3, 0x00363834, 3, 0x00373834, 3,
 106         0x00383834, 3, 0x00393834, 3, 0x00303934, 3, 0x00313934, 3, 0x00323934, 3, 0x00333934, 3, 0x00343934, 3, 0x00353934, 3,
 107         0x00363934, 3, 0x00373934, 3, 0x00383934, 3, 0x00393934, 3, 0x00303035, 3, 0x00313035, 3, 0x00323035, 3, 0x00333035, 3,
 108         0x00343035, 3, 0x00353035, 3, 0x00363035, 3, 0x00373035, 3, 0x00383035, 3, 0x00393035, 3, 0x00303135, 3, 0x00313135, 3,
 109         0x00323135, 3, 0x00333135, 3, 0x00343135, 3, 0x00353135, 3, 0x00363135, 3, 0x00373135, 3, 0x00383135, 3, 0x00393135, 3,
 110         0x00303235, 3, 0x00313235, 3, 0x00323235, 3, 0x00333235, 3, 0x00343235, 3, 0x00353235, 3, 0x00363235, 3, 0x00373235, 3,
 111         0x00383235, 3, 0x00393235, 3, 0x00303335, 3, 0x00313335, 3, 0x00323335, 3, 0x00333335, 3, 0x00343335, 3, 0x00353335, 3,
 112         0x00363335, 3, 0x00373335, 3, 0x00383335, 3, 0x00393335, 3, 0x00303435, 3, 0x00313435, 3, 0x00323435, 3, 0x00333435, 3,
 113         0x00343435, 3, 0x00353435, 3, 0x00363435, 3, 0x00373435, 3, 0x00383435, 3, 0x00393435, 3, 0x00303535, 3, 0x00313535, 3,
 114         0x00323535, 3, 0x00333535, 3, 0x00343535, 3, 0x00353535, 3, 0x00363535, 3, 0x00373535, 3, 0x00383535, 3, 0x00393535, 3,
 115         0x00303635, 3, 0x00313635, 3, 0x00323635, 3, 0x00333635, 3, 0x00343635, 3, 0x00353635, 3, 0x00363635, 3, 0x00373635, 3,
 116         0x00383635, 3, 0x00393635, 3, 0x00303735, 3, 0x00313735, 3, 0x00323735, 3, 0x00333735, 3, 0x00343735, 3, 0x00353735, 3,
 117         0x00363735, 3, 0x00373735, 3, 0x00383735, 3, 0x00393735, 3, 0x00303835, 3, 0x00313835, 3, 0x00323835, 3, 0x00333835, 3,
 118         0x00343835, 3, 0x00353835, 3, 0x00363835, 3, 0x00373835, 3, 0x00383835, 3, 0x00393835, 3, 0x00303935, 3, 0x00313935, 3,
 119         0x00323935, 3, 0x00333935, 3, 0x00343935, 3, 0x00353935, 3, 0x00363935, 3, 0x00373935, 3, 0x00383935, 3, 0x00393935, 3,
 120         0x00303036, 3, 0x00313036, 3, 0x00323036, 3, 0x00333036, 3, 0x00343036, 3, 0x00353036, 3, 0x00363036, 3, 0x00373036, 3,
 121         0x00383036, 3, 0x00393036, 3, 0x00303136, 3, 0x00313136, 3, 0x00323136, 3, 0x00333136, 3, 0x00343136, 3, 0x00353136, 3,
 122         0x00363136, 3, 0x00373136, 3, 0x00383136, 3, 0x00393136, 3, 0x00303236, 3, 0x00313236, 3, 0x00323236, 3, 0x00333236, 3,
 123         0x00343236, 3, 0x00353236, 3, 0x00363236, 3, 0x00373236, 3, 0x00383236, 3, 0x00393236, 3, 0x00303336, 3, 0x00313336, 3,
 124         0x00323336, 3, 0x00333336, 3, 0x00343336, 3, 0x00353336, 3, 0x00363336, 3, 0x00373336, 3, 0x00383336, 3, 0x00393336, 3,
 125         0x00303436, 3, 0x00313436, 3, 0x00323436, 3, 0x00333436, 3, 0x00343436, 3, 0x00353436, 3, 0x00363436, 3, 0x00373436, 3,
 126         0x00383436, 3, 0x00393436, 3, 0x00303536, 3, 0x00313536, 3, 0x00323536, 3, 0x00333536, 3, 0x00343536, 3, 0x00353536, 3,
 127         0x00363536, 3, 0x00373536, 3, 0x00383536, 3, 0x00393536, 3, 0x00303636, 3, 0x00313636, 3, 0x00323636, 3, 0x00333636, 3,
 128         0x00343636, 3, 0x00353636, 3, 0x00363636, 3, 0x00373636, 3, 0x00383636, 3, 0x00393636, 3, 0x00303736, 3, 0x00313736, 3,
 129         0x00323736, 3, 0x00333736, 3, 0x00343736, 3, 0x00353736, 3, 0x00363736, 3, 0x00373736, 3, 0x00383736, 3, 0x00393736, 3,
 130         0x00303836, 3, 0x00313836, 3, 0x00323836, 3, 0x00333836, 3, 0x00343836, 3, 0x00353836, 3, 0x00363836, 3, 0x00373836, 3,
 131         0x00383836, 3, 0x00393836, 3, 0x00303936, 3, 0x00313936, 3, 0x00323936, 3, 0x00333936, 3, 0x00343936, 3, 0x00353936, 3,
 132         0x00363936, 3, 0x00373936, 3, 0x00383936, 3, 0x00393936, 3, 0x00303037, 3, 0x00313037, 3, 0x00323037, 3, 0x00333037, 3,
 133         0x00343037, 3, 0x00353037, 3, 0x00363037, 3, 0x00373037, 3, 0x00383037, 3, 0x00393037, 3, 0x00303137, 3, 0x00313137, 3,
 134         0x00323137, 3, 0x00333137, 3, 0x00343137, 3, 0x00353137, 3, 0x00363137, 3, 0x00373137, 3, 0x00383137, 3, 0x00393137, 3,
 135         0x00303237, 3, 0x00313237, 3, 0x00323237, 3, 0x00333237, 3, 0x00343237, 3, 0x00353237, 3, 0x00363237, 3, 0x00373237, 3,
 136         0x00383237, 3, 0x00393237, 3, 0x00303337, 3, 0x00313337, 3, 0x00323337, 3, 0x00333337, 3, 0x00343337, 3, 0x00353337, 3,
 137         0x00363337, 3, 0x00373337, 3, 0x00383337, 3, 0x00393337, 3, 0x00303437, 3, 0x00313437, 3, 0x00323437, 3, 0x00333437, 3,
 138         0x00343437, 3, 0x00353437, 3, 0x00363437, 3, 0x00373437, 3, 0x00383437, 3, 0x00393437, 3, 0x00303537, 3, 0x00313537, 3,
 139         0x00323537, 3, 0x00333537, 3, 0x00343537, 3, 0x00353537, 3, 0x00363537, 3, 0x00373537, 3, 0x00383537, 3, 0x00393537, 3,
 140         0x00303637, 3, 0x00313637, 3, 0x00323637, 3, 0x00333637, 3, 0x00343637, 3, 0x00353637, 3, 0x00363637, 3, 0x00373637, 3,
 141         0x00383637, 3, 0x00393637, 3, 0x00303737, 3, 0x00313737, 3, 0x00323737, 3, 0x00333737, 3, 0x00343737, 3, 0x00353737, 3,
 142         0x00363737, 3, 0x00373737, 3, 0x00383737, 3, 0x00393737, 3, 0x00303837, 3, 0x00313837, 3, 0x00323837, 3, 0x00333837, 3,
 143         0x00343837, 3, 0x00353837, 3, 0x00363837, 3, 0x00373837, 3, 0x00383837, 3, 0x00393837, 3, 0x00303937, 3, 0x00313937, 3,
 144         0x00323937, 3, 0x00333937, 3, 0x00343937, 3, 0x00353937, 3, 0x00363937, 3, 0x00373937, 3, 0x00383937, 3, 0x00393937, 3,
 145         0x00303038, 3, 0x00313038, 3, 0x00323038, 3, 0x00333038, 3, 0x00343038, 3, 0x00353038, 3, 0x00363038, 3, 0x00373038, 3,
 146         0x00383038, 3, 0x00393038, 3, 0x00303138, 3, 0x00313138, 3, 0x00323138, 3, 0x00333138, 3, 0x00343138, 3, 0x00353138, 3,
 147         0x00363138, 3, 0x00373138, 3, 0x00383138, 3, 0x00393138, 3, 0x00303238, 3, 0x00313238, 3, 0x00323238, 3, 0x00333238, 3,
 148         0x00343238, 3, 0x00353238, 3, 0x00363238, 3, 0x00373238, 3, 0x00383238, 3, 0x00393238, 3, 0x00303338, 3, 0x00313338, 3,
 149         0x00323338, 3, 0x00333338, 3, 0x00343338, 3, 0x00353338, 3, 0x00363338, 3, 0x00373338, 3, 0x00383338, 3, 0x00393338, 3,
 150         0x00303438, 3, 0x00313438, 3, 0x00323438, 3, 0x00333438, 3, 0x00343438, 3, 0x00353438, 3, 0x00363438, 3, 0x00373438, 3,
 151         0x00383438, 3, 0x00393438, 3, 0x00303538, 3, 0x00313538, 3, 0x00323538, 3, 0x00333538, 3, 0x00343538, 3, 0x00353538, 3,
 152         0x00363538, 3, 0x00373538, 3, 0x00383538, 3, 0x00393538, 3, 0x00303638, 3, 0x00313638, 3, 0x00323638, 3, 0x00333638, 3,
 153         0x00343638, 3, 0x00353638, 3, 0x00363638, 3, 0x00373638, 3, 0x00383638, 3, 0x00393638, 3, 0x00303738, 3, 0x00313738, 3,
 154         0x00323738, 3, 0x00333738, 3, 0x00343738, 3, 0x00353738, 3, 0x00363738, 3, 0x00373738, 3, 0x00383738, 3, 0x00393738, 3,
 155         0x00303838, 3, 0x00313838, 3, 0x00323838, 3, 0x00333838, 3, 0x00343838, 3, 0x00353838, 3, 0x00363838, 3, 0x00373838, 3,
 156         0x00383838, 3, 0x00393838, 3, 0x00303938, 3, 0x00313938, 3, 0x00323938, 3, 0x00333938, 3, 0x00343938, 3, 0x00353938, 3,
 157         0x00363938, 3, 0x00373938, 3, 0x00383938, 3, 0x00393938, 3, 0x00303039, 3, 0x00313039, 3, 0x00323039, 3, 0x00333039, 3,
 158         0x00343039, 3, 0x00353039, 3, 0x00363039, 3, 0x00373039, 3, 0x00383039, 3, 0x00393039, 3, 0x00303139, 3, 0x00313139, 3,
 159         0x00323139, 3, 0x00333139, 3, 0x00343139, 3, 0x00353139, 3, 0x00363139, 3, 0x00373139, 3, 0x00383139, 3, 0x00393139, 3,
 160         0x00303239, 3, 0x00313239, 3, 0x00323239, 3, 0x00333239, 3, 0x00343239, 3, 0x00353239, 3, 0x00363239, 3, 0x00373239, 3,
 161         0x00383239, 3, 0x00393239, 3, 0x00303339, 3, 0x00313339, 3, 0x00323339, 3, 0x00333339, 3, 0x00343339, 3, 0x00353339, 3,
 162         0x00363339, 3, 0x00373339, 3, 0x00383339, 3, 0x00393339, 3, 0x00303439, 3, 0x00313439, 3, 0x00323439, 3, 0x00333439, 3,
 163         0x00343439, 3, 0x00353439, 3, 0x00363439, 3, 0x00373439, 3, 0x00383439, 3, 0x00393439, 3, 0x00303539, 3, 0x00313539, 3,
 164         0x00323539, 3, 0x00333539, 3, 0x00343539, 3, 0x00353539, 3, 0x00363539, 3, 0x00373539, 3, 0x00383539, 3, 0x00393539, 3,
 165         0x00303639, 3, 0x00313639, 3, 0x00323639, 3, 0x00333639, 3, 0x00343639, 3, 0x00353639, 3, 0x00363639, 3, 0x00373639, 3,
 166         0x00383639, 3, 0x00393639, 3, 0x00303739, 3, 0x00313739, 3, 0x00323739, 3, 0x00333739, 3, 0x00343739, 3, 0x00353739, 3,
 167         0x00363739, 3, 0x00373739, 3, 0x00383739, 3, 0x00393739, 3, 0x00303839, 3, 0x00313839, 3, 0x00323839, 3, 0x00333839, 3,
 168         0x00343839, 3, 0x00353839, 3, 0x00363839, 3, 0x00373839, 3, 0x00383839, 3, 0x00393839, 3, 0x00303939, 3, 0x00313939, 3,
 169         0x00323939, 3, 0x00333939, 3, 0x00343939, 3, 0x00353939, 3, 0x00363939, 3, 0x00373939, 3, 0x00383939, 3, 0x00393939, 3,
 170         0x30303031, 4, 0x31303031, 4, 0x32303031, 4, 0x33303031, 4, 0x34303031, 4, 0x35303031, 4, 0x36303031, 4, 0x37303031, 4,
 171         0x38303031, 4, 0x39303031, 4, 0x30313031, 4, 0x31313031, 4, 0x32313031, 4, 0x33313031, 4, 0x34313031, 4, 0x35313031, 4,
 172         0x36313031, 4, 0x37313031, 4, 0x38313031, 4, 0x39313031, 4, 0x30323031, 4, 0x31323031, 4, 0x32323031, 4, 0x33323031, 4,
 173 };
 174
 175 __device__ static void append_word (u32x w0[4], u32x w1[4], const u32x append[4], const u32 offset)
 176 {
 177   switch (offset)
 178   {
 179     case 1:
 180       w0[0] = w0[0]           | append[0] <<  8;
 181       w0[1] = append[0] >> 24 | append[1] <<  8;
 182       w0[2] = append[1] >> 24 | append[2] <<  8;
 183       w0[3] = append[2] >> 24 | append[3] <<  8;
 184       break;
 185
 186     case 2:
 187       w0[0] = w0[0]           | append[0] << 16;
 188       w0[1] = append[0] >> 16 | append[1] << 16;
 189       w0[2] = append[1] >> 16 | append[2] << 16;
 190       w0[3] = append[2] >> 16 | append[3] << 16;
 191       break;
 192
 193     case 3:
 194       w0[0] = w0[0]           | append[0] << 24;
 195       w0[1] = append[0] >>  8 | append[1] << 24;
 196       w0[2] = append[1] >>  8 | append[2] << 24;
 197       w0[3] = append[2] >>  8 | append[3] << 24;
 198       break;
 199
 200     case 4:
 201       w0[1] = append[0];
 202       w0[2] = append[1];
 203       w0[3] = append[2];
 204       w1[0] = append[3];
 205       break;
 206   }
 207 }
 208
 209 __device__ static void append_salt (u32x w0[4], u32x w1[4], u32x w2[4], const u32 append[5], const u32 offset)
 210 {
 211   switch (offset)
 212   {
 213     case 2:
 214       w0[0] = w0[0]           | append[0] << 16;
 215       w0[1] = append[0] >> 16 | append[1] << 16;
 216       w0[2] = append[1] >> 16 | append[2] << 16;
 217       w0[3] = append[2] >> 16 | append[3] << 16;
 218       w1[0] = append[3] >> 16 | append[4] << 16;
 219       w1[1] = append[4] >> 16;
 220       break;
 221
 222     case 3:
 223       w0[0] = w0[0]           | append[0] << 24;
 224       w0[1] = append[0] >>  8 | append[1] << 24;
 225       w0[2] = append[1] >>  8 | append[2] << 24;
 226       w0[3] = append[2] >>  8 | append[3] << 24;
 227       w1[0] = append[3] >>  8 | append[4] << 24;
 228       w1[1] = append[4] >>  8;
 229       break;
 230
 231     case 4:
 232       w0[1] = append[0];
 233       w0[2] = append[1];
 234       w0[3] = append[2];
 235       w1[0] = append[3];
 236       w1[1] = append[4];
 237       break;
 238
 239     case 5:
 240       w0[1] = w0[1]           | append[0] <<  8;
 241       w0[2] = append[0] >> 24 | append[1] <<  8;
 242       w0[3] = append[1] >> 24 | append[2] <<  8;
 243       w1[0] = append[2] >> 24 | append[3] <<  8;
 244       w1[1] = append[3] >> 24 | append[4] <<  8;
 245       w1[2] = append[4] >> 24;
 246       break;
 247
 248     case 6:
 249       w0[1] = w0[1]           | append[0] << 16;
 250       w0[2] = append[0] >> 16 | append[1] << 16;
 251       w0[3] = append[1] >> 16 | append[2] << 16;
 252       w1[0] = append[2] >> 16 | append[3] << 16;
 253       w1[1] = append[3] >> 16 | append[4] << 16;
 254       w1[2] = append[4] >> 16;
 255       break;
 256
 257     case 7:
 258       w0[1] = w0[1]           | append[0] << 24;
 259       w0[2] = append[0] >>  8 | append[1] << 24;
 260       w0[3] = append[1] >>  8 | append[2] << 24;
 261       w1[0] = append[2] >>  8 | append[3] << 24;
 262       w1[1] = append[3] >>  8 | append[4] << 24;
 263       w1[2] = append[4] >>  8;
 264       break;
 265
 266     case 8:
 267       w0[2] = append[0];
 268       w0[3] = append[1];
 269       w1[0] = append[2];
 270       w1[1] = append[3];
 271       w1[2] = append[4];
 272       break;
 273
 274     case 9:
 275       w0[2] = w0[2]           | append[0] <<  8;
 276       w0[3] = append[0] >> 24 | append[1] <<  8;
 277       w1[0] = append[1] >> 24 | append[2] <<  8;
 278       w1[1] = append[2] >> 24 | append[3] <<  8;
 279       w1[2] = append[3] >> 24 | append[4] <<  8;
 280       w1[3] = append[4] >> 24;
 281       break;
 282
 283     case 10:
 284       w0[2] = w0[2]           | append[0] << 16;
 285       w0[3] = append[0] >> 16 | append[1] << 16;
 286       w1[0] = append[1] >> 16 | append[2] << 16;
 287       w1[1] = append[2] >> 16 | append[3] << 16;
 288       w1[2] = append[3] >> 16 | append[4] << 16;
 289       w1[3] = append[4] >> 16;
 290       break;
 291
 292     case 11:
 293       w0[2] = w0[2]           | append[0] << 24;
 294       w0[3] = append[0] >>  8 | append[1] << 24;
 295       w1[0] = append[1] >>  8 | append[2] << 24;
 296       w1[1] = append[2] >>  8 | append[3] << 24;
 297       w1[2] = append[3] >>  8 | append[4] << 24;
 298       w1[3] = append[4] >>  8;
 299       break;
 300
 301     case 12:
 302       w0[3] = append[0];
 303       w1[0] = append[1];
 304       w1[1] = append[2];
 305       w1[2] = append[3];
 306       w1[3] = append[4];
 307       break;
 308
 309     case 13:
 310       w0[3] = w0[3]           | append[0] <<  8;
 311       w1[0] = append[0] >> 24 | append[1] <<  8;
 312       w1[1] = append[1] >> 24 | append[2] <<  8;
 313       w1[2] = append[2] >> 24 | append[3] <<  8;
 314       w1[3] = append[3] >> 24 | append[4] <<  8;
 315       w2[0] = append[4] >> 24;
 316       break;
 317
 318     case 14:
 319       w0[3] = w0[3]           | append[0] << 16;
 320       w1[0] = append[0] >> 16 | append[1] << 16;
 321       w1[1] = append[1] >> 16 | append[2] << 16;
 322       w1[2] = append[2] >> 16 | append[3] << 16;
 323       w1[3] = append[3] >> 16 | append[4] << 16;
 324       w2[0] = append[4] >> 16;
 325       break;
 326
 327     case 15:
 328       w0[3] = w0[3]           | append[0] << 24;
 329       w1[0] = append[0] >>  8 | append[1] << 24;
 330       w1[1] = append[1] >>  8 | append[2] << 24;
 331       w1[2] = append[2] >>  8 | append[3] << 24;
 332       w1[3] = append[3] >>  8 | append[4] << 24;
 333       w2[0] = append[4] >>  8;
 334       break;
 335
 336     case 16:
 337       w1[0] = append[0];
 338       w1[1] = append[1];
 339       w1[2] = append[2];
 340       w1[3] = append[3];
 341       w2[0] = append[4];
 342       break;
 343
 344     case 17:
 345       w1[0] = w1[0]           | append[0] <<  8;
 346       w1[1] = append[0] >> 24 | append[1] <<  8;
 347       w1[2] = append[1] >> 24 | append[2] <<  8;
 348       w1[3] = append[2] >> 24 | append[3] <<  8;
 349       w2[0] = append[3] >> 24 | append[4] <<  8;
 350       w2[1] = append[4] >> 24;
 351       break;
 352
 353     case 18:
 354       w1[0] = w1[0]           | append[0] << 16;
 355       w1[1] = append[0] >> 16 | append[1] << 16;
 356       w1[2] = append[1] >> 16 | append[2] << 16;
 357       w1[3] = append[2] >> 16 | append[3] << 16;
 358       w2[0] = append[3] >> 16 | append[4] << 16;
 359       w2[1] = append[4] >> 16;
 360       break;
 361
 362     case 19:
 363       w1[0] = w1[0]           | append[0] << 24;
 364       w1[1] = append[0] >>  8 | append[1] << 24;
 365       w1[2] = append[1] >>  8 | append[2] << 24;
 366       w1[3] = append[2] >>  8 | append[3] << 24;
 367       w2[0] = append[3] >>  8 | append[4] << 24;
 368       w2[1] = append[4] >>  8;
 369       break;
 370   }
 371 }
 372
 373 __device__ static void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
 374 {
 375   u32x A = digest[0];
 376   u32x B = digest[1];
 377   u32x C = digest[2];
 378   u32x D = digest[3];
 379   u32x E = digest[4];
 380
 381   u32x w0_t = w0[0];
 382   u32x w1_t = w0[1];
 383   u32x w2_t = w0[2];
 384   u32x w3_t = w0[3];
 385   u32x w4_t = w1[0];
 386   u32x w5_t = w1[1];
 387   u32x w6_t = w1[2];
 388   u32x w7_t = w1[3];
 389   u32x w8_t = w2[0];
 390   u32x w9_t = w2[1];
 391   u32x wa_t = w2[2];
 392   u32x wb_t = w2[3];
 393   u32x wc_t = w3[0];
 394   u32x wd_t = w3[1];
 395   u32x we_t = w3[2];
 396   u32x wf_t = w3[3];
 397
 398   #undef K
 399   #define K SHA1C00
 400
 401   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
 402   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
 403   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
 404   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
 405   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
 406   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
 407   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
 408   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
 409   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
 410   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
 411   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
 412   SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
 413   SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
 414   SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
 415   SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
 416   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
 417   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
 418   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
 419   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
 420   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
 421
 422   #undef K
 423   #define K SHA1C01
 424
 425   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
 426   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
 427   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
 428   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
 429   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
 430   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
 431   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
 432   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
 433   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
 434   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
 435   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
 436   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
 437   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
 438   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
 439   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
 440   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
 441   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
 442   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
 443   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
 444   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
 445
 446   #undef K
 447   #define K SHA1C02
 448
 449   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
 450   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
 451   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
 452   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
 453   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
 454   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
 455   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
 456   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
 457   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
 458   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
 459   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
 460   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
 461   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
 462   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
 463   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
 464   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
 465   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
 466   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
 467   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
 468   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
 469
 470   #undef K
 471   #define K SHA1C03
 472
 473   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
 474   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
 475   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
 476   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
 477   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
 478   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
 479   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
 480   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
 481   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
 482   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
 483   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
 484   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
 485   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
 486   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
 487   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
 488   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
 489   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
 490   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
 491   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
 492   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
 493
 494   digest[0] += A;
 495   digest[1] += B;
 496   digest[2] += C;
 497   digest[3] += D;
 498   digest[4] += E;
 499 }
 500
 501 extern "C" __global__ void __launch_bounds__ (256, 1) m05800_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, androidpin_tmp_t *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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 502 {
 503   /**
 504    * base
 505    */
 506
 507   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 508
 509   if (gid >= gid_max) return;
 510
 511   u32x word_buf[4];
 512
 513   word_buf[0] = pws[gid].i[ 0];
 514   word_buf[1] = pws[gid].i[ 1];
 515   word_buf[2] = pws[gid].i[ 2];
 516   word_buf[3] = pws[gid].i[ 3];
 517
 518   const u32 pw_len = pws[gid].pw_len;
 519
 520   /**
 521    * salt
 522    */
 523
 524   u32 salt_len = salt_bufs[salt_pos].salt_len;
 525
 526   u32 salt_buf[5];
 527
 528   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 529   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 530   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 531   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 532   salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
 533
 534   /**
 535    * init
 536    */
 537
 538   const u32 pc_len = 1;
 539   const u32 pc_dec = 0x30;
 540
 541   u32x data0[4] = { 0, 0, 0, 0 };
 542   u32x data1[4] = { 0, 0, 0, 0 };
 543   u32x data2[4] = { 0, 0, 0, 0 };
 544
 545   data0[0] = pc_dec;
 546
 547   append_word (data0, data1, word_buf, pc_len);
 548
 549   append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
 550
 551   u32x w0[4];
 552   u32x w1[4];
 553   u32x w2[4];
 554   u32x w3[4];
 555
 556   w0[0] = swap_workaround (data0[0]);
 557   w0[1] = swap_workaround (data0[1]);
 558   w0[2] = swap_workaround (data0[2]);
 559   w0[3] = swap_workaround (data0[3]);
 560   w1[0] = swap_workaround (data1[0]);
 561   w1[1] = swap_workaround (data1[1]);
 562   w1[2] = swap_workaround (data1[2]);
 563   w1[3] = swap_workaround (data1[3]);
 564   w2[0] = swap_workaround (data2[0]);
 565   w2[1] = swap_workaround (data2[1]);
 566   w2[2] = 0;
 567   w2[3] = 0;
 568   w3[0] = 0;
 569   w3[1] = 0;
 570   w3[2] = 0;
 571   w3[3] = (pc_len + pw_len + salt_len) * 8;
 572
 573   u32x digest[5];
 574
 575   digest[0] = SHA1M_A;
 576   digest[1] = SHA1M_B;
 577   digest[2] = SHA1M_C;
 578   digest[3] = SHA1M_D;
 579   digest[4] = SHA1M_E;
 580
 581   sha1_transform (w0, w1, w2, w3, digest);
 582
 583   tmps[gid].digest_buf[0] = digest[0];
 584   tmps[gid].digest_buf[1] = digest[1];
 585   tmps[gid].digest_buf[2] = digest[2];
 586   tmps[gid].digest_buf[3] = digest[3];
 587   tmps[gid].digest_buf[4] = digest[4];
 588 }
 589
 590 extern "C" __global__ void __launch_bounds__ (256, 1) m05800_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, androidpin_tmp_t *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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 591 {
 592   /**
 593    * base
 594    */
 595
 596   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 597
 598   u32x word_buf[4];
 599
 600   word_buf[0] = pws[gid].i[ 0];
 601   word_buf[1] = pws[gid].i[ 1];
 602   word_buf[2] = pws[gid].i[ 2];
 603   word_buf[3] = pws[gid].i[ 3];
 604
 605   const u32 pw_len = pws[gid].pw_len;
 606
 607   u32x digest[5];
 608
 609   digest[0] = tmps[gid].digest_buf[0];
 610   digest[1] = tmps[gid].digest_buf[1];
 611   digest[2] = tmps[gid].digest_buf[2];
 612   digest[3] = tmps[gid].digest_buf[3];
 613   digest[4] = tmps[gid].digest_buf[4];
 614
 615   /**
 616    * cache precomputed conversion table in shared memory
 617    */
 618
 619   const u32 lid = threadIdx.x;
 620
 621   __shared__ entry_t s_pc[1024];
 622
 623   const u32 lid4 = lid * 4;
 624
 625   s_pc[lid4 + 0] = pc[lid4 + 0];
 626   s_pc[lid4 + 1] = pc[lid4 + 1];
 627   s_pc[lid4 + 2] = pc[lid4 + 2];
 628   s_pc[lid4 + 3] = pc[lid4 + 3];
 629
 630   __syncthreads ();
 631
 632   if (gid >= gid_max) return;
 633
 634   /**
 635    * salt
 636    */
 637
 638   u32 salt_len = salt_bufs[salt_pos].salt_len;
 639
 640   u32 salt_buf[5];
 641
 642   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 643   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 644   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 645   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 646   salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
 647
 648   /**
 649    * loop
 650    */
 651
 652   for (u32 i = 0, j = loop_pos + 1; i < loop_cnt; i++, j++)
 653   {
 654     const u32 pc_len = s_pc[j].len;
 655     const u32 pc_dec = s_pc[j].dec;
 656
 657     u32x data0[4] = { 0, 0, 0, 0 };
 658     u32x data1[4] = { 0, 0, 0, 0 };
 659     u32x data2[4] = { 0, 0, 0, 0 };
 660
 661     data0[0] = pc_dec;
 662
 663     append_word (data0, data1, word_buf, pc_len);
 664
 665     append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
 666
 667     u32x w0[4];
 668     u32x w1[4];
 669     u32x w2[4];
 670     u32x w3[4];
 671
 672     w0[0] = digest[0];
 673     w0[1] = digest[1];
 674     w0[2] = digest[2];
 675     w0[3] = digest[3];
 676     w1[0] = digest[4];
 677     w1[1] = swap_workaround (data0[0]);
 678     w1[2] = swap_workaround (data0[1]);
 679     w1[3] = swap_workaround (data0[2]);
 680     w2[0] = swap_workaround (data0[3]);
 681     w2[1] = swap_workaround (data1[0]);
 682     w2[2] = swap_workaround (data1[1]);
 683     w2[3] = swap_workaround (data1[2]);
 684     w3[0] = swap_workaround (data1[3]);
 685     w3[1] = swap_workaround (data2[0]);
 686     w3[2] = 0;
 687     w3[3] = (20 + pc_len + pw_len + salt_len) * 8;
 688
 689     digest[0] = SHA1M_A;
 690     digest[1] = SHA1M_B;
 691     digest[2] = SHA1M_C;
 692     digest[3] = SHA1M_D;
 693     digest[4] = SHA1M_E;
 694
 695     sha1_transform (w0, w1, w2, w3, digest);
 696   }
 697
 698   tmps[gid].digest_buf[0] = digest[0];
 699   tmps[gid].digest_buf[1] = digest[1];
 700   tmps[gid].digest_buf[2] = digest[2];
 701   tmps[gid].digest_buf[3] = digest[3];
 702   tmps[gid].digest_buf[4] = digest[4];
 703 }
 704
 705 extern "C" __global__ void __launch_bounds__ (256, 1) m05800_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, androidpin_tmp_t *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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 706 {
 707   /**
 708    * modifier
 709    */
 710
 711   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 712
 713   if (gid >= gid_max) return;
 714
 715   const u32 lid = threadIdx.x;
 716
 717   /**
 718    * digest
 719    */
 720
 721   const u32x r0 = tmps[gid].digest_buf[DGST_R0];
 722   const u32x r1 = tmps[gid].digest_buf[DGST_R1];
 723   const u32x r2 = tmps[gid].digest_buf[DGST_R2];
 724   const u32x r3 = tmps[gid].digest_buf[DGST_R3];
 725
 726   #define il_pos 0
 727
 728   #include VECT_COMPARE_M
 729 }