#define _SHA1_
-#include "include/constants.h"
-#include "include/kernel_vendor.h"
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
-#define DGST_R0 0
-#define DGST_R1 1
-#define DGST_R2 2
-#define DGST_R3 3
-
-#include "include/kernel_functions.c"
-#include "types_ocl.c"
-#include "common.c"
-
-#define COMPARE_S "check_single_comp4.c"
-#define COMPARE_M "check_multi_comp4.c"
+#define COMPARE_S "inc_comp_single.cl"
+#define COMPARE_M "inc_comp_multi.cl"
typedef struct
{
__constant entry_t pc[1024] =
{
- 0x00000030, 1, 0x00000031, 1, 0x00000032, 1, 0x00000033, 1, 0x00000034, 1, 0x00000035, 1, 0x00000036, 1, 0x00000037, 1,
- 0x00000038, 1, 0x00000039, 1, 0x00003031, 2, 0x00003131, 2, 0x00003231, 2, 0x00003331, 2, 0x00003431, 2, 0x00003531, 2,
- 0x00003631, 2, 0x00003731, 2, 0x00003831, 2, 0x00003931, 2, 0x00003032, 2, 0x00003132, 2, 0x00003232, 2, 0x00003332, 2,
- 0x00003432, 2, 0x00003532, 2, 0x00003632, 2, 0x00003732, 2, 0x00003832, 2, 0x00003932, 2, 0x00003033, 2, 0x00003133, 2,
- 0x00003233, 2, 0x00003333, 2, 0x00003433, 2, 0x00003533, 2, 0x00003633, 2, 0x00003733, 2, 0x00003833, 2, 0x00003933, 2,
- 0x00003034, 2, 0x00003134, 2, 0x00003234, 2, 0x00003334, 2, 0x00003434, 2, 0x00003534, 2, 0x00003634, 2, 0x00003734, 2,
- 0x00003834, 2, 0x00003934, 2, 0x00003035, 2, 0x00003135, 2, 0x00003235, 2, 0x00003335, 2, 0x00003435, 2, 0x00003535, 2,
- 0x00003635, 2, 0x00003735, 2, 0x00003835, 2, 0x00003935, 2, 0x00003036, 2, 0x00003136, 2, 0x00003236, 2, 0x00003336, 2,
- 0x00003436, 2, 0x00003536, 2, 0x00003636, 2, 0x00003736, 2, 0x00003836, 2, 0x00003936, 2, 0x00003037, 2, 0x00003137, 2,
- 0x00003237, 2, 0x00003337, 2, 0x00003437, 2, 0x00003537, 2, 0x00003637, 2, 0x00003737, 2, 0x00003837, 2, 0x00003937, 2,
- 0x00003038, 2, 0x00003138, 2, 0x00003238, 2, 0x00003338, 2, 0x00003438, 2, 0x00003538, 2, 0x00003638, 2, 0x00003738, 2,
- 0x00003838, 2, 0x00003938, 2, 0x00003039, 2, 0x00003139, 2, 0x00003239, 2, 0x00003339, 2, 0x00003439, 2, 0x00003539, 2,
- 0x00003639, 2, 0x00003739, 2, 0x00003839, 2, 0x00003939, 2, 0x00303031, 3, 0x00313031, 3, 0x00323031, 3, 0x00333031, 3,
- 0x00343031, 3, 0x00353031, 3, 0x00363031, 3, 0x00373031, 3, 0x00383031, 3, 0x00393031, 3, 0x00303131, 3, 0x00313131, 3,
- 0x00323131, 3, 0x00333131, 3, 0x00343131, 3, 0x00353131, 3, 0x00363131, 3, 0x00373131, 3, 0x00383131, 3, 0x00393131, 3,
- 0x00303231, 3, 0x00313231, 3, 0x00323231, 3, 0x00333231, 3, 0x00343231, 3, 0x00353231, 3, 0x00363231, 3, 0x00373231, 3,
- 0x00383231, 3, 0x00393231, 3, 0x00303331, 3, 0x00313331, 3, 0x00323331, 3, 0x00333331, 3, 0x00343331, 3, 0x00353331, 3,
- 0x00363331, 3, 0x00373331, 3, 0x00383331, 3, 0x00393331, 3, 0x00303431, 3, 0x00313431, 3, 0x00323431, 3, 0x00333431, 3,
- 0x00343431, 3, 0x00353431, 3, 0x00363431, 3, 0x00373431, 3, 0x00383431, 3, 0x00393431, 3, 0x00303531, 3, 0x00313531, 3,
- 0x00323531, 3, 0x00333531, 3, 0x00343531, 3, 0x00353531, 3, 0x00363531, 3, 0x00373531, 3, 0x00383531, 3, 0x00393531, 3,
- 0x00303631, 3, 0x00313631, 3, 0x00323631, 3, 0x00333631, 3, 0x00343631, 3, 0x00353631, 3, 0x00363631, 3, 0x00373631, 3,
- 0x00383631, 3, 0x00393631, 3, 0x00303731, 3, 0x00313731, 3, 0x00323731, 3, 0x00333731, 3, 0x00343731, 3, 0x00353731, 3,
- 0x00363731, 3, 0x00373731, 3, 0x00383731, 3, 0x00393731, 3, 0x00303831, 3, 0x00313831, 3, 0x00323831, 3, 0x00333831, 3,
- 0x00343831, 3, 0x00353831, 3, 0x00363831, 3, 0x00373831, 3, 0x00383831, 3, 0x00393831, 3, 0x00303931, 3, 0x00313931, 3,
- 0x00323931, 3, 0x00333931, 3, 0x00343931, 3, 0x00353931, 3, 0x00363931, 3, 0x00373931, 3, 0x00383931, 3, 0x00393931, 3,
- 0x00303032, 3, 0x00313032, 3, 0x00323032, 3, 0x00333032, 3, 0x00343032, 3, 0x00353032, 3, 0x00363032, 3, 0x00373032, 3,
- 0x00383032, 3, 0x00393032, 3, 0x00303132, 3, 0x00313132, 3, 0x00323132, 3, 0x00333132, 3, 0x00343132, 3, 0x00353132, 3,
- 0x00363132, 3, 0x00373132, 3, 0x00383132, 3, 0x00393132, 3, 0x00303232, 3, 0x00313232, 3, 0x00323232, 3, 0x00333232, 3,
- 0x00343232, 3, 0x00353232, 3, 0x00363232, 3, 0x00373232, 3, 0x00383232, 3, 0x00393232, 3, 0x00303332, 3, 0x00313332, 3,
- 0x00323332, 3, 0x00333332, 3, 0x00343332, 3, 0x00353332, 3, 0x00363332, 3, 0x00373332, 3, 0x00383332, 3, 0x00393332, 3,
- 0x00303432, 3, 0x00313432, 3, 0x00323432, 3, 0x00333432, 3, 0x00343432, 3, 0x00353432, 3, 0x00363432, 3, 0x00373432, 3,
- 0x00383432, 3, 0x00393432, 3, 0x00303532, 3, 0x00313532, 3, 0x00323532, 3, 0x00333532, 3, 0x00343532, 3, 0x00353532, 3,
- 0x00363532, 3, 0x00373532, 3, 0x00383532, 3, 0x00393532, 3, 0x00303632, 3, 0x00313632, 3, 0x00323632, 3, 0x00333632, 3,
- 0x00343632, 3, 0x00353632, 3, 0x00363632, 3, 0x00373632, 3, 0x00383632, 3, 0x00393632, 3, 0x00303732, 3, 0x00313732, 3,
- 0x00323732, 3, 0x00333732, 3, 0x00343732, 3, 0x00353732, 3, 0x00363732, 3, 0x00373732, 3, 0x00383732, 3, 0x00393732, 3,
- 0x00303832, 3, 0x00313832, 3, 0x00323832, 3, 0x00333832, 3, 0x00343832, 3, 0x00353832, 3, 0x00363832, 3, 0x00373832, 3,
- 0x00383832, 3, 0x00393832, 3, 0x00303932, 3, 0x00313932, 3, 0x00323932, 3, 0x00333932, 3, 0x00343932, 3, 0x00353932, 3,
- 0x00363932, 3, 0x00373932, 3, 0x00383932, 3, 0x00393932, 3, 0x00303033, 3, 0x00313033, 3, 0x00323033, 3, 0x00333033, 3,
- 0x00343033, 3, 0x00353033, 3, 0x00363033, 3, 0x00373033, 3, 0x00383033, 3, 0x00393033, 3, 0x00303133, 3, 0x00313133, 3,
- 0x00323133, 3, 0x00333133, 3, 0x00343133, 3, 0x00353133, 3, 0x00363133, 3, 0x00373133, 3, 0x00383133, 3, 0x00393133, 3,
- 0x00303233, 3, 0x00313233, 3, 0x00323233, 3, 0x00333233, 3, 0x00343233, 3, 0x00353233, 3, 0x00363233, 3, 0x00373233, 3,
- 0x00383233, 3, 0x00393233, 3, 0x00303333, 3, 0x00313333, 3, 0x00323333, 3, 0x00333333, 3, 0x00343333, 3, 0x00353333, 3,
- 0x00363333, 3, 0x00373333, 3, 0x00383333, 3, 0x00393333, 3, 0x00303433, 3, 0x00313433, 3, 0x00323433, 3, 0x00333433, 3,
- 0x00343433, 3, 0x00353433, 3, 0x00363433, 3, 0x00373433, 3, 0x00383433, 3, 0x00393433, 3, 0x00303533, 3, 0x00313533, 3,
- 0x00323533, 3, 0x00333533, 3, 0x00343533, 3, 0x00353533, 3, 0x00363533, 3, 0x00373533, 3, 0x00383533, 3, 0x00393533, 3,
- 0x00303633, 3, 0x00313633, 3, 0x00323633, 3, 0x00333633, 3, 0x00343633, 3, 0x00353633, 3, 0x00363633, 3, 0x00373633, 3,
- 0x00383633, 3, 0x00393633, 3, 0x00303733, 3, 0x00313733, 3, 0x00323733, 3, 0x00333733, 3, 0x00343733, 3, 0x00353733, 3,
- 0x00363733, 3, 0x00373733, 3, 0x00383733, 3, 0x00393733, 3, 0x00303833, 3, 0x00313833, 3, 0x00323833, 3, 0x00333833, 3,
- 0x00343833, 3, 0x00353833, 3, 0x00363833, 3, 0x00373833, 3, 0x00383833, 3, 0x00393833, 3, 0x00303933, 3, 0x00313933, 3,
- 0x00323933, 3, 0x00333933, 3, 0x00343933, 3, 0x00353933, 3, 0x00363933, 3, 0x00373933, 3, 0x00383933, 3, 0x00393933, 3,
- 0x00303034, 3, 0x00313034, 3, 0x00323034, 3, 0x00333034, 3, 0x00343034, 3, 0x00353034, 3, 0x00363034, 3, 0x00373034, 3,
- 0x00383034, 3, 0x00393034, 3, 0x00303134, 3, 0x00313134, 3, 0x00323134, 3, 0x00333134, 3, 0x00343134, 3, 0x00353134, 3,
- 0x00363134, 3, 0x00373134, 3, 0x00383134, 3, 0x00393134, 3, 0x00303234, 3, 0x00313234, 3, 0x00323234, 3, 0x00333234, 3,
- 0x00343234, 3, 0x00353234, 3, 0x00363234, 3, 0x00373234, 3, 0x00383234, 3, 0x00393234, 3, 0x00303334, 3, 0x00313334, 3,
- 0x00323334, 3, 0x00333334, 3, 0x00343334, 3, 0x00353334, 3, 0x00363334, 3, 0x00373334, 3, 0x00383334, 3, 0x00393334, 3,
- 0x00303434, 3, 0x00313434, 3, 0x00323434, 3, 0x00333434, 3, 0x00343434, 3, 0x00353434, 3, 0x00363434, 3, 0x00373434, 3,
- 0x00383434, 3, 0x00393434, 3, 0x00303534, 3, 0x00313534, 3, 0x00323534, 3, 0x00333534, 3, 0x00343534, 3, 0x00353534, 3,
- 0x00363534, 3, 0x00373534, 3, 0x00383534, 3, 0x00393534, 3, 0x00303634, 3, 0x00313634, 3, 0x00323634, 3, 0x00333634, 3,
- 0x00343634, 3, 0x00353634, 3, 0x00363634, 3, 0x00373634, 3, 0x00383634, 3, 0x00393634, 3, 0x00303734, 3, 0x00313734, 3,
- 0x00323734, 3, 0x00333734, 3, 0x00343734, 3, 0x00353734, 3, 0x00363734, 3, 0x00373734, 3, 0x00383734, 3, 0x00393734, 3,
- 0x00303834, 3, 0x00313834, 3, 0x00323834, 3, 0x00333834, 3, 0x00343834, 3, 0x00353834, 3, 0x00363834, 3, 0x00373834, 3,
- 0x00383834, 3, 0x00393834, 3, 0x00303934, 3, 0x00313934, 3, 0x00323934, 3, 0x00333934, 3, 0x00343934, 3, 0x00353934, 3,
- 0x00363934, 3, 0x00373934, 3, 0x00383934, 3, 0x00393934, 3, 0x00303035, 3, 0x00313035, 3, 0x00323035, 3, 0x00333035, 3,
- 0x00343035, 3, 0x00353035, 3, 0x00363035, 3, 0x00373035, 3, 0x00383035, 3, 0x00393035, 3, 0x00303135, 3, 0x00313135, 3,
- 0x00323135, 3, 0x00333135, 3, 0x00343135, 3, 0x00353135, 3, 0x00363135, 3, 0x00373135, 3, 0x00383135, 3, 0x00393135, 3,
- 0x00303235, 3, 0x00313235, 3, 0x00323235, 3, 0x00333235, 3, 0x00343235, 3, 0x00353235, 3, 0x00363235, 3, 0x00373235, 3,
- 0x00383235, 3, 0x00393235, 3, 0x00303335, 3, 0x00313335, 3, 0x00323335, 3, 0x00333335, 3, 0x00343335, 3, 0x00353335, 3,
- 0x00363335, 3, 0x00373335, 3, 0x00383335, 3, 0x00393335, 3, 0x00303435, 3, 0x00313435, 3, 0x00323435, 3, 0x00333435, 3,
- 0x00343435, 3, 0x00353435, 3, 0x00363435, 3, 0x00373435, 3, 0x00383435, 3, 0x00393435, 3, 0x00303535, 3, 0x00313535, 3,
- 0x00323535, 3, 0x00333535, 3, 0x00343535, 3, 0x00353535, 3, 0x00363535, 3, 0x00373535, 3, 0x00383535, 3, 0x00393535, 3,
- 0x00303635, 3, 0x00313635, 3, 0x00323635, 3, 0x00333635, 3, 0x00343635, 3, 0x00353635, 3, 0x00363635, 3, 0x00373635, 3,
- 0x00383635, 3, 0x00393635, 3, 0x00303735, 3, 0x00313735, 3, 0x00323735, 3, 0x00333735, 3, 0x00343735, 3, 0x00353735, 3,
- 0x00363735, 3, 0x00373735, 3, 0x00383735, 3, 0x00393735, 3, 0x00303835, 3, 0x00313835, 3, 0x00323835, 3, 0x00333835, 3,
- 0x00343835, 3, 0x00353835, 3, 0x00363835, 3, 0x00373835, 3, 0x00383835, 3, 0x00393835, 3, 0x00303935, 3, 0x00313935, 3,
- 0x00323935, 3, 0x00333935, 3, 0x00343935, 3, 0x00353935, 3, 0x00363935, 3, 0x00373935, 3, 0x00383935, 3, 0x00393935, 3,
- 0x00303036, 3, 0x00313036, 3, 0x00323036, 3, 0x00333036, 3, 0x00343036, 3, 0x00353036, 3, 0x00363036, 3, 0x00373036, 3,
- 0x00383036, 3, 0x00393036, 3, 0x00303136, 3, 0x00313136, 3, 0x00323136, 3, 0x00333136, 3, 0x00343136, 3, 0x00353136, 3,
- 0x00363136, 3, 0x00373136, 3, 0x00383136, 3, 0x00393136, 3, 0x00303236, 3, 0x00313236, 3, 0x00323236, 3, 0x00333236, 3,
- 0x00343236, 3, 0x00353236, 3, 0x00363236, 3, 0x00373236, 3, 0x00383236, 3, 0x00393236, 3, 0x00303336, 3, 0x00313336, 3,
- 0x00323336, 3, 0x00333336, 3, 0x00343336, 3, 0x00353336, 3, 0x00363336, 3, 0x00373336, 3, 0x00383336, 3, 0x00393336, 3,
- 0x00303436, 3, 0x00313436, 3, 0x00323436, 3, 0x00333436, 3, 0x00343436, 3, 0x00353436, 3, 0x00363436, 3, 0x00373436, 3,
- 0x00383436, 3, 0x00393436, 3, 0x00303536, 3, 0x00313536, 3, 0x00323536, 3, 0x00333536, 3, 0x00343536, 3, 0x00353536, 3,
- 0x00363536, 3, 0x00373536, 3, 0x00383536, 3, 0x00393536, 3, 0x00303636, 3, 0x00313636, 3, 0x00323636, 3, 0x00333636, 3,
- 0x00343636, 3, 0x00353636, 3, 0x00363636, 3, 0x00373636, 3, 0x00383636, 3, 0x00393636, 3, 0x00303736, 3, 0x00313736, 3,
- 0x00323736, 3, 0x00333736, 3, 0x00343736, 3, 0x00353736, 3, 0x00363736, 3, 0x00373736, 3, 0x00383736, 3, 0x00393736, 3,
- 0x00303836, 3, 0x00313836, 3, 0x00323836, 3, 0x00333836, 3, 0x00343836, 3, 0x00353836, 3, 0x00363836, 3, 0x00373836, 3,
- 0x00383836, 3, 0x00393836, 3, 0x00303936, 3, 0x00313936, 3, 0x00323936, 3, 0x00333936, 3, 0x00343936, 3, 0x00353936, 3,
- 0x00363936, 3, 0x00373936, 3, 0x00383936, 3, 0x00393936, 3, 0x00303037, 3, 0x00313037, 3, 0x00323037, 3, 0x00333037, 3,
- 0x00343037, 3, 0x00353037, 3, 0x00363037, 3, 0x00373037, 3, 0x00383037, 3, 0x00393037, 3, 0x00303137, 3, 0x00313137, 3,
- 0x00323137, 3, 0x00333137, 3, 0x00343137, 3, 0x00353137, 3, 0x00363137, 3, 0x00373137, 3, 0x00383137, 3, 0x00393137, 3,
- 0x00303237, 3, 0x00313237, 3, 0x00323237, 3, 0x00333237, 3, 0x00343237, 3, 0x00353237, 3, 0x00363237, 3, 0x00373237, 3,
- 0x00383237, 3, 0x00393237, 3, 0x00303337, 3, 0x00313337, 3, 0x00323337, 3, 0x00333337, 3, 0x00343337, 3, 0x00353337, 3,
- 0x00363337, 3, 0x00373337, 3, 0x00383337, 3, 0x00393337, 3, 0x00303437, 3, 0x00313437, 3, 0x00323437, 3, 0x00333437, 3,
- 0x00343437, 3, 0x00353437, 3, 0x00363437, 3, 0x00373437, 3, 0x00383437, 3, 0x00393437, 3, 0x00303537, 3, 0x00313537, 3,
- 0x00323537, 3, 0x00333537, 3, 0x00343537, 3, 0x00353537, 3, 0x00363537, 3, 0x00373537, 3, 0x00383537, 3, 0x00393537, 3,
- 0x00303637, 3, 0x00313637, 3, 0x00323637, 3, 0x00333637, 3, 0x00343637, 3, 0x00353637, 3, 0x00363637, 3, 0x00373637, 3,
- 0x00383637, 3, 0x00393637, 3, 0x00303737, 3, 0x00313737, 3, 0x00323737, 3, 0x00333737, 3, 0x00343737, 3, 0x00353737, 3,
- 0x00363737, 3, 0x00373737, 3, 0x00383737, 3, 0x00393737, 3, 0x00303837, 3, 0x00313837, 3, 0x00323837, 3, 0x00333837, 3,
- 0x00343837, 3, 0x00353837, 3, 0x00363837, 3, 0x00373837, 3, 0x00383837, 3, 0x00393837, 3, 0x00303937, 3, 0x00313937, 3,
- 0x00323937, 3, 0x00333937, 3, 0x00343937, 3, 0x00353937, 3, 0x00363937, 3, 0x00373937, 3, 0x00383937, 3, 0x00393937, 3,
- 0x00303038, 3, 0x00313038, 3, 0x00323038, 3, 0x00333038, 3, 0x00343038, 3, 0x00353038, 3, 0x00363038, 3, 0x00373038, 3,
- 0x00383038, 3, 0x00393038, 3, 0x00303138, 3, 0x00313138, 3, 0x00323138, 3, 0x00333138, 3, 0x00343138, 3, 0x00353138, 3,
- 0x00363138, 3, 0x00373138, 3, 0x00383138, 3, 0x00393138, 3, 0x00303238, 3, 0x00313238, 3, 0x00323238, 3, 0x00333238, 3,
- 0x00343238, 3, 0x00353238, 3, 0x00363238, 3, 0x00373238, 3, 0x00383238, 3, 0x00393238, 3, 0x00303338, 3, 0x00313338, 3,
- 0x00323338, 3, 0x00333338, 3, 0x00343338, 3, 0x00353338, 3, 0x00363338, 3, 0x00373338, 3, 0x00383338, 3, 0x00393338, 3,
- 0x00303438, 3, 0x00313438, 3, 0x00323438, 3, 0x00333438, 3, 0x00343438, 3, 0x00353438, 3, 0x00363438, 3, 0x00373438, 3,
- 0x00383438, 3, 0x00393438, 3, 0x00303538, 3, 0x00313538, 3, 0x00323538, 3, 0x00333538, 3, 0x00343538, 3, 0x00353538, 3,
- 0x00363538, 3, 0x00373538, 3, 0x00383538, 3, 0x00393538, 3, 0x00303638, 3, 0x00313638, 3, 0x00323638, 3, 0x00333638, 3,
- 0x00343638, 3, 0x00353638, 3, 0x00363638, 3, 0x00373638, 3, 0x00383638, 3, 0x00393638, 3, 0x00303738, 3, 0x00313738, 3,
- 0x00323738, 3, 0x00333738, 3, 0x00343738, 3, 0x00353738, 3, 0x00363738, 3, 0x00373738, 3, 0x00383738, 3, 0x00393738, 3,
- 0x00303838, 3, 0x00313838, 3, 0x00323838, 3, 0x00333838, 3, 0x00343838, 3, 0x00353838, 3, 0x00363838, 3, 0x00373838, 3,
- 0x00383838, 3, 0x00393838, 3, 0x00303938, 3, 0x00313938, 3, 0x00323938, 3, 0x00333938, 3, 0x00343938, 3, 0x00353938, 3,
- 0x00363938, 3, 0x00373938, 3, 0x00383938, 3, 0x00393938, 3, 0x00303039, 3, 0x00313039, 3, 0x00323039, 3, 0x00333039, 3,
- 0x00343039, 3, 0x00353039, 3, 0x00363039, 3, 0x00373039, 3, 0x00383039, 3, 0x00393039, 3, 0x00303139, 3, 0x00313139, 3,
- 0x00323139, 3, 0x00333139, 3, 0x00343139, 3, 0x00353139, 3, 0x00363139, 3, 0x00373139, 3, 0x00383139, 3, 0x00393139, 3,
- 0x00303239, 3, 0x00313239, 3, 0x00323239, 3, 0x00333239, 3, 0x00343239, 3, 0x00353239, 3, 0x00363239, 3, 0x00373239, 3,
- 0x00383239, 3, 0x00393239, 3, 0x00303339, 3, 0x00313339, 3, 0x00323339, 3, 0x00333339, 3, 0x00343339, 3, 0x00353339, 3,
- 0x00363339, 3, 0x00373339, 3, 0x00383339, 3, 0x00393339, 3, 0x00303439, 3, 0x00313439, 3, 0x00323439, 3, 0x00333439, 3,
- 0x00343439, 3, 0x00353439, 3, 0x00363439, 3, 0x00373439, 3, 0x00383439, 3, 0x00393439, 3, 0x00303539, 3, 0x00313539, 3,
- 0x00323539, 3, 0x00333539, 3, 0x00343539, 3, 0x00353539, 3, 0x00363539, 3, 0x00373539, 3, 0x00383539, 3, 0x00393539, 3,
- 0x00303639, 3, 0x00313639, 3, 0x00323639, 3, 0x00333639, 3, 0x00343639, 3, 0x00353639, 3, 0x00363639, 3, 0x00373639, 3,
- 0x00383639, 3, 0x00393639, 3, 0x00303739, 3, 0x00313739, 3, 0x00323739, 3, 0x00333739, 3, 0x00343739, 3, 0x00353739, 3,
- 0x00363739, 3, 0x00373739, 3, 0x00383739, 3, 0x00393739, 3, 0x00303839, 3, 0x00313839, 3, 0x00323839, 3, 0x00333839, 3,
- 0x00343839, 3, 0x00353839, 3, 0x00363839, 3, 0x00373839, 3, 0x00383839, 3, 0x00393839, 3, 0x00303939, 3, 0x00313939, 3,
- 0x00323939, 3, 0x00333939, 3, 0x00343939, 3, 0x00353939, 3, 0x00363939, 3, 0x00373939, 3, 0x00383939, 3, 0x00393939, 3,
- 0x30303031, 4, 0x31303031, 4, 0x32303031, 4, 0x33303031, 4, 0x34303031, 4, 0x35303031, 4, 0x36303031, 4, 0x37303031, 4,
- 0x38303031, 4, 0x39303031, 4, 0x30313031, 4, 0x31313031, 4, 0x32313031, 4, 0x33313031, 4, 0x34313031, 4, 0x35313031, 4,
- 0x36313031, 4, 0x37313031, 4, 0x38313031, 4, 0x39313031, 4, 0x30323031, 4, 0x31323031, 4, 0x32323031, 4, 0x33323031, 4,
+ { 0x00000030, 1 },
+ { 0x00000031, 1 },
+ { 0x00000032, 1 },
+ { 0x00000033, 1 },
+ { 0x00000034, 1 },
+ { 0x00000035, 1 },
+ { 0x00000036, 1 },
+ { 0x00000037, 1 },
+ { 0x00000038, 1 },
+ { 0x00000039, 1 },
+ { 0x00003031, 2 },
+ { 0x00003131, 2 },
+ { 0x00003231, 2 },
+ { 0x00003331, 2 },
+ { 0x00003431, 2 },
+ { 0x00003531, 2 },
+ { 0x00003631, 2 },
+ { 0x00003731, 2 },
+ { 0x00003831, 2 },
+ { 0x00003931, 2 },
+ { 0x00003032, 2 },
+ { 0x00003132, 2 },
+ { 0x00003232, 2 },
+ { 0x00003332, 2 },
+ { 0x00003432, 2 },
+ { 0x00003532, 2 },
+ { 0x00003632, 2 },
+ { 0x00003732, 2 },
+ { 0x00003832, 2 },
+ { 0x00003932, 2 },
+ { 0x00003033, 2 },
+ { 0x00003133, 2 },
+ { 0x00003233, 2 },
+ { 0x00003333, 2 },
+ { 0x00003433, 2 },
+ { 0x00003533, 2 },
+ { 0x00003633, 2 },
+ { 0x00003733, 2 },
+ { 0x00003833, 2 },
+ { 0x00003933, 2 },
+ { 0x00003034, 2 },
+ { 0x00003134, 2 },
+ { 0x00003234, 2 },
+ { 0x00003334, 2 },
+ { 0x00003434, 2 },
+ { 0x00003534, 2 },
+ { 0x00003634, 2 },
+ { 0x00003734, 2 },
+ { 0x00003834, 2 },
+ { 0x00003934, 2 },
+ { 0x00003035, 2 },
+ { 0x00003135, 2 },
+ { 0x00003235, 2 },
+ { 0x00003335, 2 },
+ { 0x00003435, 2 },
+ { 0x00003535, 2 },
+ { 0x00003635, 2 },
+ { 0x00003735, 2 },
+ { 0x00003835, 2 },
+ { 0x00003935, 2 },
+ { 0x00003036, 2 },
+ { 0x00003136, 2 },
+ { 0x00003236, 2 },
+ { 0x00003336, 2 },
+ { 0x00003436, 2 },
+ { 0x00003536, 2 },
+ { 0x00003636, 2 },
+ { 0x00003736, 2 },
+ { 0x00003836, 2 },
+ { 0x00003936, 2 },
+ { 0x00003037, 2 },
+ { 0x00003137, 2 },
+ { 0x00003237, 2 },
+ { 0x00003337, 2 },
+ { 0x00003437, 2 },
+ { 0x00003537, 2 },
+ { 0x00003637, 2 },
+ { 0x00003737, 2 },
+ { 0x00003837, 2 },
+ { 0x00003937, 2 },
+ { 0x00003038, 2 },
+ { 0x00003138, 2 },
+ { 0x00003238, 2 },
+ { 0x00003338, 2 },
+ { 0x00003438, 2 },
+ { 0x00003538, 2 },
+ { 0x00003638, 2 },
+ { 0x00003738, 2 },
+ { 0x00003838, 2 },
+ { 0x00003938, 2 },
+ { 0x00003039, 2 },
+ { 0x00003139, 2 },
+ { 0x00003239, 2 },
+ { 0x00003339, 2 },
+ { 0x00003439, 2 },
+ { 0x00003539, 2 },
+ { 0x00003639, 2 },
+ { 0x00003739, 2 },
+ { 0x00003839, 2 },
+ { 0x00003939, 2 },
+ { 0x00303031, 3 },
+ { 0x00313031, 3 },
+ { 0x00323031, 3 },
+ { 0x00333031, 3 },
+ { 0x00343031, 3 },
+ { 0x00353031, 3 },
+ { 0x00363031, 3 },
+ { 0x00373031, 3 },
+ { 0x00383031, 3 },
+ { 0x00393031, 3 },
+ { 0x00303131, 3 },
+ { 0x00313131, 3 },
+ { 0x00323131, 3 },
+ { 0x00333131, 3 },
+ { 0x00343131, 3 },
+ { 0x00353131, 3 },
+ { 0x00363131, 3 },
+ { 0x00373131, 3 },
+ { 0x00383131, 3 },
+ { 0x00393131, 3 },
+ { 0x00303231, 3 },
+ { 0x00313231, 3 },
+ { 0x00323231, 3 },
+ { 0x00333231, 3 },
+ { 0x00343231, 3 },
+ { 0x00353231, 3 },
+ { 0x00363231, 3 },
+ { 0x00373231, 3 },
+ { 0x00383231, 3 },
+ { 0x00393231, 3 },
+ { 0x00303331, 3 },
+ { 0x00313331, 3 },
+ { 0x00323331, 3 },
+ { 0x00333331, 3 },
+ { 0x00343331, 3 },
+ { 0x00353331, 3 },
+ { 0x00363331, 3 },
+ { 0x00373331, 3 },
+ { 0x00383331, 3 },
+ { 0x00393331, 3 },
+ { 0x00303431, 3 },
+ { 0x00313431, 3 },
+ { 0x00323431, 3 },
+ { 0x00333431, 3 },
+ { 0x00343431, 3 },
+ { 0x00353431, 3 },
+ { 0x00363431, 3 },
+ { 0x00373431, 3 },
+ { 0x00383431, 3 },
+ { 0x00393431, 3 },
+ { 0x00303531, 3 },
+ { 0x00313531, 3 },
+ { 0x00323531, 3 },
+ { 0x00333531, 3 },
+ { 0x00343531, 3 },
+ { 0x00353531, 3 },
+ { 0x00363531, 3 },
+ { 0x00373531, 3 },
+ { 0x00383531, 3 },
+ { 0x00393531, 3 },
+ { 0x00303631, 3 },
+ { 0x00313631, 3 },
+ { 0x00323631, 3 },
+ { 0x00333631, 3 },
+ { 0x00343631, 3 },
+ { 0x00353631, 3 },
+ { 0x00363631, 3 },
+ { 0x00373631, 3 },
+ { 0x00383631, 3 },
+ { 0x00393631, 3 },
+ { 0x00303731, 3 },
+ { 0x00313731, 3 },
+ { 0x00323731, 3 },
+ { 0x00333731, 3 },
+ { 0x00343731, 3 },
+ { 0x00353731, 3 },
+ { 0x00363731, 3 },
+ { 0x00373731, 3 },
+ { 0x00383731, 3 },
+ { 0x00393731, 3 },
+ { 0x00303831, 3 },
+ { 0x00313831, 3 },
+ { 0x00323831, 3 },
+ { 0x00333831, 3 },
+ { 0x00343831, 3 },
+ { 0x00353831, 3 },
+ { 0x00363831, 3 },
+ { 0x00373831, 3 },
+ { 0x00383831, 3 },
+ { 0x00393831, 3 },
+ { 0x00303931, 3 },
+ { 0x00313931, 3 },
+ { 0x00323931, 3 },
+ { 0x00333931, 3 },
+ { 0x00343931, 3 },
+ { 0x00353931, 3 },
+ { 0x00363931, 3 },
+ { 0x00373931, 3 },
+ { 0x00383931, 3 },
+ { 0x00393931, 3 },
+ { 0x00303032, 3 },
+ { 0x00313032, 3 },
+ { 0x00323032, 3 },
+ { 0x00333032, 3 },
+ { 0x00343032, 3 },
+ { 0x00353032, 3 },
+ { 0x00363032, 3 },
+ { 0x00373032, 3 },
+ { 0x00383032, 3 },
+ { 0x00393032, 3 },
+ { 0x00303132, 3 },
+ { 0x00313132, 3 },
+ { 0x00323132, 3 },
+ { 0x00333132, 3 },
+ { 0x00343132, 3 },
+ { 0x00353132, 3 },
+ { 0x00363132, 3 },
+ { 0x00373132, 3 },
+ { 0x00383132, 3 },
+ { 0x00393132, 3 },
+ { 0x00303232, 3 },
+ { 0x00313232, 3 },
+ { 0x00323232, 3 },
+ { 0x00333232, 3 },
+ { 0x00343232, 3 },
+ { 0x00353232, 3 },
+ { 0x00363232, 3 },
+ { 0x00373232, 3 },
+ { 0x00383232, 3 },
+ { 0x00393232, 3 },
+ { 0x00303332, 3 },
+ { 0x00313332, 3 },
+ { 0x00323332, 3 },
+ { 0x00333332, 3 },
+ { 0x00343332, 3 },
+ { 0x00353332, 3 },
+ { 0x00363332, 3 },
+ { 0x00373332, 3 },
+ { 0x00383332, 3 },
+ { 0x00393332, 3 },
+ { 0x00303432, 3 },
+ { 0x00313432, 3 },
+ { 0x00323432, 3 },
+ { 0x00333432, 3 },
+ { 0x00343432, 3 },
+ { 0x00353432, 3 },
+ { 0x00363432, 3 },
+ { 0x00373432, 3 },
+ { 0x00383432, 3 },
+ { 0x00393432, 3 },
+ { 0x00303532, 3 },
+ { 0x00313532, 3 },
+ { 0x00323532, 3 },
+ { 0x00333532, 3 },
+ { 0x00343532, 3 },
+ { 0x00353532, 3 },
+ { 0x00363532, 3 },
+ { 0x00373532, 3 },
+ { 0x00383532, 3 },
+ { 0x00393532, 3 },
+ { 0x00303632, 3 },
+ { 0x00313632, 3 },
+ { 0x00323632, 3 },
+ { 0x00333632, 3 },
+ { 0x00343632, 3 },
+ { 0x00353632, 3 },
+ { 0x00363632, 3 },
+ { 0x00373632, 3 },
+ { 0x00383632, 3 },
+ { 0x00393632, 3 },
+ { 0x00303732, 3 },
+ { 0x00313732, 3 },
+ { 0x00323732, 3 },
+ { 0x00333732, 3 },
+ { 0x00343732, 3 },
+ { 0x00353732, 3 },
+ { 0x00363732, 3 },
+ { 0x00373732, 3 },
+ { 0x00383732, 3 },
+ { 0x00393732, 3 },
+ { 0x00303832, 3 },
+ { 0x00313832, 3 },
+ { 0x00323832, 3 },
+ { 0x00333832, 3 },
+ { 0x00343832, 3 },
+ { 0x00353832, 3 },
+ { 0x00363832, 3 },
+ { 0x00373832, 3 },
+ { 0x00383832, 3 },
+ { 0x00393832, 3 },
+ { 0x00303932, 3 },
+ { 0x00313932, 3 },
+ { 0x00323932, 3 },
+ { 0x00333932, 3 },
+ { 0x00343932, 3 },
+ { 0x00353932, 3 },
+ { 0x00363932, 3 },
+ { 0x00373932, 3 },
+ { 0x00383932, 3 },
+ { 0x00393932, 3 },
+ { 0x00303033, 3 },
+ { 0x00313033, 3 },
+ { 0x00323033, 3 },
+ { 0x00333033, 3 },
+ { 0x00343033, 3 },
+ { 0x00353033, 3 },
+ { 0x00363033, 3 },
+ { 0x00373033, 3 },
+ { 0x00383033, 3 },
+ { 0x00393033, 3 },
+ { 0x00303133, 3 },
+ { 0x00313133, 3 },
+ { 0x00323133, 3 },
+ { 0x00333133, 3 },
+ { 0x00343133, 3 },
+ { 0x00353133, 3 },
+ { 0x00363133, 3 },
+ { 0x00373133, 3 },
+ { 0x00383133, 3 },
+ { 0x00393133, 3 },
+ { 0x00303233, 3 },
+ { 0x00313233, 3 },
+ { 0x00323233, 3 },
+ { 0x00333233, 3 },
+ { 0x00343233, 3 },
+ { 0x00353233, 3 },
+ { 0x00363233, 3 },
+ { 0x00373233, 3 },
+ { 0x00383233, 3 },
+ { 0x00393233, 3 },
+ { 0x00303333, 3 },
+ { 0x00313333, 3 },
+ { 0x00323333, 3 },
+ { 0x00333333, 3 },
+ { 0x00343333, 3 },
+ { 0x00353333, 3 },
+ { 0x00363333, 3 },
+ { 0x00373333, 3 },
+ { 0x00383333, 3 },
+ { 0x00393333, 3 },
+ { 0x00303433, 3 },
+ { 0x00313433, 3 },
+ { 0x00323433, 3 },
+ { 0x00333433, 3 },
+ { 0x00343433, 3 },
+ { 0x00353433, 3 },
+ { 0x00363433, 3 },
+ { 0x00373433, 3 },
+ { 0x00383433, 3 },
+ { 0x00393433, 3 },
+ { 0x00303533, 3 },
+ { 0x00313533, 3 },
+ { 0x00323533, 3 },
+ { 0x00333533, 3 },
+ { 0x00343533, 3 },
+ { 0x00353533, 3 },
+ { 0x00363533, 3 },
+ { 0x00373533, 3 },
+ { 0x00383533, 3 },
+ { 0x00393533, 3 },
+ { 0x00303633, 3 },
+ { 0x00313633, 3 },
+ { 0x00323633, 3 },
+ { 0x00333633, 3 },
+ { 0x00343633, 3 },
+ { 0x00353633, 3 },
+ { 0x00363633, 3 },
+ { 0x00373633, 3 },
+ { 0x00383633, 3 },
+ { 0x00393633, 3 },
+ { 0x00303733, 3 },
+ { 0x00313733, 3 },
+ { 0x00323733, 3 },
+ { 0x00333733, 3 },
+ { 0x00343733, 3 },
+ { 0x00353733, 3 },
+ { 0x00363733, 3 },
+ { 0x00373733, 3 },
+ { 0x00383733, 3 },
+ { 0x00393733, 3 },
+ { 0x00303833, 3 },
+ { 0x00313833, 3 },
+ { 0x00323833, 3 },
+ { 0x00333833, 3 },
+ { 0x00343833, 3 },
+ { 0x00353833, 3 },
+ { 0x00363833, 3 },
+ { 0x00373833, 3 },
+ { 0x00383833, 3 },
+ { 0x00393833, 3 },
+ { 0x00303933, 3 },
+ { 0x00313933, 3 },
+ { 0x00323933, 3 },
+ { 0x00333933, 3 },
+ { 0x00343933, 3 },
+ { 0x00353933, 3 },
+ { 0x00363933, 3 },
+ { 0x00373933, 3 },
+ { 0x00383933, 3 },
+ { 0x00393933, 3 },
+ { 0x00303034, 3 },
+ { 0x00313034, 3 },
+ { 0x00323034, 3 },
+ { 0x00333034, 3 },
+ { 0x00343034, 3 },
+ { 0x00353034, 3 },
+ { 0x00363034, 3 },
+ { 0x00373034, 3 },
+ { 0x00383034, 3 },
+ { 0x00393034, 3 },
+ { 0x00303134, 3 },
+ { 0x00313134, 3 },
+ { 0x00323134, 3 },
+ { 0x00333134, 3 },
+ { 0x00343134, 3 },
+ { 0x00353134, 3 },
+ { 0x00363134, 3 },
+ { 0x00373134, 3 },
+ { 0x00383134, 3 },
+ { 0x00393134, 3 },
+ { 0x00303234, 3 },
+ { 0x00313234, 3 },
+ { 0x00323234, 3 },
+ { 0x00333234, 3 },
+ { 0x00343234, 3 },
+ { 0x00353234, 3 },
+ { 0x00363234, 3 },
+ { 0x00373234, 3 },
+ { 0x00383234, 3 },
+ { 0x00393234, 3 },
+ { 0x00303334, 3 },
+ { 0x00313334, 3 },
+ { 0x00323334, 3 },
+ { 0x00333334, 3 },
+ { 0x00343334, 3 },
+ { 0x00353334, 3 },
+ { 0x00363334, 3 },
+ { 0x00373334, 3 },
+ { 0x00383334, 3 },
+ { 0x00393334, 3 },
+ { 0x00303434, 3 },
+ { 0x00313434, 3 },
+ { 0x00323434, 3 },
+ { 0x00333434, 3 },
+ { 0x00343434, 3 },
+ { 0x00353434, 3 },
+ { 0x00363434, 3 },
+ { 0x00373434, 3 },
+ { 0x00383434, 3 },
+ { 0x00393434, 3 },
+ { 0x00303534, 3 },
+ { 0x00313534, 3 },
+ { 0x00323534, 3 },
+ { 0x00333534, 3 },
+ { 0x00343534, 3 },
+ { 0x00353534, 3 },
+ { 0x00363534, 3 },
+ { 0x00373534, 3 },
+ { 0x00383534, 3 },
+ { 0x00393534, 3 },
+ { 0x00303634, 3 },
+ { 0x00313634, 3 },
+ { 0x00323634, 3 },
+ { 0x00333634, 3 },
+ { 0x00343634, 3 },
+ { 0x00353634, 3 },
+ { 0x00363634, 3 },
+ { 0x00373634, 3 },
+ { 0x00383634, 3 },
+ { 0x00393634, 3 },
+ { 0x00303734, 3 },
+ { 0x00313734, 3 },
+ { 0x00323734, 3 },
+ { 0x00333734, 3 },
+ { 0x00343734, 3 },
+ { 0x00353734, 3 },
+ { 0x00363734, 3 },
+ { 0x00373734, 3 },
+ { 0x00383734, 3 },
+ { 0x00393734, 3 },
+ { 0x00303834, 3 },
+ { 0x00313834, 3 },
+ { 0x00323834, 3 },
+ { 0x00333834, 3 },
+ { 0x00343834, 3 },
+ { 0x00353834, 3 },
+ { 0x00363834, 3 },
+ { 0x00373834, 3 },
+ { 0x00383834, 3 },
+ { 0x00393834, 3 },
+ { 0x00303934, 3 },
+ { 0x00313934, 3 },
+ { 0x00323934, 3 },
+ { 0x00333934, 3 },
+ { 0x00343934, 3 },
+ { 0x00353934, 3 },
+ { 0x00363934, 3 },
+ { 0x00373934, 3 },
+ { 0x00383934, 3 },
+ { 0x00393934, 3 },
+ { 0x00303035, 3 },
+ { 0x00313035, 3 },
+ { 0x00323035, 3 },
+ { 0x00333035, 3 },
+ { 0x00343035, 3 },
+ { 0x00353035, 3 },
+ { 0x00363035, 3 },
+ { 0x00373035, 3 },
+ { 0x00383035, 3 },
+ { 0x00393035, 3 },
+ { 0x00303135, 3 },
+ { 0x00313135, 3 },
+ { 0x00323135, 3 },
+ { 0x00333135, 3 },
+ { 0x00343135, 3 },
+ { 0x00353135, 3 },
+ { 0x00363135, 3 },
+ { 0x00373135, 3 },
+ { 0x00383135, 3 },
+ { 0x00393135, 3 },
+ { 0x00303235, 3 },
+ { 0x00313235, 3 },
+ { 0x00323235, 3 },
+ { 0x00333235, 3 },
+ { 0x00343235, 3 },
+ { 0x00353235, 3 },
+ { 0x00363235, 3 },
+ { 0x00373235, 3 },
+ { 0x00383235, 3 },
+ { 0x00393235, 3 },
+ { 0x00303335, 3 },
+ { 0x00313335, 3 },
+ { 0x00323335, 3 },
+ { 0x00333335, 3 },
+ { 0x00343335, 3 },
+ { 0x00353335, 3 },
+ { 0x00363335, 3 },
+ { 0x00373335, 3 },
+ { 0x00383335, 3 },
+ { 0x00393335, 3 },
+ { 0x00303435, 3 },
+ { 0x00313435, 3 },
+ { 0x00323435, 3 },
+ { 0x00333435, 3 },
+ { 0x00343435, 3 },
+ { 0x00353435, 3 },
+ { 0x00363435, 3 },
+ { 0x00373435, 3 },
+ { 0x00383435, 3 },
+ { 0x00393435, 3 },
+ { 0x00303535, 3 },
+ { 0x00313535, 3 },
+ { 0x00323535, 3 },
+ { 0x00333535, 3 },
+ { 0x00343535, 3 },
+ { 0x00353535, 3 },
+ { 0x00363535, 3 },
+ { 0x00373535, 3 },
+ { 0x00383535, 3 },
+ { 0x00393535, 3 },
+ { 0x00303635, 3 },
+ { 0x00313635, 3 },
+ { 0x00323635, 3 },
+ { 0x00333635, 3 },
+ { 0x00343635, 3 },
+ { 0x00353635, 3 },
+ { 0x00363635, 3 },
+ { 0x00373635, 3 },
+ { 0x00383635, 3 },
+ { 0x00393635, 3 },
+ { 0x00303735, 3 },
+ { 0x00313735, 3 },
+ { 0x00323735, 3 },
+ { 0x00333735, 3 },
+ { 0x00343735, 3 },
+ { 0x00353735, 3 },
+ { 0x00363735, 3 },
+ { 0x00373735, 3 },
+ { 0x00383735, 3 },
+ { 0x00393735, 3 },
+ { 0x00303835, 3 },
+ { 0x00313835, 3 },
+ { 0x00323835, 3 },
+ { 0x00333835, 3 },
+ { 0x00343835, 3 },
+ { 0x00353835, 3 },
+ { 0x00363835, 3 },
+ { 0x00373835, 3 },
+ { 0x00383835, 3 },
+ { 0x00393835, 3 },
+ { 0x00303935, 3 },
+ { 0x00313935, 3 },
+ { 0x00323935, 3 },
+ { 0x00333935, 3 },
+ { 0x00343935, 3 },
+ { 0x00353935, 3 },
+ { 0x00363935, 3 },
+ { 0x00373935, 3 },
+ { 0x00383935, 3 },
+ { 0x00393935, 3 },
+ { 0x00303036, 3 },
+ { 0x00313036, 3 },
+ { 0x00323036, 3 },
+ { 0x00333036, 3 },
+ { 0x00343036, 3 },
+ { 0x00353036, 3 },
+ { 0x00363036, 3 },
+ { 0x00373036, 3 },
+ { 0x00383036, 3 },
+ { 0x00393036, 3 },
+ { 0x00303136, 3 },
+ { 0x00313136, 3 },
+ { 0x00323136, 3 },
+ { 0x00333136, 3 },
+ { 0x00343136, 3 },
+ { 0x00353136, 3 },
+ { 0x00363136, 3 },
+ { 0x00373136, 3 },
+ { 0x00383136, 3 },
+ { 0x00393136, 3 },
+ { 0x00303236, 3 },
+ { 0x00313236, 3 },
+ { 0x00323236, 3 },
+ { 0x00333236, 3 },
+ { 0x00343236, 3 },
+ { 0x00353236, 3 },
+ { 0x00363236, 3 },
+ { 0x00373236, 3 },
+ { 0x00383236, 3 },
+ { 0x00393236, 3 },
+ { 0x00303336, 3 },
+ { 0x00313336, 3 },
+ { 0x00323336, 3 },
+ { 0x00333336, 3 },
+ { 0x00343336, 3 },
+ { 0x00353336, 3 },
+ { 0x00363336, 3 },
+ { 0x00373336, 3 },
+ { 0x00383336, 3 },
+ { 0x00393336, 3 },
+ { 0x00303436, 3 },
+ { 0x00313436, 3 },
+ { 0x00323436, 3 },
+ { 0x00333436, 3 },
+ { 0x00343436, 3 },
+ { 0x00353436, 3 },
+ { 0x00363436, 3 },
+ { 0x00373436, 3 },
+ { 0x00383436, 3 },
+ { 0x00393436, 3 },
+ { 0x00303536, 3 },
+ { 0x00313536, 3 },
+ { 0x00323536, 3 },
+ { 0x00333536, 3 },
+ { 0x00343536, 3 },
+ { 0x00353536, 3 },
+ { 0x00363536, 3 },
+ { 0x00373536, 3 },
+ { 0x00383536, 3 },
+ { 0x00393536, 3 },
+ { 0x00303636, 3 },
+ { 0x00313636, 3 },
+ { 0x00323636, 3 },
+ { 0x00333636, 3 },
+ { 0x00343636, 3 },
+ { 0x00353636, 3 },
+ { 0x00363636, 3 },
+ { 0x00373636, 3 },
+ { 0x00383636, 3 },
+ { 0x00393636, 3 },
+ { 0x00303736, 3 },
+ { 0x00313736, 3 },
+ { 0x00323736, 3 },
+ { 0x00333736, 3 },
+ { 0x00343736, 3 },
+ { 0x00353736, 3 },
+ { 0x00363736, 3 },
+ { 0x00373736, 3 },
+ { 0x00383736, 3 },
+ { 0x00393736, 3 },
+ { 0x00303836, 3 },
+ { 0x00313836, 3 },
+ { 0x00323836, 3 },
+ { 0x00333836, 3 },
+ { 0x00343836, 3 },
+ { 0x00353836, 3 },
+ { 0x00363836, 3 },
+ { 0x00373836, 3 },
+ { 0x00383836, 3 },
+ { 0x00393836, 3 },
+ { 0x00303936, 3 },
+ { 0x00313936, 3 },
+ { 0x00323936, 3 },
+ { 0x00333936, 3 },
+ { 0x00343936, 3 },
+ { 0x00353936, 3 },
+ { 0x00363936, 3 },
+ { 0x00373936, 3 },
+ { 0x00383936, 3 },
+ { 0x00393936, 3 },
+ { 0x00303037, 3 },
+ { 0x00313037, 3 },
+ { 0x00323037, 3 },
+ { 0x00333037, 3 },
+ { 0x00343037, 3 },
+ { 0x00353037, 3 },
+ { 0x00363037, 3 },
+ { 0x00373037, 3 },
+ { 0x00383037, 3 },
+ { 0x00393037, 3 },
+ { 0x00303137, 3 },
+ { 0x00313137, 3 },
+ { 0x00323137, 3 },
+ { 0x00333137, 3 },
+ { 0x00343137, 3 },
+ { 0x00353137, 3 },
+ { 0x00363137, 3 },
+ { 0x00373137, 3 },
+ { 0x00383137, 3 },
+ { 0x00393137, 3 },
+ { 0x00303237, 3 },
+ { 0x00313237, 3 },
+ { 0x00323237, 3 },
+ { 0x00333237, 3 },
+ { 0x00343237, 3 },
+ { 0x00353237, 3 },
+ { 0x00363237, 3 },
+ { 0x00373237, 3 },
+ { 0x00383237, 3 },
+ { 0x00393237, 3 },
+ { 0x00303337, 3 },
+ { 0x00313337, 3 },
+ { 0x00323337, 3 },
+ { 0x00333337, 3 },
+ { 0x00343337, 3 },
+ { 0x00353337, 3 },
+ { 0x00363337, 3 },
+ { 0x00373337, 3 },
+ { 0x00383337, 3 },
+ { 0x00393337, 3 },
+ { 0x00303437, 3 },
+ { 0x00313437, 3 },
+ { 0x00323437, 3 },
+ { 0x00333437, 3 },
+ { 0x00343437, 3 },
+ { 0x00353437, 3 },
+ { 0x00363437, 3 },
+ { 0x00373437, 3 },
+ { 0x00383437, 3 },
+ { 0x00393437, 3 },
+ { 0x00303537, 3 },
+ { 0x00313537, 3 },
+ { 0x00323537, 3 },
+ { 0x00333537, 3 },
+ { 0x00343537, 3 },
+ { 0x00353537, 3 },
+ { 0x00363537, 3 },
+ { 0x00373537, 3 },
+ { 0x00383537, 3 },
+ { 0x00393537, 3 },
+ { 0x00303637, 3 },
+ { 0x00313637, 3 },
+ { 0x00323637, 3 },
+ { 0x00333637, 3 },
+ { 0x00343637, 3 },
+ { 0x00353637, 3 },
+ { 0x00363637, 3 },
+ { 0x00373637, 3 },
+ { 0x00383637, 3 },
+ { 0x00393637, 3 },
+ { 0x00303737, 3 },
+ { 0x00313737, 3 },
+ { 0x00323737, 3 },
+ { 0x00333737, 3 },
+ { 0x00343737, 3 },
+ { 0x00353737, 3 },
+ { 0x00363737, 3 },
+ { 0x00373737, 3 },
+ { 0x00383737, 3 },
+ { 0x00393737, 3 },
+ { 0x00303837, 3 },
+ { 0x00313837, 3 },
+ { 0x00323837, 3 },
+ { 0x00333837, 3 },
+ { 0x00343837, 3 },
+ { 0x00353837, 3 },
+ { 0x00363837, 3 },
+ { 0x00373837, 3 },
+ { 0x00383837, 3 },
+ { 0x00393837, 3 },
+ { 0x00303937, 3 },
+ { 0x00313937, 3 },
+ { 0x00323937, 3 },
+ { 0x00333937, 3 },
+ { 0x00343937, 3 },
+ { 0x00353937, 3 },
+ { 0x00363937, 3 },
+ { 0x00373937, 3 },
+ { 0x00383937, 3 },
+ { 0x00393937, 3 },
+ { 0x00303038, 3 },
+ { 0x00313038, 3 },
+ { 0x00323038, 3 },
+ { 0x00333038, 3 },
+ { 0x00343038, 3 },
+ { 0x00353038, 3 },
+ { 0x00363038, 3 },
+ { 0x00373038, 3 },
+ { 0x00383038, 3 },
+ { 0x00393038, 3 },
+ { 0x00303138, 3 },
+ { 0x00313138, 3 },
+ { 0x00323138, 3 },
+ { 0x00333138, 3 },
+ { 0x00343138, 3 },
+ { 0x00353138, 3 },
+ { 0x00363138, 3 },
+ { 0x00373138, 3 },
+ { 0x00383138, 3 },
+ { 0x00393138, 3 },
+ { 0x00303238, 3 },
+ { 0x00313238, 3 },
+ { 0x00323238, 3 },
+ { 0x00333238, 3 },
+ { 0x00343238, 3 },
+ { 0x00353238, 3 },
+ { 0x00363238, 3 },
+ { 0x00373238, 3 },
+ { 0x00383238, 3 },
+ { 0x00393238, 3 },
+ { 0x00303338, 3 },
+ { 0x00313338, 3 },
+ { 0x00323338, 3 },
+ { 0x00333338, 3 },
+ { 0x00343338, 3 },
+ { 0x00353338, 3 },
+ { 0x00363338, 3 },
+ { 0x00373338, 3 },
+ { 0x00383338, 3 },
+ { 0x00393338, 3 },
+ { 0x00303438, 3 },
+ { 0x00313438, 3 },
+ { 0x00323438, 3 },
+ { 0x00333438, 3 },
+ { 0x00343438, 3 },
+ { 0x00353438, 3 },
+ { 0x00363438, 3 },
+ { 0x00373438, 3 },
+ { 0x00383438, 3 },
+ { 0x00393438, 3 },
+ { 0x00303538, 3 },
+ { 0x00313538, 3 },
+ { 0x00323538, 3 },
+ { 0x00333538, 3 },
+ { 0x00343538, 3 },
+ { 0x00353538, 3 },
+ { 0x00363538, 3 },
+ { 0x00373538, 3 },
+ { 0x00383538, 3 },
+ { 0x00393538, 3 },
+ { 0x00303638, 3 },
+ { 0x00313638, 3 },
+ { 0x00323638, 3 },
+ { 0x00333638, 3 },
+ { 0x00343638, 3 },
+ { 0x00353638, 3 },
+ { 0x00363638, 3 },
+ { 0x00373638, 3 },
+ { 0x00383638, 3 },
+ { 0x00393638, 3 },
+ { 0x00303738, 3 },
+ { 0x00313738, 3 },
+ { 0x00323738, 3 },
+ { 0x00333738, 3 },
+ { 0x00343738, 3 },
+ { 0x00353738, 3 },
+ { 0x00363738, 3 },
+ { 0x00373738, 3 },
+ { 0x00383738, 3 },
+ { 0x00393738, 3 },
+ { 0x00303838, 3 },
+ { 0x00313838, 3 },
+ { 0x00323838, 3 },
+ { 0x00333838, 3 },
+ { 0x00343838, 3 },
+ { 0x00353838, 3 },
+ { 0x00363838, 3 },
+ { 0x00373838, 3 },
+ { 0x00383838, 3 },
+ { 0x00393838, 3 },
+ { 0x00303938, 3 },
+ { 0x00313938, 3 },
+ { 0x00323938, 3 },
+ { 0x00333938, 3 },
+ { 0x00343938, 3 },
+ { 0x00353938, 3 },
+ { 0x00363938, 3 },
+ { 0x00373938, 3 },
+ { 0x00383938, 3 },
+ { 0x00393938, 3 },
+ { 0x00303039, 3 },
+ { 0x00313039, 3 },
+ { 0x00323039, 3 },
+ { 0x00333039, 3 },
+ { 0x00343039, 3 },
+ { 0x00353039, 3 },
+ { 0x00363039, 3 },
+ { 0x00373039, 3 },
+ { 0x00383039, 3 },
+ { 0x00393039, 3 },
+ { 0x00303139, 3 },
+ { 0x00313139, 3 },
+ { 0x00323139, 3 },
+ { 0x00333139, 3 },
+ { 0x00343139, 3 },
+ { 0x00353139, 3 },
+ { 0x00363139, 3 },
+ { 0x00373139, 3 },
+ { 0x00383139, 3 },
+ { 0x00393139, 3 },
+ { 0x00303239, 3 },
+ { 0x00313239, 3 },
+ { 0x00323239, 3 },
+ { 0x00333239, 3 },
+ { 0x00343239, 3 },
+ { 0x00353239, 3 },
+ { 0x00363239, 3 },
+ { 0x00373239, 3 },
+ { 0x00383239, 3 },
+ { 0x00393239, 3 },
+ { 0x00303339, 3 },
+ { 0x00313339, 3 },
+ { 0x00323339, 3 },
+ { 0x00333339, 3 },
+ { 0x00343339, 3 },
+ { 0x00353339, 3 },
+ { 0x00363339, 3 },
+ { 0x00373339, 3 },
+ { 0x00383339, 3 },
+ { 0x00393339, 3 },
+ { 0x00303439, 3 },
+ { 0x00313439, 3 },
+ { 0x00323439, 3 },
+ { 0x00333439, 3 },
+ { 0x00343439, 3 },
+ { 0x00353439, 3 },
+ { 0x00363439, 3 },
+ { 0x00373439, 3 },
+ { 0x00383439, 3 },
+ { 0x00393439, 3 },
+ { 0x00303539, 3 },
+ { 0x00313539, 3 },
+ { 0x00323539, 3 },
+ { 0x00333539, 3 },
+ { 0x00343539, 3 },
+ { 0x00353539, 3 },
+ { 0x00363539, 3 },
+ { 0x00373539, 3 },
+ { 0x00383539, 3 },
+ { 0x00393539, 3 },
+ { 0x00303639, 3 },
+ { 0x00313639, 3 },
+ { 0x00323639, 3 },
+ { 0x00333639, 3 },
+ { 0x00343639, 3 },
+ { 0x00353639, 3 },
+ { 0x00363639, 3 },
+ { 0x00373639, 3 },
+ { 0x00383639, 3 },
+ { 0x00393639, 3 },
+ { 0x00303739, 3 },
+ { 0x00313739, 3 },
+ { 0x00323739, 3 },
+ { 0x00333739, 3 },
+ { 0x00343739, 3 },
+ { 0x00353739, 3 },
+ { 0x00363739, 3 },
+ { 0x00373739, 3 },
+ { 0x00383739, 3 },
+ { 0x00393739, 3 },
+ { 0x00303839, 3 },
+ { 0x00313839, 3 },
+ { 0x00323839, 3 },
+ { 0x00333839, 3 },
+ { 0x00343839, 3 },
+ { 0x00353839, 3 },
+ { 0x00363839, 3 },
+ { 0x00373839, 3 },
+ { 0x00383839, 3 },
+ { 0x00393839, 3 },
+ { 0x00303939, 3 },
+ { 0x00313939, 3 },
+ { 0x00323939, 3 },
+ { 0x00333939, 3 },
+ { 0x00343939, 3 },
+ { 0x00353939, 3 },
+ { 0x00363939, 3 },
+ { 0x00373939, 3 },
+ { 0x00383939, 3 },
+ { 0x00393939, 3 },
+ { 0x30303031, 4 },
+ { 0x31303031, 4 },
+ { 0x32303031, 4 },
+ { 0x33303031, 4 },
+ { 0x34303031, 4 },
+ { 0x35303031, 4 },
+ { 0x36303031, 4 },
+ { 0x37303031, 4 },
+ { 0x38303031, 4 },
+ { 0x39303031, 4 },
+ { 0x30313031, 4 },
+ { 0x31313031, 4 },
+ { 0x32313031, 4 },
+ { 0x33313031, 4 },
+ { 0x34313031, 4 },
+ { 0x35313031, 4 },
+ { 0x36313031, 4 },
+ { 0x37313031, 4 },
+ { 0x38313031, 4 },
+ { 0x39313031, 4 },
+ { 0x30323031, 4 },
+ { 0x31323031, 4 },
+ { 0x32323031, 4 },
+ { 0x33323031, 4 }
};
-static void append_word (u32 w0[4], u32 w1[4], const u32 append[4], const u32 offset)
+void append_word (u32 w0[4], u32 w1[4], const u32 append[4], const u32 offset)
{
switch (offset)
{
}
}
-static void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], const u32 offset)
+void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], const u32 offset)
{
- switch (offset)
- {
- case 2:
- w0[0] = w0[0] | append[0] << 16;
- w0[1] = append[0] >> 16 | append[1] << 16;
- w0[2] = append[1] >> 16 | append[2] << 16;
- w0[3] = append[2] >> 16 | append[3] << 16;
- w1[0] = append[3] >> 16 | append[4] << 16;
- w1[1] = append[4] >> 16;
- break;
-
- case 3:
- w0[0] = w0[0] | append[0] << 24;
- w0[1] = append[0] >> 8 | append[1] << 24;
- w0[2] = append[1] >> 8 | append[2] << 24;
- w0[3] = append[2] >> 8 | append[3] << 24;
- w1[0] = append[3] >> 8 | append[4] << 24;
- w1[1] = append[4] >> 8;
- break;
-
- case 4:
- w0[1] = append[0];
- w0[2] = append[1];
- w0[3] = append[2];
- w1[0] = append[3];
- w1[1] = append[4];
- break;
+ u32 tmp0;
+ u32 tmp1;
+ u32 tmp2;
+ u32 tmp3;
+ u32 tmp4;
+ u32 tmp5;
- case 5:
- w0[1] = w0[1] | append[0] << 8;
- w0[2] = append[0] >> 24 | append[1] << 8;
- w0[3] = append[1] >> 24 | append[2] << 8;
- w1[0] = append[2] >> 24 | append[3] << 8;
- w1[1] = append[3] >> 24 | append[4] << 8;
- w1[2] = append[4] >> 24;
- break;
+ #if defined IS_AMD || defined IS_GENERIC
- case 6:
- w0[1] = w0[1] | append[0] << 16;
- w0[2] = append[0] >> 16 | append[1] << 16;
- w0[3] = append[1] >> 16 | append[2] << 16;
- w1[0] = append[2] >> 16 | append[3] << 16;
- w1[1] = append[3] >> 16 | append[4] << 16;
- w1[2] = append[4] >> 16;
- break;
+ const int offset_minus_4 = 4 - (offset & 3);
- case 7:
- w0[1] = w0[1] | append[0] << 24;
- w0[2] = append[0] >> 8 | append[1] << 24;
- w0[3] = append[1] >> 8 | append[2] << 24;
- w1[0] = append[2] >> 8 | append[3] << 24;
- w1[1] = append[3] >> 8 | append[4] << 24;
- w1[2] = append[4] >> 8;
- break;
+ tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
+ tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
+ tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
+ tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
+ tmp4 = amd_bytealign (append[4], append[3], offset_minus_4);
+ tmp5 = amd_bytealign ( 0, append[4], offset_minus_4);
- case 8:
- w0[2] = append[0];
- w0[3] = append[1];
- w1[0] = append[2];
- w1[1] = append[3];
- w1[2] = append[4];
- break;
+ const u32 mod = offset & 3;
- case 9:
- w0[2] = w0[2] | append[0] << 8;
- w0[3] = append[0] >> 24 | append[1] << 8;
- w1[0] = append[1] >> 24 | append[2] << 8;
- w1[1] = append[2] >> 24 | append[3] << 8;
- w1[2] = append[3] >> 24 | append[4] << 8;
- w1[3] = append[4] >> 24;
- break;
+ if (mod == 0)
+ {
+ tmp0 = tmp1;
+ tmp1 = tmp2;
+ tmp2 = tmp3;
+ tmp3 = tmp4;
+ tmp4 = tmp5;
+ tmp5 = 0;
+ }
- case 10:
- w0[2] = w0[2] | append[0] << 16;
- w0[3] = append[0] >> 16 | append[1] << 16;
- w1[0] = append[1] >> 16 | append[2] << 16;
- w1[1] = append[2] >> 16 | append[3] << 16;
- w1[2] = append[3] >> 16 | append[4] << 16;
- w1[3] = append[4] >> 16;
- break;
+ #endif
- case 11:
- w0[2] = w0[2] | append[0] << 24;
- w0[3] = append[0] >> 8 | append[1] << 24;
- w1[0] = append[1] >> 8 | append[2] << 24;
- w1[1] = append[2] >> 8 | append[3] << 24;
- w1[2] = append[3] >> 8 | append[4] << 24;
- w1[3] = append[4] >> 8;
- break;
+ #ifdef IS_NV
- case 12:
- w0[3] = append[0];
- w1[0] = append[1];
- w1[1] = append[2];
- w1[2] = append[3];
- w1[3] = append[4];
- break;
+ const int offset_minus_4 = 4 - (offset & 3);
- case 13:
- w0[3] = w0[3] | append[0] << 8;
- w1[0] = append[0] >> 24 | append[1] << 8;
- w1[1] = append[1] >> 24 | append[2] << 8;
- w1[2] = append[2] >> 24 | append[3] << 8;
- w1[3] = append[3] >> 24 | append[4] << 8;
- w2[0] = append[4] >> 24;
- break;
+ const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
- case 14:
- w0[3] = w0[3] | append[0] << 16;
- w1[0] = append[0] >> 16 | append[1] << 16;
- w1[1] = append[1] >> 16 | append[2] << 16;
- w1[2] = append[2] >> 16 | append[3] << 16;
- w1[3] = append[3] >> 16 | append[4] << 16;
- w2[0] = append[4] >> 16;
- break;
+ tmp0 = __byte_perm ( 0, append[0], selector);
+ tmp1 = __byte_perm (append[0], append[1], selector);
+ tmp2 = __byte_perm (append[1], append[2], selector);
+ tmp3 = __byte_perm (append[2], append[3], selector);
+ tmp4 = __byte_perm (append[3], append[4], selector);
+ tmp5 = __byte_perm (append[4], 0, selector);
- case 15:
- w0[3] = w0[3] | append[0] << 24;
- w1[0] = append[0] >> 8 | append[1] << 24;
- w1[1] = append[1] >> 8 | append[2] << 24;
- w1[2] = append[2] >> 8 | append[3] << 24;
- w1[3] = append[3] >> 8 | append[4] << 24;
- w2[0] = append[4] >> 8;
- break;
+ #endif
- case 16:
- w1[0] = append[0];
- w1[1] = append[1];
- w1[2] = append[2];
- w1[3] = append[3];
- w2[0] = append[4];
- break;
+ const u32 div = offset / 4;
- case 17:
- w1[0] = w1[0] | append[0] << 8;
- w1[1] = append[0] >> 24 | append[1] << 8;
- w1[2] = append[1] >> 24 | append[2] << 8;
- w1[3] = append[2] >> 24 | append[3] << 8;
- w2[0] = append[3] >> 24 | append[4] << 8;
- w2[1] = append[4] >> 24;
- break;
-
- case 18:
- w1[0] = w1[0] | append[0] << 16;
- w1[1] = append[0] >> 16 | append[1] << 16;
- w1[2] = append[1] >> 16 | append[2] << 16;
- w1[3] = append[2] >> 16 | append[3] << 16;
- w2[0] = append[3] >> 16 | append[4] << 16;
- w2[1] = append[4] >> 16;
- break;
-
- case 19:
- w1[0] = w1[0] | append[0] << 24;
- w1[1] = append[0] >> 8 | append[1] << 24;
- w1[2] = append[1] >> 8 | append[2] << 24;
- w1[3] = append[2] >> 8 | append[3] << 24;
- w2[0] = append[3] >> 8 | append[4] << 24;
- w2[1] = append[4] >> 8;
- break;
+ switch (div)
+ {
+ case 0: w0[0] |= tmp0;
+ w0[1] = tmp1;
+ w0[2] = tmp2;
+ w0[3] = tmp3;
+ w1[0] = tmp4;
+ w1[1] = tmp5;
+ break;
+ case 1: w0[1] |= tmp0;
+ w0[2] = tmp1;
+ w0[3] = tmp2;
+ w1[0] = tmp3;
+ w1[1] = tmp4;
+ w1[2] = tmp5;
+ break;
+ case 2: w0[2] |= tmp0;
+ w0[3] = tmp1;
+ w1[0] = tmp2;
+ w1[1] = tmp3;
+ w1[2] = tmp4;
+ w1[3] = tmp5;
+ break;
+ case 3: w0[3] |= tmp0;
+ w1[0] = tmp1;
+ w1[1] = tmp2;
+ w1[2] = tmp3;
+ w1[3] = tmp4;
+ w2[0] = tmp5;
+ break;
+ case 4: w1[0] |= tmp0;
+ w1[1] = tmp1;
+ w1[2] = tmp2;
+ w1[3] = tmp3;
+ w2[0] = tmp4;
+ w2[1] = tmp5;
+ break;
}
}
-static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
+void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
{
u32 A = digest[0];
u32 B = digest[1];
digest[4] += E;
}
-__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)
+__kernel void m05800_init (__global pw_t *pws, __global kernel_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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
* base
tmps[gid].digest_buf[4] = digest[4];
}
-__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)
+__kernel void m05800_loop (__global pw_t *pws, __global kernel_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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
* base
*/
const u32 gid = get_global_id (0);
+ const u32 lid = get_local_id (0);
+ const u32 lsz = get_local_size (0);
+
+ /**
+ * cache precomputed conversion table in shared memory
+ */
+
+ __local entry_t s_pc[1024];
+
+ for (u32 i = lid; i < 1024; i += lsz)
+ {
+ s_pc[i] = pc[i];
+ }
+
+ barrier (CLK_LOCAL_MEM_FENCE);
+
+ if (gid >= gid_max) return;
+
+ /**
+ * base
+ */
u32 word_buf[4];
const u32 pw_len = pws[gid].pw_len;
- const u32 lid = get_local_id (0);
-
u32 digest[5];
digest[0] = tmps[gid].digest_buf[0];
digest[3] = tmps[gid].digest_buf[3];
digest[4] = tmps[gid].digest_buf[4];
- /**
- * cache precomputed conversion table in shared memory
- */
-
- __local entry_t s_pc[1024];
-
- const u32 lid16 = lid * 16;
-
- s_pc[lid16 + 0] = pc[lid16 + 0];
- s_pc[lid16 + 1] = pc[lid16 + 1];
- s_pc[lid16 + 2] = pc[lid16 + 2];
- s_pc[lid16 + 3] = pc[lid16 + 3];
- s_pc[lid16 + 4] = pc[lid16 + 4];
- s_pc[lid16 + 5] = pc[lid16 + 5];
- s_pc[lid16 + 6] = pc[lid16 + 6];
- s_pc[lid16 + 7] = pc[lid16 + 7];
- s_pc[lid16 + 8] = pc[lid16 + 8];
- s_pc[lid16 + 9] = pc[lid16 + 9];
- s_pc[lid16 + 10] = pc[lid16 + 10];
- s_pc[lid16 + 11] = pc[lid16 + 11];
- s_pc[lid16 + 12] = pc[lid16 + 12];
- s_pc[lid16 + 13] = pc[lid16 + 13];
- s_pc[lid16 + 14] = pc[lid16 + 14];
- s_pc[lid16 + 15] = pc[lid16 + 15];
-
- barrier (CLK_LOCAL_MEM_FENCE);
-
- if (gid >= gid_max) return;
-
/**
* salt
*/
tmps[gid].digest_buf[4] = digest[4];
}
-__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)
+__kernel void m05800_comp (__global pw_t *pws, __global kernel_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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
* modifier
projects / hashcat.git / blobdiff