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