/** * Author......: Jens Steube * License.....: MIT */ #define _WHIRLPOOL_ #include "include/constants.h" #include "include/kernel_vendor.h" #define DGST_R0 0 #define DGST_R1 1 #define DGST_R2 2 #define DGST_R3 3 #include "include/kernel_functions.c" #include "OpenCL/types_ocl.c" #include "OpenCL/common.c" #define COMPARE_S "OpenCL/check_single_comp4.c" #define COMPARE_M "OpenCL/check_multi_comp4.c" #define R 10 #define BOX(S,n,i) (S)[(n)][(i)] __constant u32 Ch[8][256] = { { 0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8, 0x87872687, 0xb8b8dab8, 0x01010401, 0x4f4f214f, 0x3636d836, 0xa6a6a2a6, 0xd2d26fd2, 0xf5f5f3f5, 0x7979f979, 0x6f6fa16f, 0x91917e91, 0x52525552, 0x60609d60, 0xbcbccabc, 0x9b9b569b, 0x8e8e028e, 0xa3a3b6a3, 0x0c0c300c, 0x7b7bf17b, 0x3535d435, 0x1d1d741d, 0xe0e0a7e0, 0xd7d77bd7, 0xc2c22fc2, 0x2e2eb82e, 0x4b4b314b, 0xfefedffe, 0x57574157, 0x15155415, 0x7777c177, 0x3737dc37, 0xe5e5b3e5, 0x9f9f469f, 0xf0f0e7f0, 0x4a4a354a, 0xdada4fda, 0x58587d58, 0xc9c903c9, 0x2929a429, 0x0a0a280a, 0xb1b1feb1, 0xa0a0baa0, 0x6b6bb16b, 0x85852e85, 0xbdbdcebd, 0x5d5d695d, 0x10104010, 0xf4f4f7f4, 0xcbcb0bcb, 0x3e3ef83e, 0x05051405, 0x67678167, 0xe4e4b7e4, 0x27279c27, 0x41411941, 0x8b8b168b, 0xa7a7a6a7, 0x7d7de97d, 0x95956e95, 0xd8d847d8, 0xfbfbcbfb, 0xeeee9fee, 0x7c7ced7c, 0x66668566, 0xdddd53dd, 0x17175c17, 0x47470147, 0x9e9e429e, 0xcaca0fca, 0x2d2db42d, 0xbfbfc6bf, 0x07071c07, 0xadad8ead, 0x5a5a755a, 0x83833683, 0x3333cc33, 0x63639163, 0x02020802, 0xaaaa92aa, 0x7171d971, 0xc8c807c8, 0x19196419, 0x49493949, 0xd9d943d9, 0xf2f2eff2, 0xe3e3abe3, 0x5b5b715b, 0x88881a88, 0x9a9a529a, 0x26269826, 0x3232c832, 0xb0b0fab0, 0xe9e983e9, 0x0f0f3c0f, 0xd5d573d5, 0x80803a80, 0xbebec2be, 0xcdcd13cd, 0x3434d034, 0x48483d48, 0xffffdbff, 0x7a7af57a, 0x90907a90, 0x5f5f615f, 0x20208020, 0x6868bd68, 0x1a1a681a, 0xaeae82ae, 0xb4b4eab4, 0x54544d54, 0x93937693, 0x22228822, 0x64648d64, 0xf1f1e3f1, 0x7373d173, 0x12124812, 0x40401d40, 0x08082008, 0xc3c32bc3, 0xecec97ec, 0xdbdb4bdb, 0xa1a1bea1, 0x8d8d0e8d, 0x3d3df43d, 0x97976697, 0x00000000, 0xcfcf1bcf, 0x2b2bac2b, 0x7676c576, 0x82823282, 0xd6d67fd6, 0x1b1b6c1b, 0xb5b5eeb5, 0xafaf86af, 0x6a6ab56a, 0x50505d50, 0x45450945, 0xf3f3ebf3, 0x3030c030, 0xefef9bef, 0x3f3ffc3f, 0x55554955, 0xa2a2b2a2, 0xeaea8fea, 0x65658965, 0xbabad2ba, 0x2f2fbc2f, 0xc0c027c0, 0xdede5fde, 0x1c1c701c, 0xfdfdd3fd, 0x4d4d294d, 0x92927292, 0x7575c975, 0x06061806, 0x8a8a128a, 0xb2b2f2b2, 0xe6e6bfe6, 0x0e0e380e, 0x1f1f7c1f, 0x62629562, 0xd4d477d4, 0xa8a89aa8, 0x96966296, 0xf9f9c3f9, 0xc5c533c5, 0x25259425, 0x59597959, 0x84842a84, 0x7272d572, 0x3939e439, 0x4c4c2d4c, 0x5e5e655e, 0x7878fd78, 0x3838e038, 0x8c8c0a8c, 0xd1d163d1, 0xa5a5aea5, 0xe2e2afe2, 0x61619961, 0xb3b3f6b3, 0x21218421, 0x9c9c4a9c, 0x1e1e781e, 0x43431143, 0xc7c73bc7, 0xfcfcd7fc, 0x04041004, 0x51515951, 0x99995e99, 0x6d6da96d, 0x0d0d340d, 0xfafacffa, 0xdfdf5bdf, 0x7e7ee57e, 0x24249024, 0x3b3bec3b, 0xabab96ab, 0xcece1fce, 0x11114411, 0x8f8f068f, 0x4e4e254e, 0xb7b7e6b7, 0xebeb8beb, 0x3c3cf03c, 0x81813e81, 0x94946a94, 0xf7f7fbf7, 0xb9b9deb9, 0x13134c13, 0x2c2cb02c, 0xd3d36bd3, 0xe7e7bbe7, 0x6e6ea56e, 0xc4c437c4, 0x03030c03, 0x56564556, 0x44440d44, 0x7f7fe17f, 0xa9a99ea9, 0x2a2aa82a, 0xbbbbd6bb, 0xc1c123c1, 0x53535153, 0xdcdc57dc, 0x0b0b2c0b, 0x9d9d4e9d, 0x6c6cad6c, 0x3131c431, 0x7474cd74, 0xf6f6fff6, 0x46460546, 0xacac8aac, 0x89891e89, 0x14145014, 0xe1e1a3e1, 0x16165816, 0x3a3ae83a, 0x6969b969, 0x09092409, 0x7070dd70, 0xb6b6e2b6, 0xd0d067d0, 0xeded93ed, 0xcccc17cc, 0x42421542, 0x98985a98, 0xa4a4aaa4, 0x2828a028, 0x5c5c6d5c, 0xf8f8c7f8, 0x86862286, }, { 0xd8181860, 0x2623238c, 0xb8c6c63f, 0xfbe8e887, 0xcb878726, 0x11b8b8da, 0x09010104, 0x0d4f4f21, 0x9b3636d8, 0xffa6a6a2, 0x0cd2d26f, 0x0ef5f5f3, 0x967979f9, 0x306f6fa1, 0x6d91917e, 0xf8525255, 0x4760609d, 0x35bcbcca, 0x379b9b56, 0x8a8e8e02, 0xd2a3a3b6, 0x6c0c0c30, 0x847b7bf1, 0x803535d4, 0xf51d1d74, 0xb3e0e0a7, 0x21d7d77b, 0x9cc2c22f, 0x432e2eb8, 0x294b4b31, 0x5dfefedf, 0xd5575741, 0xbd151554, 0xe87777c1, 0x923737dc, 0x9ee5e5b3, 0x139f9f46, 0x23f0f0e7, 0x204a4a35, 0x44dada4f, 0xa258587d, 0xcfc9c903, 0x7c2929a4, 0x5a0a0a28, 0x50b1b1fe, 0xc9a0a0ba, 0x146b6bb1, 0xd985852e, 0x3cbdbdce, 0x8f5d5d69, 0x90101040, 0x07f4f4f7, 0xddcbcb0b, 0xd33e3ef8, 0x2d050514, 0x78676781, 0x97e4e4b7, 0x0227279c, 0x73414119, 0xa78b8b16, 0xf6a7a7a6, 0xb27d7de9, 0x4995956e, 0x56d8d847, 0x70fbfbcb, 0xcdeeee9f, 0xbb7c7ced, 0x71666685, 0x7bdddd53, 0xaf17175c, 0x45474701, 0x1a9e9e42, 0xd4caca0f, 0x582d2db4, 0x2ebfbfc6, 0x3f07071c, 0xacadad8e, 0xb05a5a75, 0xef838336, 0xb63333cc, 0x5c636391, 0x12020208, 0x93aaaa92, 0xde7171d9, 0xc6c8c807, 0xd1191964, 0x3b494939, 0x5fd9d943, 0x31f2f2ef, 0xa8e3e3ab, 0xb95b5b71, 0xbc88881a, 0x3e9a9a52, 0x0b262698, 0xbf3232c8, 0x59b0b0fa, 0xf2e9e983, 0x770f0f3c, 0x33d5d573, 0xf480803a, 0x27bebec2, 0xebcdcd13, 0x893434d0, 0x3248483d, 0x54ffffdb, 0x8d7a7af5, 0x6490907a, 0x9d5f5f61, 0x3d202080, 0x0f6868bd, 0xca1a1a68, 0xb7aeae82, 0x7db4b4ea, 0xce54544d, 0x7f939376, 0x2f222288, 0x6364648d, 0x2af1f1e3, 0xcc7373d1, 0x82121248, 0x7a40401d, 0x48080820, 0x95c3c32b, 0xdfecec97, 0x4ddbdb4b, 0xc0a1a1be, 0x918d8d0e, 0xc83d3df4, 0x5b979766, 0x00000000, 0xf9cfcf1b, 0x6e2b2bac, 0xe17676c5, 0xe6828232, 0x28d6d67f, 0xc31b1b6c, 0x74b5b5ee, 0xbeafaf86, 0x1d6a6ab5, 0xea50505d, 0x57454509, 0x38f3f3eb, 0xad3030c0, 0xc4efef9b, 0xda3f3ffc, 0xc7555549, 0xdba2a2b2, 0xe9eaea8f, 0x6a656589, 0x03babad2, 0x4a2f2fbc, 0x8ec0c027, 0x60dede5f, 0xfc1c1c70, 0x46fdfdd3, 0x1f4d4d29, 0x76929272, 0xfa7575c9, 0x36060618, 0xae8a8a12, 0x4bb2b2f2, 0x85e6e6bf, 0x7e0e0e38, 0xe71f1f7c, 0x55626295, 0x3ad4d477, 0x81a8a89a, 0x52969662, 0x62f9f9c3, 0xa3c5c533, 0x10252594, 0xab595979, 0xd084842a, 0xc57272d5, 0xec3939e4, 0x164c4c2d, 0x945e5e65, 0x9f7878fd, 0xe53838e0, 0x988c8c0a, 0x17d1d163, 0xe4a5a5ae, 0xa1e2e2af, 0x4e616199, 0x42b3b3f6, 0x34212184, 0x089c9c4a, 0xee1e1e78, 0x61434311, 0xb1c7c73b, 0x4ffcfcd7, 0x24040410, 0xe3515159, 0x2599995e, 0x226d6da9, 0x650d0d34, 0x79fafacf, 0x69dfdf5b, 0xa97e7ee5, 0x19242490, 0xfe3b3bec, 0x9aabab96, 0xf0cece1f, 0x99111144, 0x838f8f06, 0x044e4e25, 0x66b7b7e6, 0xe0ebeb8b, 0xc13c3cf0, 0xfd81813e, 0x4094946a, 0x1cf7f7fb, 0x18b9b9de, 0x8b13134c, 0x512c2cb0, 0x05d3d36b, 0x8ce7e7bb, 0x396e6ea5, 0xaac4c437, 0x1b03030c, 0xdc565645, 0x5e44440d, 0xa07f7fe1, 0x88a9a99e, 0x672a2aa8, 0x0abbbbd6, 0x87c1c123, 0xf1535351, 0x72dcdc57, 0x530b0b2c, 0x019d9d4e, 0x2b6c6cad, 0xa43131c4, 0xf37474cd, 0x15f6f6ff, 0x4c464605, 0xa5acac8a, 0xb589891e, 0xb4141450, 0xbae1e1a3, 0xa6161658, 0xf73a3ae8, 0x066969b9, 0x41090924, 0xd77070dd, 0x6fb6b6e2, 0x1ed0d067, 0xd6eded93, 0xe2cccc17, 0x68424215, 0x2c98985a, 0xeda4a4aa, 0x752828a0, 0x865c5c6d, 0x6bf8f8c7, 0xc2868622, }, { 0x30d81818, 0x46262323, 0x91b8c6c6, 0xcdfbe8e8, 0x13cb8787, 0x6d11b8b8, 0x02090101, 0x9e0d4f4f, 0x6c9b3636, 0x51ffa6a6, 0xb90cd2d2, 0xf70ef5f5, 0xf2967979, 0xde306f6f, 0x3f6d9191, 0xa4f85252, 0xc0476060, 0x6535bcbc, 0x2b379b9b, 0x018a8e8e, 0x5bd2a3a3, 0x186c0c0c, 0xf6847b7b, 0x6a803535, 0x3af51d1d, 0xddb3e0e0, 0xb321d7d7, 0x999cc2c2, 0x5c432e2e, 0x96294b4b, 0xe15dfefe, 0xaed55757, 0x2abd1515, 0xeee87777, 0x6e923737, 0xd79ee5e5, 0x23139f9f, 0xfd23f0f0, 0x94204a4a, 0xa944dada, 0xb0a25858, 0x8fcfc9c9, 0x527c2929, 0x145a0a0a, 0x7f50b1b1, 0x5dc9a0a0, 0xd6146b6b, 0x17d98585, 0x673cbdbd, 0xba8f5d5d, 0x20901010, 0xf507f4f4, 0x8bddcbcb, 0x7cd33e3e, 0x0a2d0505, 0xce786767, 0xd597e4e4, 0x4e022727, 0x82734141, 0x0ba78b8b, 0x53f6a7a7, 0xfab27d7d, 0x37499595, 0xad56d8d8, 0xeb70fbfb, 0xc1cdeeee, 0xf8bb7c7c, 0xcc716666, 0xa77bdddd, 0x2eaf1717, 0x8e454747, 0x211a9e9e, 0x89d4caca, 0x5a582d2d, 0x632ebfbf, 0x0e3f0707, 0x47acadad, 0xb4b05a5a, 0x1bef8383, 0x66b63333, 0xc65c6363, 0x04120202, 0x4993aaaa, 0xe2de7171, 0x8dc6c8c8, 0x32d11919, 0x923b4949, 0xaf5fd9d9, 0xf931f2f2, 0xdba8e3e3, 0xb6b95b5b, 0x0dbc8888, 0x293e9a9a, 0x4c0b2626, 0x64bf3232, 0x7d59b0b0, 0xcff2e9e9, 0x1e770f0f, 0xb733d5d5, 0x1df48080, 0x6127bebe, 0x87ebcdcd, 0x68893434, 0x90324848, 0xe354ffff, 0xf48d7a7a, 0x3d649090, 0xbe9d5f5f, 0x403d2020, 0xd00f6868, 0x34ca1a1a, 0x41b7aeae, 0x757db4b4, 0xa8ce5454, 0x3b7f9393, 0x442f2222, 0xc8636464, 0xff2af1f1, 0xe6cc7373, 0x24821212, 0x807a4040, 0x10480808, 0x9b95c3c3, 0xc5dfecec, 0xab4ddbdb, 0x5fc0a1a1, 0x07918d8d, 0x7ac83d3d, 0x335b9797, 0x00000000, 0x83f9cfcf, 0x566e2b2b, 0xece17676, 0x19e68282, 0xb128d6d6, 0x36c31b1b, 0x7774b5b5, 0x43beafaf, 0xd41d6a6a, 0xa0ea5050, 0x8a574545, 0xfb38f3f3, 0x60ad3030, 0xc3c4efef, 0x7eda3f3f, 0xaac75555, 0x59dba2a2, 0xc9e9eaea, 0xca6a6565, 0x6903baba, 0x5e4a2f2f, 0x9d8ec0c0, 0xa160dede, 0x38fc1c1c, 0xe746fdfd, 0x9a1f4d4d, 0x39769292, 0xeafa7575, 0x0c360606, 0x09ae8a8a, 0x794bb2b2, 0xd185e6e6, 0x1c7e0e0e, 0x3ee71f1f, 0xc4556262, 0xb53ad4d4, 0x4d81a8a8, 0x31529696, 0xef62f9f9, 0x97a3c5c5, 0x4a102525, 0xb2ab5959, 0x15d08484, 0xe4c57272, 0x72ec3939, 0x98164c4c, 0xbc945e5e, 0xf09f7878, 0x70e53838, 0x05988c8c, 0xbf17d1d1, 0x57e4a5a5, 0xd9a1e2e2, 0xc24e6161, 0x7b42b3b3, 0x42342121, 0x25089c9c, 0x3cee1e1e, 0x86614343, 0x93b1c7c7, 0xe54ffcfc, 0x08240404, 0xa2e35151, 0x2f259999, 0xda226d6d, 0x1a650d0d, 0xe979fafa, 0xa369dfdf, 0xfca97e7e, 0x48192424, 0x76fe3b3b, 0x4b9aabab, 0x81f0cece, 0x22991111, 0x03838f8f, 0x9c044e4e, 0x7366b7b7, 0xcbe0ebeb, 0x78c13c3c, 0x1ffd8181, 0x35409494, 0xf31cf7f7, 0x6f18b9b9, 0x268b1313, 0x58512c2c, 0xbb05d3d3, 0xd38ce7e7, 0xdc396e6e, 0x95aac4c4, 0x061b0303, 0xacdc5656, 0x885e4444, 0xfea07f7f, 0x4f88a9a9, 0x54672a2a, 0x6b0abbbb, 0x9f87c1c1, 0xa6f15353, 0xa572dcdc, 0x16530b0b, 0x27019d9d, 0xd82b6c6c, 0x62a43131, 0xe8f37474, 0xf115f6f6, 0x8c4c4646, 0x45a5acac, 0x0fb58989, 0x28b41414, 0xdfbae1e1, 0x2ca61616, 0x74f73a3a, 0xd2066969, 0x12410909, 0xe0d77070, 0x716fb6b6, 0xbd1ed0d0, 0xc7d6eded, 0x85e2cccc, 0x84684242, 0x2d2c9898, 0x55eda4a4, 0x50752828, 0xb8865c5c, 0xed6bf8f8, 0x11c28686, }, { 0x7830d818, 0xaf462623, 0xf991b8c6, 0x6fcdfbe8, 0xa113cb87, 0x626d11b8, 0x05020901, 0x6e9e0d4f, 0xee6c9b36, 0x0451ffa6, 0xbdb90cd2, 0x06f70ef5, 0x80f29679, 0xcede306f, 0xef3f6d91, 0x07a4f852, 0xfdc04760, 0x766535bc, 0xcd2b379b, 0x8c018a8e, 0x155bd2a3, 0x3c186c0c, 0x8af6847b, 0xe16a8035, 0x693af51d, 0x47ddb3e0, 0xacb321d7, 0xed999cc2, 0x965c432e, 0x7a96294b, 0x21e15dfe, 0x16aed557, 0x412abd15, 0xb6eee877, 0xeb6e9237, 0x56d79ee5, 0xd923139f, 0x17fd23f0, 0x7f94204a, 0x95a944da, 0x25b0a258, 0xca8fcfc9, 0x8d527c29, 0x22145a0a, 0x4f7f50b1, 0x1a5dc9a0, 0xdad6146b, 0xab17d985, 0x73673cbd, 0x34ba8f5d, 0x50209010, 0x03f507f4, 0xc08bddcb, 0xc67cd33e, 0x110a2d05, 0xe6ce7867, 0x53d597e4, 0xbb4e0227, 0x58827341, 0x9d0ba78b, 0x0153f6a7, 0x94fab27d, 0xfb374995, 0x9fad56d8, 0x30eb70fb, 0x71c1cdee, 0x91f8bb7c, 0xe3cc7166, 0x8ea77bdd, 0x4b2eaf17, 0x468e4547, 0xdc211a9e, 0xc589d4ca, 0x995a582d, 0x79632ebf, 0x1b0e3f07, 0x2347acad, 0x2fb4b05a, 0xb51bef83, 0xff66b633, 0xf2c65c63, 0x0a041202, 0x384993aa, 0xa8e2de71, 0xcf8dc6c8, 0x7d32d119, 0x70923b49, 0x9aaf5fd9, 0x1df931f2, 0x48dba8e3, 0x2ab6b95b, 0x920dbc88, 0xc8293e9a, 0xbe4c0b26, 0xfa64bf32, 0x4a7d59b0, 0x6acff2e9, 0x331e770f, 0xa6b733d5, 0xba1df480, 0x7c6127be, 0xde87ebcd, 0xe4688934, 0x75903248, 0x24e354ff, 0x8ff48d7a, 0xea3d6490, 0x3ebe9d5f, 0xa0403d20, 0xd5d00f68, 0x7234ca1a, 0x2c41b7ae, 0x5e757db4, 0x19a8ce54, 0xe53b7f93, 0xaa442f22, 0xe9c86364, 0x12ff2af1, 0xa2e6cc73, 0x5a248212, 0x5d807a40, 0x28104808, 0xe89b95c3, 0x7bc5dfec, 0x90ab4ddb, 0x1f5fc0a1, 0x8307918d, 0xc97ac83d, 0xf1335b97, 0x00000000, 0xd483f9cf, 0x87566e2b, 0xb3ece176, 0xb019e682, 0xa9b128d6, 0x7736c31b, 0x5b7774b5, 0x2943beaf, 0xdfd41d6a, 0x0da0ea50, 0x4c8a5745, 0x18fb38f3, 0xf060ad30, 0x74c3c4ef, 0xc37eda3f, 0x1caac755, 0x1059dba2, 0x65c9e9ea, 0xecca6a65, 0x686903ba, 0x935e4a2f, 0xe79d8ec0, 0x81a160de, 0x6c38fc1c, 0x2ee746fd, 0x649a1f4d, 0xe0397692, 0xbceafa75, 0x1e0c3606, 0x9809ae8a, 0x40794bb2, 0x59d185e6, 0x361c7e0e, 0x633ee71f, 0xf7c45562, 0xa3b53ad4, 0x324d81a8, 0xf4315296, 0x3aef62f9, 0xf697a3c5, 0xb14a1025, 0x20b2ab59, 0xae15d084, 0xa7e4c572, 0xdd72ec39, 0x6198164c, 0x3bbc945e, 0x85f09f78, 0xd870e538, 0x8605988c, 0xb2bf17d1, 0x0b57e4a5, 0x4dd9a1e2, 0xf8c24e61, 0x457b42b3, 0xa5423421, 0xd625089c, 0x663cee1e, 0x52866143, 0xfc93b1c7, 0x2be54ffc, 0x14082404, 0x08a2e351, 0xc72f2599, 0xc4da226d, 0x391a650d, 0x35e979fa, 0x84a369df, 0x9bfca97e, 0xb4481924, 0xd776fe3b, 0x3d4b9aab, 0xd181f0ce, 0x55229911, 0x8903838f, 0x6b9c044e, 0x517366b7, 0x60cbe0eb, 0xcc78c13c, 0xbf1ffd81, 0xfe354094, 0x0cf31cf7, 0x676f18b9, 0x5f268b13, 0x9c58512c, 0xb8bb05d3, 0x5cd38ce7, 0xcbdc396e, 0xf395aac4, 0x0f061b03, 0x13acdc56, 0x49885e44, 0x9efea07f, 0x374f88a9, 0x8254672a, 0x6d6b0abb, 0xe29f87c1, 0x02a6f153, 0x8ba572dc, 0x2716530b, 0xd327019d, 0xc1d82b6c, 0xf562a431, 0xb9e8f374, 0x09f115f6, 0x438c4c46, 0x2645a5ac, 0x970fb589, 0x4428b414, 0x42dfbae1, 0x4e2ca616, 0xd274f73a, 0xd0d20669, 0x2d124109, 0xade0d770, 0x54716fb6, 0xb7bd1ed0, 0x7ec7d6ed, 0xdb85e2cc, 0x57846842, 0xc22d2c98, 0x0e55eda4, 0x88507528, 0x31b8865c, 0x3fed6bf8, 0xa411c286, }, { 0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb, 0x4ca113cb, 0xa9626d11, 0x08050209, 0x426e9e0d, 0xadee6c9b, 0x590451ff, 0xdebdb90c, 0xfb06f70e, 0xef80f296, 0x5fcede30, 0xfcef3f6d, 0xaa07a4f8, 0x27fdc047, 0x89766535, 0xaccd2b37, 0x048c018a, 0x71155bd2, 0x603c186c, 0xff8af684, 0xb5e16a80, 0xe8693af5, 0x5347ddb3, 0xf6acb321, 0x5eed999c, 0x6d965c43, 0x627a9629, 0xa321e15d, 0x8216aed5, 0xa8412abd, 0x9fb6eee8, 0xa5eb6e92, 0x7b56d79e, 0x8cd92313, 0xd317fd23, 0x6a7f9420, 0x9e95a944, 0xfa25b0a2, 0x06ca8fcf, 0x558d527c, 0x5022145a, 0xe14f7f50, 0x691a5dc9, 0x7fdad614, 0x5cab17d9, 0x8173673c, 0xd234ba8f, 0x80502090, 0xf303f507, 0x16c08bdd, 0xedc67cd3, 0x28110a2d, 0x1fe6ce78, 0x7353d597, 0x25bb4e02, 0x32588273, 0x2c9d0ba7, 0x510153f6, 0xcf94fab2, 0xdcfb3749, 0x8e9fad56, 0x8b30eb70, 0x2371c1cd, 0xc791f8bb, 0x17e3cc71, 0xa68ea77b, 0xb84b2eaf, 0x02468e45, 0x84dc211a, 0x1ec589d4, 0x75995a58, 0x9179632e, 0x381b0e3f, 0x012347ac, 0xea2fb4b0, 0x6cb51bef, 0x85ff66b6, 0x3ff2c65c, 0x100a0412, 0x39384993, 0xafa8e2de, 0x0ecf8dc6, 0xc87d32d1, 0x7270923b, 0x869aaf5f, 0xc31df931, 0x4b48dba8, 0xe22ab6b9, 0x34920dbc, 0xa4c8293e, 0x2dbe4c0b, 0x8dfa64bf, 0xe94a7d59, 0x1b6acff2, 0x78331e77, 0xe6a6b733, 0x74ba1df4, 0x997c6127, 0x26de87eb, 0xbde46889, 0x7a759032, 0xab24e354, 0xf78ff48d, 0xf4ea3d64, 0xc23ebe9d, 0x1da0403d, 0x67d5d00f, 0xd07234ca, 0x192c41b7, 0xc95e757d, 0x9a19a8ce, 0xece53b7f, 0x0daa442f, 0x07e9c863, 0xdb12ff2a, 0xbfa2e6cc, 0x905a2482, 0x3a5d807a, 0x40281048, 0x56e89b95, 0x337bc5df, 0x9690ab4d, 0x611f5fc0, 0x1c830791, 0xf5c97ac8, 0xccf1335b, 0x00000000, 0x36d483f9, 0x4587566e, 0x97b3ece1, 0x64b019e6, 0xfea9b128, 0xd87736c3, 0xc15b7774, 0x112943be, 0x77dfd41d, 0xba0da0ea, 0x124c8a57, 0xcb18fb38, 0x9df060ad, 0x2b74c3c4, 0xe5c37eda, 0x921caac7, 0x791059db, 0x0365c9e9, 0x0fecca6a, 0xb9686903, 0x65935e4a, 0x4ee79d8e, 0xbe81a160, 0xe06c38fc, 0xbb2ee746, 0x52649a1f, 0xe4e03976, 0x8fbceafa, 0x301e0c36, 0x249809ae, 0xf940794b, 0x6359d185, 0x70361c7e, 0xf8633ee7, 0x37f7c455, 0xeea3b53a, 0x29324d81, 0xc4f43152, 0x9b3aef62, 0x66f697a3, 0x35b14a10, 0xf220b2ab, 0x54ae15d0, 0xb7a7e4c5, 0xd5dd72ec, 0x5a619816, 0xca3bbc94, 0xe785f09f, 0xddd870e5, 0x14860598, 0xc6b2bf17, 0x410b57e4, 0x434dd9a1, 0x2ff8c24e, 0xf1457b42, 0x15a54234, 0x94d62508, 0xf0663cee, 0x22528661, 0x76fc93b1, 0xb32be54f, 0x20140824, 0xb208a2e3, 0xbcc72f25, 0x4fc4da22, 0x68391a65, 0x8335e979, 0xb684a369, 0xd79bfca9, 0x3db44819, 0xc5d776fe, 0x313d4b9a, 0x3ed181f0, 0x88552299, 0x0c890383, 0x4a6b9c04, 0xd1517366, 0x0b60cbe0, 0xfdcc78c1, 0x7cbf1ffd, 0xd4fe3540, 0xeb0cf31c, 0xa1676f18, 0x985f268b, 0x7d9c5851, 0xd6b8bb05, 0x6b5cd38c, 0x57cbdc39, 0x6ef395aa, 0x180f061b, 0x8a13acdc, 0x1a49885e, 0xdf9efea0, 0x21374f88, 0x4d825467, 0xb16d6b0a, 0x46e29f87, 0xa202a6f1, 0xae8ba572, 0x58271653, 0x9cd32701, 0x47c1d82b, 0x95f562a4, 0x87b9e8f3, 0xe309f115, 0x0a438c4c, 0x092645a5, 0x3c970fb5, 0xa04428b4, 0x5b42dfba, 0xb04e2ca6, 0xcdd274f7, 0x6fd0d206, 0x482d1241, 0xa7ade0d7, 0xd954716f, 0xceb7bd1e, 0x3b7ec7d6, 0x2edb85e2, 0x2a578468, 0xb4c22d2c, 0x490e55ed, 0x5d885075, 0xda31b886, 0x933fed6b, 0x44a411c2, }, { 0x18c07830, 0x2305af46, 0xc67ef991, 0xe8136fcd, 0x874ca113, 0xb8a9626d, 0x01080502, 0x4f426e9e, 0x36adee6c, 0xa6590451, 0xd2debdb9, 0xf5fb06f7, 0x79ef80f2, 0x6f5fcede, 0x91fcef3f, 0x52aa07a4, 0x6027fdc0, 0xbc897665, 0x9baccd2b, 0x8e048c01, 0xa371155b, 0x0c603c18, 0x7bff8af6, 0x35b5e16a, 0x1de8693a, 0xe05347dd, 0xd7f6acb3, 0xc25eed99, 0x2e6d965c, 0x4b627a96, 0xfea321e1, 0x578216ae, 0x15a8412a, 0x779fb6ee, 0x37a5eb6e, 0xe57b56d7, 0x9f8cd923, 0xf0d317fd, 0x4a6a7f94, 0xda9e95a9, 0x58fa25b0, 0xc906ca8f, 0x29558d52, 0x0a502214, 0xb1e14f7f, 0xa0691a5d, 0x6b7fdad6, 0x855cab17, 0xbd817367, 0x5dd234ba, 0x10805020, 0xf4f303f5, 0xcb16c08b, 0x3eedc67c, 0x0528110a, 0x671fe6ce, 0xe47353d5, 0x2725bb4e, 0x41325882, 0x8b2c9d0b, 0xa7510153, 0x7dcf94fa, 0x95dcfb37, 0xd88e9fad, 0xfb8b30eb, 0xee2371c1, 0x7cc791f8, 0x6617e3cc, 0xdda68ea7, 0x17b84b2e, 0x4702468e, 0x9e84dc21, 0xca1ec589, 0x2d75995a, 0xbf917963, 0x07381b0e, 0xad012347, 0x5aea2fb4, 0x836cb51b, 0x3385ff66, 0x633ff2c6, 0x02100a04, 0xaa393849, 0x71afa8e2, 0xc80ecf8d, 0x19c87d32, 0x49727092, 0xd9869aaf, 0xf2c31df9, 0xe34b48db, 0x5be22ab6, 0x8834920d, 0x9aa4c829, 0x262dbe4c, 0x328dfa64, 0xb0e94a7d, 0xe91b6acf, 0x0f78331e, 0xd5e6a6b7, 0x8074ba1d, 0xbe997c61, 0xcd26de87, 0x34bde468, 0x487a7590, 0xffab24e3, 0x7af78ff4, 0x90f4ea3d, 0x5fc23ebe, 0x201da040, 0x6867d5d0, 0x1ad07234, 0xae192c41, 0xb4c95e75, 0x549a19a8, 0x93ece53b, 0x220daa44, 0x6407e9c8, 0xf1db12ff, 0x73bfa2e6, 0x12905a24, 0x403a5d80, 0x08402810, 0xc356e89b, 0xec337bc5, 0xdb9690ab, 0xa1611f5f, 0x8d1c8307, 0x3df5c97a, 0x97ccf133, 0x00000000, 0xcf36d483, 0x2b458756, 0x7697b3ec, 0x8264b019, 0xd6fea9b1, 0x1bd87736, 0xb5c15b77, 0xaf112943, 0x6a77dfd4, 0x50ba0da0, 0x45124c8a, 0xf3cb18fb, 0x309df060, 0xef2b74c3, 0x3fe5c37e, 0x55921caa, 0xa2791059, 0xea0365c9, 0x650fecca, 0xbab96869, 0x2f65935e, 0xc04ee79d, 0xdebe81a1, 0x1ce06c38, 0xfdbb2ee7, 0x4d52649a, 0x92e4e039, 0x758fbcea, 0x06301e0c, 0x8a249809, 0xb2f94079, 0xe66359d1, 0x0e70361c, 0x1ff8633e, 0x6237f7c4, 0xd4eea3b5, 0xa829324d, 0x96c4f431, 0xf99b3aef, 0xc566f697, 0x2535b14a, 0x59f220b2, 0x8454ae15, 0x72b7a7e4, 0x39d5dd72, 0x4c5a6198, 0x5eca3bbc, 0x78e785f0, 0x38ddd870, 0x8c148605, 0xd1c6b2bf, 0xa5410b57, 0xe2434dd9, 0x612ff8c2, 0xb3f1457b, 0x2115a542, 0x9c94d625, 0x1ef0663c, 0x43225286, 0xc776fc93, 0xfcb32be5, 0x04201408, 0x51b208a2, 0x99bcc72f, 0x6d4fc4da, 0x0d68391a, 0xfa8335e9, 0xdfb684a3, 0x7ed79bfc, 0x243db448, 0x3bc5d776, 0xab313d4b, 0xce3ed181, 0x11885522, 0x8f0c8903, 0x4e4a6b9c, 0xb7d15173, 0xeb0b60cb, 0x3cfdcc78, 0x817cbf1f, 0x94d4fe35, 0xf7eb0cf3, 0xb9a1676f, 0x13985f26, 0x2c7d9c58, 0xd3d6b8bb, 0xe76b5cd3, 0x6e57cbdc, 0xc46ef395, 0x03180f06, 0x568a13ac, 0x441a4988, 0x7fdf9efe, 0xa921374f, 0x2a4d8254, 0xbbb16d6b, 0xc146e29f, 0x53a202a6, 0xdcae8ba5, 0x0b582716, 0x9d9cd327, 0x6c47c1d8, 0x3195f562, 0x7487b9e8, 0xf6e309f1, 0x460a438c, 0xac092645, 0x893c970f, 0x14a04428, 0xe15b42df, 0x16b04e2c, 0x3acdd274, 0x696fd0d2, 0x09482d12, 0x70a7ade0, 0xb6d95471, 0xd0ceb7bd, 0xed3b7ec7, 0xcc2edb85, 0x422a5784, 0x98b4c22d, 0xa4490e55, 0x285d8850, 0x5cda31b8, 0xf8933fed, 0x8644a411, }, { 0x6018c078, 0x8c2305af, 0x3fc67ef9, 0x87e8136f, 0x26874ca1, 0xdab8a962, 0x04010805, 0x214f426e, 0xd836adee, 0xa2a65904, 0x6fd2debd, 0xf3f5fb06, 0xf979ef80, 0xa16f5fce, 0x7e91fcef, 0x5552aa07, 0x9d6027fd, 0xcabc8976, 0x569baccd, 0x028e048c, 0xb6a37115, 0x300c603c, 0xf17bff8a, 0xd435b5e1, 0x741de869, 0xa7e05347, 0x7bd7f6ac, 0x2fc25eed, 0xb82e6d96, 0x314b627a, 0xdffea321, 0x41578216, 0x5415a841, 0xc1779fb6, 0xdc37a5eb, 0xb3e57b56, 0x469f8cd9, 0xe7f0d317, 0x354a6a7f, 0x4fda9e95, 0x7d58fa25, 0x03c906ca, 0xa429558d, 0x280a5022, 0xfeb1e14f, 0xbaa0691a, 0xb16b7fda, 0x2e855cab, 0xcebd8173, 0x695dd234, 0x40108050, 0xf7f4f303, 0x0bcb16c0, 0xf83eedc6, 0x14052811, 0x81671fe6, 0xb7e47353, 0x9c2725bb, 0x19413258, 0x168b2c9d, 0xa6a75101, 0xe97dcf94, 0x6e95dcfb, 0x47d88e9f, 0xcbfb8b30, 0x9fee2371, 0xed7cc791, 0x856617e3, 0x53dda68e, 0x5c17b84b, 0x01470246, 0x429e84dc, 0x0fca1ec5, 0xb42d7599, 0xc6bf9179, 0x1c07381b, 0x8ead0123, 0x755aea2f, 0x36836cb5, 0xcc3385ff, 0x91633ff2, 0x0802100a, 0x92aa3938, 0xd971afa8, 0x07c80ecf, 0x6419c87d, 0x39497270, 0x43d9869a, 0xeff2c31d, 0xabe34b48, 0x715be22a, 0x1a883492, 0x529aa4c8, 0x98262dbe, 0xc8328dfa, 0xfab0e94a, 0x83e91b6a, 0x3c0f7833, 0x73d5e6a6, 0x3a8074ba, 0xc2be997c, 0x13cd26de, 0xd034bde4, 0x3d487a75, 0xdbffab24, 0xf57af78f, 0x7a90f4ea, 0x615fc23e, 0x80201da0, 0xbd6867d5, 0x681ad072, 0x82ae192c, 0xeab4c95e, 0x4d549a19, 0x7693ece5, 0x88220daa, 0x8d6407e9, 0xe3f1db12, 0xd173bfa2, 0x4812905a, 0x1d403a5d, 0x20084028, 0x2bc356e8, 0x97ec337b, 0x4bdb9690, 0xbea1611f, 0x0e8d1c83, 0xf43df5c9, 0x6697ccf1, 0x00000000, 0x1bcf36d4, 0xac2b4587, 0xc57697b3, 0x328264b0, 0x7fd6fea9, 0x6c1bd877, 0xeeb5c15b, 0x86af1129, 0xb56a77df, 0x5d50ba0d, 0x0945124c, 0xebf3cb18, 0xc0309df0, 0x9bef2b74, 0xfc3fe5c3, 0x4955921c, 0xb2a27910, 0x8fea0365, 0x89650fec, 0xd2bab968, 0xbc2f6593, 0x27c04ee7, 0x5fdebe81, 0x701ce06c, 0xd3fdbb2e, 0x294d5264, 0x7292e4e0, 0xc9758fbc, 0x1806301e, 0x128a2498, 0xf2b2f940, 0xbfe66359, 0x380e7036, 0x7c1ff863, 0x956237f7, 0x77d4eea3, 0x9aa82932, 0x6296c4f4, 0xc3f99b3a, 0x33c566f6, 0x942535b1, 0x7959f220, 0x2a8454ae, 0xd572b7a7, 0xe439d5dd, 0x2d4c5a61, 0x655eca3b, 0xfd78e785, 0xe038ddd8, 0x0a8c1486, 0x63d1c6b2, 0xaea5410b, 0xafe2434d, 0x99612ff8, 0xf6b3f145, 0x842115a5, 0x4a9c94d6, 0x781ef066, 0x11432252, 0x3bc776fc, 0xd7fcb32b, 0x10042014, 0x5951b208, 0x5e99bcc7, 0xa96d4fc4, 0x340d6839, 0xcffa8335, 0x5bdfb684, 0xe57ed79b, 0x90243db4, 0xec3bc5d7, 0x96ab313d, 0x1fce3ed1, 0x44118855, 0x068f0c89, 0x254e4a6b, 0xe6b7d151, 0x8beb0b60, 0xf03cfdcc, 0x3e817cbf, 0x6a94d4fe, 0xfbf7eb0c, 0xdeb9a167, 0x4c13985f, 0xb02c7d9c, 0x6bd3d6b8, 0xbbe76b5c, 0xa56e57cb, 0x37c46ef3, 0x0c03180f, 0x45568a13, 0x0d441a49, 0xe17fdf9e, 0x9ea92137, 0xa82a4d82, 0xd6bbb16d, 0x23c146e2, 0x5153a202, 0x57dcae8b, 0x2c0b5827, 0x4e9d9cd3, 0xad6c47c1, 0xc43195f5, 0xcd7487b9, 0xfff6e309, 0x05460a43, 0x8aac0926, 0x1e893c97, 0x5014a044, 0xa3e15b42, 0x5816b04e, 0xe83acdd2, 0xb9696fd0, 0x2409482d, 0xdd70a7ad, 0xe2b6d954, 0x67d0ceb7, 0x93ed3b7e, 0x17cc2edb, 0x15422a57, 0x5a98b4c2, 0xaaa4490e, 0xa0285d88, 0x6d5cda31, 0xc7f8933f, 0x228644a4, }, { 0x186018c0, 0x238c2305, 0xc63fc67e, 0xe887e813, 0x8726874c, 0xb8dab8a9, 0x01040108, 0x4f214f42, 0x36d836ad, 0xa6a2a659, 0xd26fd2de, 0xf5f3f5fb, 0x79f979ef, 0x6fa16f5f, 0x917e91fc, 0x525552aa, 0x609d6027, 0xbccabc89, 0x9b569bac, 0x8e028e04, 0xa3b6a371, 0x0c300c60, 0x7bf17bff, 0x35d435b5, 0x1d741de8, 0xe0a7e053, 0xd77bd7f6, 0xc22fc25e, 0x2eb82e6d, 0x4b314b62, 0xfedffea3, 0x57415782, 0x155415a8, 0x77c1779f, 0x37dc37a5, 0xe5b3e57b, 0x9f469f8c, 0xf0e7f0d3, 0x4a354a6a, 0xda4fda9e, 0x587d58fa, 0xc903c906, 0x29a42955, 0x0a280a50, 0xb1feb1e1, 0xa0baa069, 0x6bb16b7f, 0x852e855c, 0xbdcebd81, 0x5d695dd2, 0x10401080, 0xf4f7f4f3, 0xcb0bcb16, 0x3ef83eed, 0x05140528, 0x6781671f, 0xe4b7e473, 0x279c2725, 0x41194132, 0x8b168b2c, 0xa7a6a751, 0x7de97dcf, 0x956e95dc, 0xd847d88e, 0xfbcbfb8b, 0xee9fee23, 0x7ced7cc7, 0x66856617, 0xdd53dda6, 0x175c17b8, 0x47014702, 0x9e429e84, 0xca0fca1e, 0x2db42d75, 0xbfc6bf91, 0x071c0738, 0xad8ead01, 0x5a755aea, 0x8336836c, 0x33cc3385, 0x6391633f, 0x02080210, 0xaa92aa39, 0x71d971af, 0xc807c80e, 0x196419c8, 0x49394972, 0xd943d986, 0xf2eff2c3, 0xe3abe34b, 0x5b715be2, 0x881a8834, 0x9a529aa4, 0x2698262d, 0x32c8328d, 0xb0fab0e9, 0xe983e91b, 0x0f3c0f78, 0xd573d5e6, 0x803a8074, 0xbec2be99, 0xcd13cd26, 0x34d034bd, 0x483d487a, 0xffdbffab, 0x7af57af7, 0x907a90f4, 0x5f615fc2, 0x2080201d, 0x68bd6867, 0x1a681ad0, 0xae82ae19, 0xb4eab4c9, 0x544d549a, 0x937693ec, 0x2288220d, 0x648d6407, 0xf1e3f1db, 0x73d173bf, 0x12481290, 0x401d403a, 0x08200840, 0xc32bc356, 0xec97ec33, 0xdb4bdb96, 0xa1bea161, 0x8d0e8d1c, 0x3df43df5, 0x976697cc, 0x00000000, 0xcf1bcf36, 0x2bac2b45, 0x76c57697, 0x82328264, 0xd67fd6fe, 0x1b6c1bd8, 0xb5eeb5c1, 0xaf86af11, 0x6ab56a77, 0x505d50ba, 0x45094512, 0xf3ebf3cb, 0x30c0309d, 0xef9bef2b, 0x3ffc3fe5, 0x55495592, 0xa2b2a279, 0xea8fea03, 0x6589650f, 0xbad2bab9, 0x2fbc2f65, 0xc027c04e, 0xde5fdebe, 0x1c701ce0, 0xfdd3fdbb, 0x4d294d52, 0x927292e4, 0x75c9758f, 0x06180630, 0x8a128a24, 0xb2f2b2f9, 0xe6bfe663, 0x0e380e70, 0x1f7c1ff8, 0x62956237, 0xd477d4ee, 0xa89aa829, 0x966296c4, 0xf9c3f99b, 0xc533c566, 0x25942535, 0x597959f2, 0x842a8454, 0x72d572b7, 0x39e439d5, 0x4c2d4c5a, 0x5e655eca, 0x78fd78e7, 0x38e038dd, 0x8c0a8c14, 0xd163d1c6, 0xa5aea541, 0xe2afe243, 0x6199612f, 0xb3f6b3f1, 0x21842115, 0x9c4a9c94, 0x1e781ef0, 0x43114322, 0xc73bc776, 0xfcd7fcb3, 0x04100420, 0x515951b2, 0x995e99bc, 0x6da96d4f, 0x0d340d68, 0xfacffa83, 0xdf5bdfb6, 0x7ee57ed7, 0x2490243d, 0x3bec3bc5, 0xab96ab31, 0xce1fce3e, 0x11441188, 0x8f068f0c, 0x4e254e4a, 0xb7e6b7d1, 0xeb8beb0b, 0x3cf03cfd, 0x813e817c, 0x946a94d4, 0xf7fbf7eb, 0xb9deb9a1, 0x134c1398, 0x2cb02c7d, 0xd36bd3d6, 0xe7bbe76b, 0x6ea56e57, 0xc437c46e, 0x030c0318, 0x5645568a, 0x440d441a, 0x7fe17fdf, 0xa99ea921, 0x2aa82a4d, 0xbbd6bbb1, 0xc123c146, 0x535153a2, 0xdc57dcae, 0x0b2c0b58, 0x9d4e9d9c, 0x6cad6c47, 0x31c43195, 0x74cd7487, 0xf6fff6e3, 0x4605460a, 0xac8aac09, 0x891e893c, 0x145014a0, 0xe1a3e15b, 0x165816b0, 0x3ae83acd, 0x69b9696f, 0x09240948, 0x70dd70a7, 0xb6e2b6d9, 0xd067d0ce, 0xed93ed3b, 0xcc17cc2e, 0x4215422a, 0x985a98b4, 0xa4aaa449, 0x28a0285d, 0x5c6d5cda, 0xf8c7f893, 0x86228644, } }; __constant u32 Cl[8][256] = { { 0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb, 0x4ca113cb, 0xa9626d11, 0x08050209, 0x426e9e0d, 0xadee6c9b, 0x590451ff, 0xdebdb90c, 0xfb06f70e, 0xef80f296, 0x5fcede30, 0xfcef3f6d, 0xaa07a4f8, 0x27fdc047, 0x89766535, 0xaccd2b37, 0x048c018a, 0x71155bd2, 0x603c186c, 0xff8af684, 0xb5e16a80, 0xe8693af5, 0x5347ddb3, 0xf6acb321, 0x5eed999c, 0x6d965c43, 0x627a9629, 0xa321e15d, 0x8216aed5, 0xa8412abd, 0x9fb6eee8, 0xa5eb6e92, 0x7b56d79e, 0x8cd92313, 0xd317fd23, 0x6a7f9420, 0x9e95a944, 0xfa25b0a2, 0x06ca8fcf, 0x558d527c, 0x5022145a, 0xe14f7f50, 0x691a5dc9, 0x7fdad614, 0x5cab17d9, 0x8173673c, 0xd234ba8f, 0x80502090, 0xf303f507, 0x16c08bdd, 0xedc67cd3, 0x28110a2d, 0x1fe6ce78, 0x7353d597, 0x25bb4e02, 0x32588273, 0x2c9d0ba7, 0x510153f6, 0xcf94fab2, 0xdcfb3749, 0x8e9fad56, 0x8b30eb70, 0x2371c1cd, 0xc791f8bb, 0x17e3cc71, 0xa68ea77b, 0xb84b2eaf, 0x02468e45, 0x84dc211a, 0x1ec589d4, 0x75995a58, 0x9179632e, 0x381b0e3f, 0x012347ac, 0xea2fb4b0, 0x6cb51bef, 0x85ff66b6, 0x3ff2c65c, 0x100a0412, 0x39384993, 0xafa8e2de, 0x0ecf8dc6, 0xc87d32d1, 0x7270923b, 0x869aaf5f, 0xc31df931, 0x4b48dba8, 0xe22ab6b9, 0x34920dbc, 0xa4c8293e, 0x2dbe4c0b, 0x8dfa64bf, 0xe94a7d59, 0x1b6acff2, 0x78331e77, 0xe6a6b733, 0x74ba1df4, 0x997c6127, 0x26de87eb, 0xbde46889, 0x7a759032, 0xab24e354, 0xf78ff48d, 0xf4ea3d64, 0xc23ebe9d, 0x1da0403d, 0x67d5d00f, 0xd07234ca, 0x192c41b7, 0xc95e757d, 0x9a19a8ce, 0xece53b7f, 0x0daa442f, 0x07e9c863, 0xdb12ff2a, 0xbfa2e6cc, 0x905a2482, 0x3a5d807a, 0x40281048, 0x56e89b95, 0x337bc5df, 0x9690ab4d, 0x611f5fc0, 0x1c830791, 0xf5c97ac8, 0xccf1335b, 0x00000000, 0x36d483f9, 0x4587566e, 0x97b3ece1, 0x64b019e6, 0xfea9b128, 0xd87736c3, 0xc15b7774, 0x112943be, 0x77dfd41d, 0xba0da0ea, 0x124c8a57, 0xcb18fb38, 0x9df060ad, 0x2b74c3c4, 0xe5c37eda, 0x921caac7, 0x791059db, 0x0365c9e9, 0x0fecca6a, 0xb9686903, 0x65935e4a, 0x4ee79d8e, 0xbe81a160, 0xe06c38fc, 0xbb2ee746, 0x52649a1f, 0xe4e03976, 0x8fbceafa, 0x301e0c36, 0x249809ae, 0xf940794b, 0x6359d185, 0x70361c7e, 0xf8633ee7, 0x37f7c455, 0xeea3b53a, 0x29324d81, 0xc4f43152, 0x9b3aef62, 0x66f697a3, 0x35b14a10, 0xf220b2ab, 0x54ae15d0, 0xb7a7e4c5, 0xd5dd72ec, 0x5a619816, 0xca3bbc94, 0xe785f09f, 0xddd870e5, 0x14860598, 0xc6b2bf17, 0x410b57e4, 0x434dd9a1, 0x2ff8c24e, 0xf1457b42, 0x15a54234, 0x94d62508, 0xf0663cee, 0x22528661, 0x76fc93b1, 0xb32be54f, 0x20140824, 0xb208a2e3, 0xbcc72f25, 0x4fc4da22, 0x68391a65, 0x8335e979, 0xb684a369, 0xd79bfca9, 0x3db44819, 0xc5d776fe, 0x313d4b9a, 0x3ed181f0, 0x88552299, 0x0c890383, 0x4a6b9c04, 0xd1517366, 0x0b60cbe0, 0xfdcc78c1, 0x7cbf1ffd, 0xd4fe3540, 0xeb0cf31c, 0xa1676f18, 0x985f268b, 0x7d9c5851, 0xd6b8bb05, 0x6b5cd38c, 0x57cbdc39, 0x6ef395aa, 0x180f061b, 0x8a13acdc, 0x1a49885e, 0xdf9efea0, 0x21374f88, 0x4d825467, 0xb16d6b0a, 0x46e29f87, 0xa202a6f1, 0xae8ba572, 0x58271653, 0x9cd32701, 0x47c1d82b, 0x95f562a4, 0x87b9e8f3, 0xe309f115, 0x0a438c4c, 0x092645a5, 0x3c970fb5, 0xa04428b4, 0x5b42dfba, 0xb04e2ca6, 0xcdd274f7, 0x6fd0d206, 0x482d1241, 0xa7ade0d7, 0xd954716f, 0xceb7bd1e, 0x3b7ec7d6, 0x2edb85e2, 0x2a578468, 0xb4c22d2c, 0x490e55ed, 0x5d885075, 0xda31b886, 0x933fed6b, 0x44a411c2, }, { 0x18c07830, 0x2305af46, 0xc67ef991, 0xe8136fcd, 0x874ca113, 0xb8a9626d, 0x01080502, 0x4f426e9e, 0x36adee6c, 0xa6590451, 0xd2debdb9, 0xf5fb06f7, 0x79ef80f2, 0x6f5fcede, 0x91fcef3f, 0x52aa07a4, 0x6027fdc0, 0xbc897665, 0x9baccd2b, 0x8e048c01, 0xa371155b, 0x0c603c18, 0x7bff8af6, 0x35b5e16a, 0x1de8693a, 0xe05347dd, 0xd7f6acb3, 0xc25eed99, 0x2e6d965c, 0x4b627a96, 0xfea321e1, 0x578216ae, 0x15a8412a, 0x779fb6ee, 0x37a5eb6e, 0xe57b56d7, 0x9f8cd923, 0xf0d317fd, 0x4a6a7f94, 0xda9e95a9, 0x58fa25b0, 0xc906ca8f, 0x29558d52, 0x0a502214, 0xb1e14f7f, 0xa0691a5d, 0x6b7fdad6, 0x855cab17, 0xbd817367, 0x5dd234ba, 0x10805020, 0xf4f303f5, 0xcb16c08b, 0x3eedc67c, 0x0528110a, 0x671fe6ce, 0xe47353d5, 0x2725bb4e, 0x41325882, 0x8b2c9d0b, 0xa7510153, 0x7dcf94fa, 0x95dcfb37, 0xd88e9fad, 0xfb8b30eb, 0xee2371c1, 0x7cc791f8, 0x6617e3cc, 0xdda68ea7, 0x17b84b2e, 0x4702468e, 0x9e84dc21, 0xca1ec589, 0x2d75995a, 0xbf917963, 0x07381b0e, 0xad012347, 0x5aea2fb4, 0x836cb51b, 0x3385ff66, 0x633ff2c6, 0x02100a04, 0xaa393849, 0x71afa8e2, 0xc80ecf8d, 0x19c87d32, 0x49727092, 0xd9869aaf, 0xf2c31df9, 0xe34b48db, 0x5be22ab6, 0x8834920d, 0x9aa4c829, 0x262dbe4c, 0x328dfa64, 0xb0e94a7d, 0xe91b6acf, 0x0f78331e, 0xd5e6a6b7, 0x8074ba1d, 0xbe997c61, 0xcd26de87, 0x34bde468, 0x487a7590, 0xffab24e3, 0x7af78ff4, 0x90f4ea3d, 0x5fc23ebe, 0x201da040, 0x6867d5d0, 0x1ad07234, 0xae192c41, 0xb4c95e75, 0x549a19a8, 0x93ece53b, 0x220daa44, 0x6407e9c8, 0xf1db12ff, 0x73bfa2e6, 0x12905a24, 0x403a5d80, 0x08402810, 0xc356e89b, 0xec337bc5, 0xdb9690ab, 0xa1611f5f, 0x8d1c8307, 0x3df5c97a, 0x97ccf133, 0x00000000, 0xcf36d483, 0x2b458756, 0x7697b3ec, 0x8264b019, 0xd6fea9b1, 0x1bd87736, 0xb5c15b77, 0xaf112943, 0x6a77dfd4, 0x50ba0da0, 0x45124c8a, 0xf3cb18fb, 0x309df060, 0xef2b74c3, 0x3fe5c37e, 0x55921caa, 0xa2791059, 0xea0365c9, 0x650fecca, 0xbab96869, 0x2f65935e, 0xc04ee79d, 0xdebe81a1, 0x1ce06c38, 0xfdbb2ee7, 0x4d52649a, 0x92e4e039, 0x758fbcea, 0x06301e0c, 0x8a249809, 0xb2f94079, 0xe66359d1, 0x0e70361c, 0x1ff8633e, 0x6237f7c4, 0xd4eea3b5, 0xa829324d, 0x96c4f431, 0xf99b3aef, 0xc566f697, 0x2535b14a, 0x59f220b2, 0x8454ae15, 0x72b7a7e4, 0x39d5dd72, 0x4c5a6198, 0x5eca3bbc, 0x78e785f0, 0x38ddd870, 0x8c148605, 0xd1c6b2bf, 0xa5410b57, 0xe2434dd9, 0x612ff8c2, 0xb3f1457b, 0x2115a542, 0x9c94d625, 0x1ef0663c, 0x43225286, 0xc776fc93, 0xfcb32be5, 0x04201408, 0x51b208a2, 0x99bcc72f, 0x6d4fc4da, 0x0d68391a, 0xfa8335e9, 0xdfb684a3, 0x7ed79bfc, 0x243db448, 0x3bc5d776, 0xab313d4b, 0xce3ed181, 0x11885522, 0x8f0c8903, 0x4e4a6b9c, 0xb7d15173, 0xeb0b60cb, 0x3cfdcc78, 0x817cbf1f, 0x94d4fe35, 0xf7eb0cf3, 0xb9a1676f, 0x13985f26, 0x2c7d9c58, 0xd3d6b8bb, 0xe76b5cd3, 0x6e57cbdc, 0xc46ef395, 0x03180f06, 0x568a13ac, 0x441a4988, 0x7fdf9efe, 0xa921374f, 0x2a4d8254, 0xbbb16d6b, 0xc146e29f, 0x53a202a6, 0xdcae8ba5, 0x0b582716, 0x9d9cd327, 0x6c47c1d8, 0x3195f562, 0x7487b9e8, 0xf6e309f1, 0x460a438c, 0xac092645, 0x893c970f, 0x14a04428, 0xe15b42df, 0x16b04e2c, 0x3acdd274, 0x696fd0d2, 0x09482d12, 0x70a7ade0, 0xb6d95471, 0xd0ceb7bd, 0xed3b7ec7, 0xcc2edb85, 0x422a5784, 0x98b4c22d, 0xa4490e55, 0x285d8850, 0x5cda31b8, 0xf8933fed, 0x8644a411, }, { 0x6018c078, 0x8c2305af, 0x3fc67ef9, 0x87e8136f, 0x26874ca1, 0xdab8a962, 0x04010805, 0x214f426e, 0xd836adee, 0xa2a65904, 0x6fd2debd, 0xf3f5fb06, 0xf979ef80, 0xa16f5fce, 0x7e91fcef, 0x5552aa07, 0x9d6027fd, 0xcabc8976, 0x569baccd, 0x028e048c, 0xb6a37115, 0x300c603c, 0xf17bff8a, 0xd435b5e1, 0x741de869, 0xa7e05347, 0x7bd7f6ac, 0x2fc25eed, 0xb82e6d96, 0x314b627a, 0xdffea321, 0x41578216, 0x5415a841, 0xc1779fb6, 0xdc37a5eb, 0xb3e57b56, 0x469f8cd9, 0xe7f0d317, 0x354a6a7f, 0x4fda9e95, 0x7d58fa25, 0x03c906ca, 0xa429558d, 0x280a5022, 0xfeb1e14f, 0xbaa0691a, 0xb16b7fda, 0x2e855cab, 0xcebd8173, 0x695dd234, 0x40108050, 0xf7f4f303, 0x0bcb16c0, 0xf83eedc6, 0x14052811, 0x81671fe6, 0xb7e47353, 0x9c2725bb, 0x19413258, 0x168b2c9d, 0xa6a75101, 0xe97dcf94, 0x6e95dcfb, 0x47d88e9f, 0xcbfb8b30, 0x9fee2371, 0xed7cc791, 0x856617e3, 0x53dda68e, 0x5c17b84b, 0x01470246, 0x429e84dc, 0x0fca1ec5, 0xb42d7599, 0xc6bf9179, 0x1c07381b, 0x8ead0123, 0x755aea2f, 0x36836cb5, 0xcc3385ff, 0x91633ff2, 0x0802100a, 0x92aa3938, 0xd971afa8, 0x07c80ecf, 0x6419c87d, 0x39497270, 0x43d9869a, 0xeff2c31d, 0xabe34b48, 0x715be22a, 0x1a883492, 0x529aa4c8, 0x98262dbe, 0xc8328dfa, 0xfab0e94a, 0x83e91b6a, 0x3c0f7833, 0x73d5e6a6, 0x3a8074ba, 0xc2be997c, 0x13cd26de, 0xd034bde4, 0x3d487a75, 0xdbffab24, 0xf57af78f, 0x7a90f4ea, 0x615fc23e, 0x80201da0, 0xbd6867d5, 0x681ad072, 0x82ae192c, 0xeab4c95e, 0x4d549a19, 0x7693ece5, 0x88220daa, 0x8d6407e9, 0xe3f1db12, 0xd173bfa2, 0x4812905a, 0x1d403a5d, 0x20084028, 0x2bc356e8, 0x97ec337b, 0x4bdb9690, 0xbea1611f, 0x0e8d1c83, 0xf43df5c9, 0x6697ccf1, 0x00000000, 0x1bcf36d4, 0xac2b4587, 0xc57697b3, 0x328264b0, 0x7fd6fea9, 0x6c1bd877, 0xeeb5c15b, 0x86af1129, 0xb56a77df, 0x5d50ba0d, 0x0945124c, 0xebf3cb18, 0xc0309df0, 0x9bef2b74, 0xfc3fe5c3, 0x4955921c, 0xb2a27910, 0x8fea0365, 0x89650fec, 0xd2bab968, 0xbc2f6593, 0x27c04ee7, 0x5fdebe81, 0x701ce06c, 0xd3fdbb2e, 0x294d5264, 0x7292e4e0, 0xc9758fbc, 0x1806301e, 0x128a2498, 0xf2b2f940, 0xbfe66359, 0x380e7036, 0x7c1ff863, 0x956237f7, 0x77d4eea3, 0x9aa82932, 0x6296c4f4, 0xc3f99b3a, 0x33c566f6, 0x942535b1, 0x7959f220, 0x2a8454ae, 0xd572b7a7, 0xe439d5dd, 0x2d4c5a61, 0x655eca3b, 0xfd78e785, 0xe038ddd8, 0x0a8c1486, 0x63d1c6b2, 0xaea5410b, 0xafe2434d, 0x99612ff8, 0xf6b3f145, 0x842115a5, 0x4a9c94d6, 0x781ef066, 0x11432252, 0x3bc776fc, 0xd7fcb32b, 0x10042014, 0x5951b208, 0x5e99bcc7, 0xa96d4fc4, 0x340d6839, 0xcffa8335, 0x5bdfb684, 0xe57ed79b, 0x90243db4, 0xec3bc5d7, 0x96ab313d, 0x1fce3ed1, 0x44118855, 0x068f0c89, 0x254e4a6b, 0xe6b7d151, 0x8beb0b60, 0xf03cfdcc, 0x3e817cbf, 0x6a94d4fe, 0xfbf7eb0c, 0xdeb9a167, 0x4c13985f, 0xb02c7d9c, 0x6bd3d6b8, 0xbbe76b5c, 0xa56e57cb, 0x37c46ef3, 0x0c03180f, 0x45568a13, 0x0d441a49, 0xe17fdf9e, 0x9ea92137, 0xa82a4d82, 0xd6bbb16d, 0x23c146e2, 0x5153a202, 0x57dcae8b, 0x2c0b5827, 0x4e9d9cd3, 0xad6c47c1, 0xc43195f5, 0xcd7487b9, 0xfff6e309, 0x05460a43, 0x8aac0926, 0x1e893c97, 0x5014a044, 0xa3e15b42, 0x5816b04e, 0xe83acdd2, 0xb9696fd0, 0x2409482d, 0xdd70a7ad, 0xe2b6d954, 0x67d0ceb7, 0x93ed3b7e, 0x17cc2edb, 0x15422a57, 0x5a98b4c2, 0xaaa4490e, 0xa0285d88, 0x6d5cda31, 0xc7f8933f, 0x228644a4, }, { 0x186018c0, 0x238c2305, 0xc63fc67e, 0xe887e813, 0x8726874c, 0xb8dab8a9, 0x01040108, 0x4f214f42, 0x36d836ad, 0xa6a2a659, 0xd26fd2de, 0xf5f3f5fb, 0x79f979ef, 0x6fa16f5f, 0x917e91fc, 0x525552aa, 0x609d6027, 0xbccabc89, 0x9b569bac, 0x8e028e04, 0xa3b6a371, 0x0c300c60, 0x7bf17bff, 0x35d435b5, 0x1d741de8, 0xe0a7e053, 0xd77bd7f6, 0xc22fc25e, 0x2eb82e6d, 0x4b314b62, 0xfedffea3, 0x57415782, 0x155415a8, 0x77c1779f, 0x37dc37a5, 0xe5b3e57b, 0x9f469f8c, 0xf0e7f0d3, 0x4a354a6a, 0xda4fda9e, 0x587d58fa, 0xc903c906, 0x29a42955, 0x0a280a50, 0xb1feb1e1, 0xa0baa069, 0x6bb16b7f, 0x852e855c, 0xbdcebd81, 0x5d695dd2, 0x10401080, 0xf4f7f4f3, 0xcb0bcb16, 0x3ef83eed, 0x05140528, 0x6781671f, 0xe4b7e473, 0x279c2725, 0x41194132, 0x8b168b2c, 0xa7a6a751, 0x7de97dcf, 0x956e95dc, 0xd847d88e, 0xfbcbfb8b, 0xee9fee23, 0x7ced7cc7, 0x66856617, 0xdd53dda6, 0x175c17b8, 0x47014702, 0x9e429e84, 0xca0fca1e, 0x2db42d75, 0xbfc6bf91, 0x071c0738, 0xad8ead01, 0x5a755aea, 0x8336836c, 0x33cc3385, 0x6391633f, 0x02080210, 0xaa92aa39, 0x71d971af, 0xc807c80e, 0x196419c8, 0x49394972, 0xd943d986, 0xf2eff2c3, 0xe3abe34b, 0x5b715be2, 0x881a8834, 0x9a529aa4, 0x2698262d, 0x32c8328d, 0xb0fab0e9, 0xe983e91b, 0x0f3c0f78, 0xd573d5e6, 0x803a8074, 0xbec2be99, 0xcd13cd26, 0x34d034bd, 0x483d487a, 0xffdbffab, 0x7af57af7, 0x907a90f4, 0x5f615fc2, 0x2080201d, 0x68bd6867, 0x1a681ad0, 0xae82ae19, 0xb4eab4c9, 0x544d549a, 0x937693ec, 0x2288220d, 0x648d6407, 0xf1e3f1db, 0x73d173bf, 0x12481290, 0x401d403a, 0x08200840, 0xc32bc356, 0xec97ec33, 0xdb4bdb96, 0xa1bea161, 0x8d0e8d1c, 0x3df43df5, 0x976697cc, 0x00000000, 0xcf1bcf36, 0x2bac2b45, 0x76c57697, 0x82328264, 0xd67fd6fe, 0x1b6c1bd8, 0xb5eeb5c1, 0xaf86af11, 0x6ab56a77, 0x505d50ba, 0x45094512, 0xf3ebf3cb, 0x30c0309d, 0xef9bef2b, 0x3ffc3fe5, 0x55495592, 0xa2b2a279, 0xea8fea03, 0x6589650f, 0xbad2bab9, 0x2fbc2f65, 0xc027c04e, 0xde5fdebe, 0x1c701ce0, 0xfdd3fdbb, 0x4d294d52, 0x927292e4, 0x75c9758f, 0x06180630, 0x8a128a24, 0xb2f2b2f9, 0xe6bfe663, 0x0e380e70, 0x1f7c1ff8, 0x62956237, 0xd477d4ee, 0xa89aa829, 0x966296c4, 0xf9c3f99b, 0xc533c566, 0x25942535, 0x597959f2, 0x842a8454, 0x72d572b7, 0x39e439d5, 0x4c2d4c5a, 0x5e655eca, 0x78fd78e7, 0x38e038dd, 0x8c0a8c14, 0xd163d1c6, 0xa5aea541, 0xe2afe243, 0x6199612f, 0xb3f6b3f1, 0x21842115, 0x9c4a9c94, 0x1e781ef0, 0x43114322, 0xc73bc776, 0xfcd7fcb3, 0x04100420, 0x515951b2, 0x995e99bc, 0x6da96d4f, 0x0d340d68, 0xfacffa83, 0xdf5bdfb6, 0x7ee57ed7, 0x2490243d, 0x3bec3bc5, 0xab96ab31, 0xce1fce3e, 0x11441188, 0x8f068f0c, 0x4e254e4a, 0xb7e6b7d1, 0xeb8beb0b, 0x3cf03cfd, 0x813e817c, 0x946a94d4, 0xf7fbf7eb, 0xb9deb9a1, 0x134c1398, 0x2cb02c7d, 0xd36bd3d6, 0xe7bbe76b, 0x6ea56e57, 0xc437c46e, 0x030c0318, 0x5645568a, 0x440d441a, 0x7fe17fdf, 0xa99ea921, 0x2aa82a4d, 0xbbd6bbb1, 0xc123c146, 0x535153a2, 0xdc57dcae, 0x0b2c0b58, 0x9d4e9d9c, 0x6cad6c47, 0x31c43195, 0x74cd7487, 0xf6fff6e3, 0x4605460a, 0xac8aac09, 0x891e893c, 0x145014a0, 0xe1a3e15b, 0x165816b0, 0x3ae83acd, 0x69b9696f, 0x09240948, 0x70dd70a7, 0xb6e2b6d9, 0xd067d0ce, 0xed93ed3b, 0xcc17cc2e, 0x4215422a, 0x985a98b4, 0xa4aaa449, 0x28a0285d, 0x5c6d5cda, 0xf8c7f893, 0x86228644, }, { 0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8, 0x87872687, 0xb8b8dab8, 0x01010401, 0x4f4f214f, 0x3636d836, 0xa6a6a2a6, 0xd2d26fd2, 0xf5f5f3f5, 0x7979f979, 0x6f6fa16f, 0x91917e91, 0x52525552, 0x60609d60, 0xbcbccabc, 0x9b9b569b, 0x8e8e028e, 0xa3a3b6a3, 0x0c0c300c, 0x7b7bf17b, 0x3535d435, 0x1d1d741d, 0xe0e0a7e0, 0xd7d77bd7, 0xc2c22fc2, 0x2e2eb82e, 0x4b4b314b, 0xfefedffe, 0x57574157, 0x15155415, 0x7777c177, 0x3737dc37, 0xe5e5b3e5, 0x9f9f469f, 0xf0f0e7f0, 0x4a4a354a, 0xdada4fda, 0x58587d58, 0xc9c903c9, 0x2929a429, 0x0a0a280a, 0xb1b1feb1, 0xa0a0baa0, 0x6b6bb16b, 0x85852e85, 0xbdbdcebd, 0x5d5d695d, 0x10104010, 0xf4f4f7f4, 0xcbcb0bcb, 0x3e3ef83e, 0x05051405, 0x67678167, 0xe4e4b7e4, 0x27279c27, 0x41411941, 0x8b8b168b, 0xa7a7a6a7, 0x7d7de97d, 0x95956e95, 0xd8d847d8, 0xfbfbcbfb, 0xeeee9fee, 0x7c7ced7c, 0x66668566, 0xdddd53dd, 0x17175c17, 0x47470147, 0x9e9e429e, 0xcaca0fca, 0x2d2db42d, 0xbfbfc6bf, 0x07071c07, 0xadad8ead, 0x5a5a755a, 0x83833683, 0x3333cc33, 0x63639163, 0x02020802, 0xaaaa92aa, 0x7171d971, 0xc8c807c8, 0x19196419, 0x49493949, 0xd9d943d9, 0xf2f2eff2, 0xe3e3abe3, 0x5b5b715b, 0x88881a88, 0x9a9a529a, 0x26269826, 0x3232c832, 0xb0b0fab0, 0xe9e983e9, 0x0f0f3c0f, 0xd5d573d5, 0x80803a80, 0xbebec2be, 0xcdcd13cd, 0x3434d034, 0x48483d48, 0xffffdbff, 0x7a7af57a, 0x90907a90, 0x5f5f615f, 0x20208020, 0x6868bd68, 0x1a1a681a, 0xaeae82ae, 0xb4b4eab4, 0x54544d54, 0x93937693, 0x22228822, 0x64648d64, 0xf1f1e3f1, 0x7373d173, 0x12124812, 0x40401d40, 0x08082008, 0xc3c32bc3, 0xecec97ec, 0xdbdb4bdb, 0xa1a1bea1, 0x8d8d0e8d, 0x3d3df43d, 0x97976697, 0x00000000, 0xcfcf1bcf, 0x2b2bac2b, 0x7676c576, 0x82823282, 0xd6d67fd6, 0x1b1b6c1b, 0xb5b5eeb5, 0xafaf86af, 0x6a6ab56a, 0x50505d50, 0x45450945, 0xf3f3ebf3, 0x3030c030, 0xefef9bef, 0x3f3ffc3f, 0x55554955, 0xa2a2b2a2, 0xeaea8fea, 0x65658965, 0xbabad2ba, 0x2f2fbc2f, 0xc0c027c0, 0xdede5fde, 0x1c1c701c, 0xfdfdd3fd, 0x4d4d294d, 0x92927292, 0x7575c975, 0x06061806, 0x8a8a128a, 0xb2b2f2b2, 0xe6e6bfe6, 0x0e0e380e, 0x1f1f7c1f, 0x62629562, 0xd4d477d4, 0xa8a89aa8, 0x96966296, 0xf9f9c3f9, 0xc5c533c5, 0x25259425, 0x59597959, 0x84842a84, 0x7272d572, 0x3939e439, 0x4c4c2d4c, 0x5e5e655e, 0x7878fd78, 0x3838e038, 0x8c8c0a8c, 0xd1d163d1, 0xa5a5aea5, 0xe2e2afe2, 0x61619961, 0xb3b3f6b3, 0x21218421, 0x9c9c4a9c, 0x1e1e781e, 0x43431143, 0xc7c73bc7, 0xfcfcd7fc, 0x04041004, 0x51515951, 0x99995e99, 0x6d6da96d, 0x0d0d340d, 0xfafacffa, 0xdfdf5bdf, 0x7e7ee57e, 0x24249024, 0x3b3bec3b, 0xabab96ab, 0xcece1fce, 0x11114411, 0x8f8f068f, 0x4e4e254e, 0xb7b7e6b7, 0xebeb8beb, 0x3c3cf03c, 0x81813e81, 0x94946a94, 0xf7f7fbf7, 0xb9b9deb9, 0x13134c13, 0x2c2cb02c, 0xd3d36bd3, 0xe7e7bbe7, 0x6e6ea56e, 0xc4c437c4, 0x03030c03, 0x56564556, 0x44440d44, 0x7f7fe17f, 0xa9a99ea9, 0x2a2aa82a, 0xbbbbd6bb, 0xc1c123c1, 0x53535153, 0xdcdc57dc, 0x0b0b2c0b, 0x9d9d4e9d, 0x6c6cad6c, 0x3131c431, 0x7474cd74, 0xf6f6fff6, 0x46460546, 0xacac8aac, 0x89891e89, 0x14145014, 0xe1e1a3e1, 0x16165816, 0x3a3ae83a, 0x6969b969, 0x09092409, 0x7070dd70, 0xb6b6e2b6, 0xd0d067d0, 0xeded93ed, 0xcccc17cc, 0x42421542, 0x98985a98, 0xa4a4aaa4, 0x2828a028, 0x5c5c6d5c, 0xf8f8c7f8, 0x86862286, }, { 0xd8181860, 0x2623238c, 0xb8c6c63f, 0xfbe8e887, 0xcb878726, 0x11b8b8da, 0x09010104, 0x0d4f4f21, 0x9b3636d8, 0xffa6a6a2, 0x0cd2d26f, 0x0ef5f5f3, 0x967979f9, 0x306f6fa1, 0x6d91917e, 0xf8525255, 0x4760609d, 0x35bcbcca, 0x379b9b56, 0x8a8e8e02, 0xd2a3a3b6, 0x6c0c0c30, 0x847b7bf1, 0x803535d4, 0xf51d1d74, 0xb3e0e0a7, 0x21d7d77b, 0x9cc2c22f, 0x432e2eb8, 0x294b4b31, 0x5dfefedf, 0xd5575741, 0xbd151554, 0xe87777c1, 0x923737dc, 0x9ee5e5b3, 0x139f9f46, 0x23f0f0e7, 0x204a4a35, 0x44dada4f, 0xa258587d, 0xcfc9c903, 0x7c2929a4, 0x5a0a0a28, 0x50b1b1fe, 0xc9a0a0ba, 0x146b6bb1, 0xd985852e, 0x3cbdbdce, 0x8f5d5d69, 0x90101040, 0x07f4f4f7, 0xddcbcb0b, 0xd33e3ef8, 0x2d050514, 0x78676781, 0x97e4e4b7, 0x0227279c, 0x73414119, 0xa78b8b16, 0xf6a7a7a6, 0xb27d7de9, 0x4995956e, 0x56d8d847, 0x70fbfbcb, 0xcdeeee9f, 0xbb7c7ced, 0x71666685, 0x7bdddd53, 0xaf17175c, 0x45474701, 0x1a9e9e42, 0xd4caca0f, 0x582d2db4, 0x2ebfbfc6, 0x3f07071c, 0xacadad8e, 0xb05a5a75, 0xef838336, 0xb63333cc, 0x5c636391, 0x12020208, 0x93aaaa92, 0xde7171d9, 0xc6c8c807, 0xd1191964, 0x3b494939, 0x5fd9d943, 0x31f2f2ef, 0xa8e3e3ab, 0xb95b5b71, 0xbc88881a, 0x3e9a9a52, 0x0b262698, 0xbf3232c8, 0x59b0b0fa, 0xf2e9e983, 0x770f0f3c, 0x33d5d573, 0xf480803a, 0x27bebec2, 0xebcdcd13, 0x893434d0, 0x3248483d, 0x54ffffdb, 0x8d7a7af5, 0x6490907a, 0x9d5f5f61, 0x3d202080, 0x0f6868bd, 0xca1a1a68, 0xb7aeae82, 0x7db4b4ea, 0xce54544d, 0x7f939376, 0x2f222288, 0x6364648d, 0x2af1f1e3, 0xcc7373d1, 0x82121248, 0x7a40401d, 0x48080820, 0x95c3c32b, 0xdfecec97, 0x4ddbdb4b, 0xc0a1a1be, 0x918d8d0e, 0xc83d3df4, 0x5b979766, 0x00000000, 0xf9cfcf1b, 0x6e2b2bac, 0xe17676c5, 0xe6828232, 0x28d6d67f, 0xc31b1b6c, 0x74b5b5ee, 0xbeafaf86, 0x1d6a6ab5, 0xea50505d, 0x57454509, 0x38f3f3eb, 0xad3030c0, 0xc4efef9b, 0xda3f3ffc, 0xc7555549, 0xdba2a2b2, 0xe9eaea8f, 0x6a656589, 0x03babad2, 0x4a2f2fbc, 0x8ec0c027, 0x60dede5f, 0xfc1c1c70, 0x46fdfdd3, 0x1f4d4d29, 0x76929272, 0xfa7575c9, 0x36060618, 0xae8a8a12, 0x4bb2b2f2, 0x85e6e6bf, 0x7e0e0e38, 0xe71f1f7c, 0x55626295, 0x3ad4d477, 0x81a8a89a, 0x52969662, 0x62f9f9c3, 0xa3c5c533, 0x10252594, 0xab595979, 0xd084842a, 0xc57272d5, 0xec3939e4, 0x164c4c2d, 0x945e5e65, 0x9f7878fd, 0xe53838e0, 0x988c8c0a, 0x17d1d163, 0xe4a5a5ae, 0xa1e2e2af, 0x4e616199, 0x42b3b3f6, 0x34212184, 0x089c9c4a, 0xee1e1e78, 0x61434311, 0xb1c7c73b, 0x4ffcfcd7, 0x24040410, 0xe3515159, 0x2599995e, 0x226d6da9, 0x650d0d34, 0x79fafacf, 0x69dfdf5b, 0xa97e7ee5, 0x19242490, 0xfe3b3bec, 0x9aabab96, 0xf0cece1f, 0x99111144, 0x838f8f06, 0x044e4e25, 0x66b7b7e6, 0xe0ebeb8b, 0xc13c3cf0, 0xfd81813e, 0x4094946a, 0x1cf7f7fb, 0x18b9b9de, 0x8b13134c, 0x512c2cb0, 0x05d3d36b, 0x8ce7e7bb, 0x396e6ea5, 0xaac4c437, 0x1b03030c, 0xdc565645, 0x5e44440d, 0xa07f7fe1, 0x88a9a99e, 0x672a2aa8, 0x0abbbbd6, 0x87c1c123, 0xf1535351, 0x72dcdc57, 0x530b0b2c, 0x019d9d4e, 0x2b6c6cad, 0xa43131c4, 0xf37474cd, 0x15f6f6ff, 0x4c464605, 0xa5acac8a, 0xb589891e, 0xb4141450, 0xbae1e1a3, 0xa6161658, 0xf73a3ae8, 0x066969b9, 0x41090924, 0xd77070dd, 0x6fb6b6e2, 0x1ed0d067, 0xd6eded93, 0xe2cccc17, 0x68424215, 0x2c98985a, 0xeda4a4aa, 0x752828a0, 0x865c5c6d, 0x6bf8f8c7, 0xc2868622, }, { 0x30d81818, 0x46262323, 0x91b8c6c6, 0xcdfbe8e8, 0x13cb8787, 0x6d11b8b8, 0x02090101, 0x9e0d4f4f, 0x6c9b3636, 0x51ffa6a6, 0xb90cd2d2, 0xf70ef5f5, 0xf2967979, 0xde306f6f, 0x3f6d9191, 0xa4f85252, 0xc0476060, 0x6535bcbc, 0x2b379b9b, 0x018a8e8e, 0x5bd2a3a3, 0x186c0c0c, 0xf6847b7b, 0x6a803535, 0x3af51d1d, 0xddb3e0e0, 0xb321d7d7, 0x999cc2c2, 0x5c432e2e, 0x96294b4b, 0xe15dfefe, 0xaed55757, 0x2abd1515, 0xeee87777, 0x6e923737, 0xd79ee5e5, 0x23139f9f, 0xfd23f0f0, 0x94204a4a, 0xa944dada, 0xb0a25858, 0x8fcfc9c9, 0x527c2929, 0x145a0a0a, 0x7f50b1b1, 0x5dc9a0a0, 0xd6146b6b, 0x17d98585, 0x673cbdbd, 0xba8f5d5d, 0x20901010, 0xf507f4f4, 0x8bddcbcb, 0x7cd33e3e, 0x0a2d0505, 0xce786767, 0xd597e4e4, 0x4e022727, 0x82734141, 0x0ba78b8b, 0x53f6a7a7, 0xfab27d7d, 0x37499595, 0xad56d8d8, 0xeb70fbfb, 0xc1cdeeee, 0xf8bb7c7c, 0xcc716666, 0xa77bdddd, 0x2eaf1717, 0x8e454747, 0x211a9e9e, 0x89d4caca, 0x5a582d2d, 0x632ebfbf, 0x0e3f0707, 0x47acadad, 0xb4b05a5a, 0x1bef8383, 0x66b63333, 0xc65c6363, 0x04120202, 0x4993aaaa, 0xe2de7171, 0x8dc6c8c8, 0x32d11919, 0x923b4949, 0xaf5fd9d9, 0xf931f2f2, 0xdba8e3e3, 0xb6b95b5b, 0x0dbc8888, 0x293e9a9a, 0x4c0b2626, 0x64bf3232, 0x7d59b0b0, 0xcff2e9e9, 0x1e770f0f, 0xb733d5d5, 0x1df48080, 0x6127bebe, 0x87ebcdcd, 0x68893434, 0x90324848, 0xe354ffff, 0xf48d7a7a, 0x3d649090, 0xbe9d5f5f, 0x403d2020, 0xd00f6868, 0x34ca1a1a, 0x41b7aeae, 0x757db4b4, 0xa8ce5454, 0x3b7f9393, 0x442f2222, 0xc8636464, 0xff2af1f1, 0xe6cc7373, 0x24821212, 0x807a4040, 0x10480808, 0x9b95c3c3, 0xc5dfecec, 0xab4ddbdb, 0x5fc0a1a1, 0x07918d8d, 0x7ac83d3d, 0x335b9797, 0x00000000, 0x83f9cfcf, 0x566e2b2b, 0xece17676, 0x19e68282, 0xb128d6d6, 0x36c31b1b, 0x7774b5b5, 0x43beafaf, 0xd41d6a6a, 0xa0ea5050, 0x8a574545, 0xfb38f3f3, 0x60ad3030, 0xc3c4efef, 0x7eda3f3f, 0xaac75555, 0x59dba2a2, 0xc9e9eaea, 0xca6a6565, 0x6903baba, 0x5e4a2f2f, 0x9d8ec0c0, 0xa160dede, 0x38fc1c1c, 0xe746fdfd, 0x9a1f4d4d, 0x39769292, 0xeafa7575, 0x0c360606, 0x09ae8a8a, 0x794bb2b2, 0xd185e6e6, 0x1c7e0e0e, 0x3ee71f1f, 0xc4556262, 0xb53ad4d4, 0x4d81a8a8, 0x31529696, 0xef62f9f9, 0x97a3c5c5, 0x4a102525, 0xb2ab5959, 0x15d08484, 0xe4c57272, 0x72ec3939, 0x98164c4c, 0xbc945e5e, 0xf09f7878, 0x70e53838, 0x05988c8c, 0xbf17d1d1, 0x57e4a5a5, 0xd9a1e2e2, 0xc24e6161, 0x7b42b3b3, 0x42342121, 0x25089c9c, 0x3cee1e1e, 0x86614343, 0x93b1c7c7, 0xe54ffcfc, 0x08240404, 0xa2e35151, 0x2f259999, 0xda226d6d, 0x1a650d0d, 0xe979fafa, 0xa369dfdf, 0xfca97e7e, 0x48192424, 0x76fe3b3b, 0x4b9aabab, 0x81f0cece, 0x22991111, 0x03838f8f, 0x9c044e4e, 0x7366b7b7, 0xcbe0ebeb, 0x78c13c3c, 0x1ffd8181, 0x35409494, 0xf31cf7f7, 0x6f18b9b9, 0x268b1313, 0x58512c2c, 0xbb05d3d3, 0xd38ce7e7, 0xdc396e6e, 0x95aac4c4, 0x061b0303, 0xacdc5656, 0x885e4444, 0xfea07f7f, 0x4f88a9a9, 0x54672a2a, 0x6b0abbbb, 0x9f87c1c1, 0xa6f15353, 0xa572dcdc, 0x16530b0b, 0x27019d9d, 0xd82b6c6c, 0x62a43131, 0xe8f37474, 0xf115f6f6, 0x8c4c4646, 0x45a5acac, 0x0fb58989, 0x28b41414, 0xdfbae1e1, 0x2ca61616, 0x74f73a3a, 0xd2066969, 0x12410909, 0xe0d77070, 0x716fb6b6, 0xbd1ed0d0, 0xc7d6eded, 0x85e2cccc, 0x84684242, 0x2d2c9898, 0x55eda4a4, 0x50752828, 0xb8865c5c, 0xed6bf8f8, 0x11c28686, }, { 0x7830d818, 0xaf462623, 0xf991b8c6, 0x6fcdfbe8, 0xa113cb87, 0x626d11b8, 0x05020901, 0x6e9e0d4f, 0xee6c9b36, 0x0451ffa6, 0xbdb90cd2, 0x06f70ef5, 0x80f29679, 0xcede306f, 0xef3f6d91, 0x07a4f852, 0xfdc04760, 0x766535bc, 0xcd2b379b, 0x8c018a8e, 0x155bd2a3, 0x3c186c0c, 0x8af6847b, 0xe16a8035, 0x693af51d, 0x47ddb3e0, 0xacb321d7, 0xed999cc2, 0x965c432e, 0x7a96294b, 0x21e15dfe, 0x16aed557, 0x412abd15, 0xb6eee877, 0xeb6e9237, 0x56d79ee5, 0xd923139f, 0x17fd23f0, 0x7f94204a, 0x95a944da, 0x25b0a258, 0xca8fcfc9, 0x8d527c29, 0x22145a0a, 0x4f7f50b1, 0x1a5dc9a0, 0xdad6146b, 0xab17d985, 0x73673cbd, 0x34ba8f5d, 0x50209010, 0x03f507f4, 0xc08bddcb, 0xc67cd33e, 0x110a2d05, 0xe6ce7867, 0x53d597e4, 0xbb4e0227, 0x58827341, 0x9d0ba78b, 0x0153f6a7, 0x94fab27d, 0xfb374995, 0x9fad56d8, 0x30eb70fb, 0x71c1cdee, 0x91f8bb7c, 0xe3cc7166, 0x8ea77bdd, 0x4b2eaf17, 0x468e4547, 0xdc211a9e, 0xc589d4ca, 0x995a582d, 0x79632ebf, 0x1b0e3f07, 0x2347acad, 0x2fb4b05a, 0xb51bef83, 0xff66b633, 0xf2c65c63, 0x0a041202, 0x384993aa, 0xa8e2de71, 0xcf8dc6c8, 0x7d32d119, 0x70923b49, 0x9aaf5fd9, 0x1df931f2, 0x48dba8e3, 0x2ab6b95b, 0x920dbc88, 0xc8293e9a, 0xbe4c0b26, 0xfa64bf32, 0x4a7d59b0, 0x6acff2e9, 0x331e770f, 0xa6b733d5, 0xba1df480, 0x7c6127be, 0xde87ebcd, 0xe4688934, 0x75903248, 0x24e354ff, 0x8ff48d7a, 0xea3d6490, 0x3ebe9d5f, 0xa0403d20, 0xd5d00f68, 0x7234ca1a, 0x2c41b7ae, 0x5e757db4, 0x19a8ce54, 0xe53b7f93, 0xaa442f22, 0xe9c86364, 0x12ff2af1, 0xa2e6cc73, 0x5a248212, 0x5d807a40, 0x28104808, 0xe89b95c3, 0x7bc5dfec, 0x90ab4ddb, 0x1f5fc0a1, 0x8307918d, 0xc97ac83d, 0xf1335b97, 0x00000000, 0xd483f9cf, 0x87566e2b, 0xb3ece176, 0xb019e682, 0xa9b128d6, 0x7736c31b, 0x5b7774b5, 0x2943beaf, 0xdfd41d6a, 0x0da0ea50, 0x4c8a5745, 0x18fb38f3, 0xf060ad30, 0x74c3c4ef, 0xc37eda3f, 0x1caac755, 0x1059dba2, 0x65c9e9ea, 0xecca6a65, 0x686903ba, 0x935e4a2f, 0xe79d8ec0, 0x81a160de, 0x6c38fc1c, 0x2ee746fd, 0x649a1f4d, 0xe0397692, 0xbceafa75, 0x1e0c3606, 0x9809ae8a, 0x40794bb2, 0x59d185e6, 0x361c7e0e, 0x633ee71f, 0xf7c45562, 0xa3b53ad4, 0x324d81a8, 0xf4315296, 0x3aef62f9, 0xf697a3c5, 0xb14a1025, 0x20b2ab59, 0xae15d084, 0xa7e4c572, 0xdd72ec39, 0x6198164c, 0x3bbc945e, 0x85f09f78, 0xd870e538, 0x8605988c, 0xb2bf17d1, 0x0b57e4a5, 0x4dd9a1e2, 0xf8c24e61, 0x457b42b3, 0xa5423421, 0xd625089c, 0x663cee1e, 0x52866143, 0xfc93b1c7, 0x2be54ffc, 0x14082404, 0x08a2e351, 0xc72f2599, 0xc4da226d, 0x391a650d, 0x35e979fa, 0x84a369df, 0x9bfca97e, 0xb4481924, 0xd776fe3b, 0x3d4b9aab, 0xd181f0ce, 0x55229911, 0x8903838f, 0x6b9c044e, 0x517366b7, 0x60cbe0eb, 0xcc78c13c, 0xbf1ffd81, 0xfe354094, 0x0cf31cf7, 0x676f18b9, 0x5f268b13, 0x9c58512c, 0xb8bb05d3, 0x5cd38ce7, 0xcbdc396e, 0xf395aac4, 0x0f061b03, 0x13acdc56, 0x49885e44, 0x9efea07f, 0x374f88a9, 0x8254672a, 0x6d6b0abb, 0xe29f87c1, 0x02a6f153, 0x8ba572dc, 0x2716530b, 0xd327019d, 0xc1d82b6c, 0xf562a431, 0xb9e8f374, 0x09f115f6, 0x438c4c46, 0x2645a5ac, 0x970fb589, 0x4428b414, 0x42dfbae1, 0x4e2ca616, 0xd274f73a, 0xd0d20669, 0x2d124109, 0xade0d770, 0x54716fb6, 0xb7bd1ed0, 0x7ec7d6ed, 0xdb85e2cc, 0x57846842, 0xc22d2c98, 0x0e55eda4, 0x88507528, 0x31b8865c, 0x3fed6bf8, 0xa411c286, }, }; __constant u32 rch[R + 1] = { 0x00000000, 0x1823c6e8, 0x36a6d2f5, 0x60bc9b8e, 0x1de0d7c2, 0x157737e5, 0x58c9290a, 0xbd5d10f4, 0xe427418b, 0xfbee7c66, 0xca2dbf07, }; __constant u32 rcl[R + 1] = { 0x00000000, 0x87b8014f, 0x796f9152, 0xa30c7b35, 0x2e4bfe57, 0x9ff04ada, 0xb1a06b85, 0xcb3e0567, 0xa77d95d8, 0xdd17479e, 0xad5a8333, }; // this is a highly optimized that assumes dgst[16] = { 0 }; only reuse of no 2nd transform is needed static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256]) { u32 Kh[8]; u32 Kl[8]; Kh[0] = 0x300beec0; Kl[0] = 0xaf902967; Kh[1] = 0x28282828; Kl[1] = 0x28282828; Kh[2] = 0x28282828; Kl[2] = 0x28282828; Kh[3] = 0x28282828; Kl[3] = 0x28282828; Kh[4] = 0x28282828; Kl[4] = 0x28282828; Kh[5] = 0x28282828; Kl[5] = 0x28282828; Kh[6] = 0x28282828; Kl[6] = 0x28282828; Kh[7] = 0x28282828; Kl[7] = 0x28282828; u32 stateh[8]; u32 statel[8]; stateh[0] = w[ 0]; statel[0] = w[ 1]; stateh[1] = w[ 2]; statel[1] = w[ 3]; stateh[2] = w[ 4]; statel[2] = w[ 5]; stateh[3] = w[ 6]; statel[3] = w[ 7]; stateh[4] = w[ 8]; statel[4] = w[ 9]; stateh[5] = w[10]; statel[5] = w[11]; stateh[6] = w[12]; statel[6] = w[13]; stateh[7] = w[14]; statel[7] = w[15]; u32 Lh[8]; u32 Ll[8]; #pragma unroll for (int i = 0; i < 8; i++) { const u32 Lp0 = stateh[(i + 8) & 7] >> 24; const u32 Lp1 = stateh[(i + 7) & 7] >> 16; const u32 Lp2 = stateh[(i + 6) & 7] >> 8; const u32 Lp3 = stateh[(i + 5) & 7] >> 0; const u32 Lp4 = statel[(i + 4) & 7] >> 24; const u32 Lp5 = statel[(i + 3) & 7] >> 16; const u32 Lp6 = statel[(i + 2) & 7] >> 8; const u32 Lp7 = statel[(i + 1) & 7] >> 0; Lh[i] = BOX (s_Ch, 0, Lp0 & 0xff) ^ BOX (s_Ch, 1, Lp1 & 0xff) ^ BOX (s_Ch, 2, Lp2 & 0xff) ^ BOX (s_Ch, 3, Lp3 & 0xff) ^ BOX (s_Ch, 4, Lp4 & 0xff) ^ BOX (s_Ch, 5, Lp5 & 0xff) ^ BOX (s_Ch, 6, Lp6 & 0xff) ^ BOX (s_Ch, 7, Lp7 & 0xff); Ll[i] = BOX (s_Cl, 0, Lp0 & 0xff) ^ BOX (s_Cl, 1, Lp1 & 0xff) ^ BOX (s_Cl, 2, Lp2 & 0xff) ^ BOX (s_Cl, 3, Lp3 & 0xff) ^ BOX (s_Cl, 4, Lp4 & 0xff) ^ BOX (s_Cl, 5, Lp5 & 0xff) ^ BOX (s_Cl, 6, Lp6 & 0xff) ^ BOX (s_Cl, 7, Lp7 & 0xff); } stateh[0] = Lh[0] ^ Kh[0]; statel[0] = Ll[0] ^ Kl[0]; stateh[1] = Lh[1] ^ Kh[1]; statel[1] = Ll[1] ^ Kl[1]; stateh[2] = Lh[2] ^ Kh[2]; statel[2] = Ll[2] ^ Kl[2]; stateh[3] = Lh[3] ^ Kh[3]; statel[3] = Ll[3] ^ Kl[3]; stateh[4] = Lh[4] ^ Kh[4]; statel[4] = Ll[4] ^ Kl[4]; stateh[5] = Lh[5] ^ Kh[5]; statel[5] = Ll[5] ^ Kl[5]; stateh[6] = Lh[6] ^ Kh[6]; statel[6] = Ll[6] ^ Kl[6]; stateh[7] = Lh[7] ^ Kh[7]; statel[7] = Ll[7] ^ Kl[7]; for (int r = 2; r <= R; r++) { u32 Lh[8]; u32 Ll[8]; #pragma unroll for (int i = 0; i < 8; i++) { const u32 Lp0 = Kh[(i + 8) & 7] >> 24; const u32 Lp1 = Kh[(i + 7) & 7] >> 16; const u32 Lp2 = Kh[(i + 6) & 7] >> 8; const u32 Lp3 = Kh[(i + 5) & 7] >> 0; const u32 Lp4 = Kl[(i + 4) & 7] >> 24; const u32 Lp5 = Kl[(i + 3) & 7] >> 16; const u32 Lp6 = Kl[(i + 2) & 7] >> 8; const u32 Lp7 = Kl[(i + 1) & 7] >> 0; Lh[i] = BOX (s_Ch, 0, Lp0 & 0xff) ^ BOX (s_Ch, 1, Lp1 & 0xff) ^ BOX (s_Ch, 2, Lp2 & 0xff) ^ BOX (s_Ch, 3, Lp3 & 0xff) ^ BOX (s_Ch, 4, Lp4 & 0xff) ^ BOX (s_Ch, 5, Lp5 & 0xff) ^ BOX (s_Ch, 6, Lp6 & 0xff) ^ BOX (s_Ch, 7, Lp7 & 0xff); Ll[i] = BOX (s_Cl, 0, Lp0 & 0xff) ^ BOX (s_Cl, 1, Lp1 & 0xff) ^ BOX (s_Cl, 2, Lp2 & 0xff) ^ BOX (s_Cl, 3, Lp3 & 0xff) ^ BOX (s_Cl, 4, Lp4 & 0xff) ^ BOX (s_Cl, 5, Lp5 & 0xff) ^ BOX (s_Cl, 6, Lp6 & 0xff) ^ BOX (s_Cl, 7, Lp7 & 0xff); } Kh[0] = Lh[0] ^ rch[r]; Kl[0] = Ll[0] ^ rcl[r]; Kh[1] = Lh[1]; Kl[1] = Ll[1]; Kh[2] = Lh[2]; Kl[2] = Ll[2]; Kh[3] = Lh[3]; Kl[3] = Ll[3]; Kh[4] = Lh[4]; Kl[4] = Ll[4]; Kh[5] = Lh[5]; Kl[5] = Ll[5]; Kh[6] = Lh[6]; Kl[6] = Ll[6]; Kh[7] = Lh[7]; Kl[7] = Ll[7]; #pragma unroll 8 for (int i = 0; i < 8; i++) { const u32 Lp0 = stateh[(i + 8) & 7] >> 24; const u32 Lp1 = stateh[(i + 7) & 7] >> 16; const u32 Lp2 = stateh[(i + 6) & 7] >> 8; const u32 Lp3 = stateh[(i + 5) & 7] >> 0; const u32 Lp4 = statel[(i + 4) & 7] >> 24; const u32 Lp5 = statel[(i + 3) & 7] >> 16; const u32 Lp6 = statel[(i + 2) & 7] >> 8; const u32 Lp7 = statel[(i + 1) & 7] >> 0; Lh[i] = BOX (s_Ch, 0, Lp0 & 0xff) ^ BOX (s_Ch, 1, Lp1 & 0xff) ^ BOX (s_Ch, 2, Lp2 & 0xff) ^ BOX (s_Ch, 3, Lp3 & 0xff) ^ BOX (s_Ch, 4, Lp4 & 0xff) ^ BOX (s_Ch, 5, Lp5 & 0xff) ^ BOX (s_Ch, 6, Lp6 & 0xff) ^ BOX (s_Ch, 7, Lp7 & 0xff); Ll[i] = BOX (s_Cl, 0, Lp0 & 0xff) ^ BOX (s_Cl, 1, Lp1 & 0xff) ^ BOX (s_Cl, 2, Lp2 & 0xff) ^ BOX (s_Cl, 3, Lp3 & 0xff) ^ BOX (s_Cl, 4, Lp4 & 0xff) ^ BOX (s_Cl, 5, Lp5 & 0xff) ^ BOX (s_Cl, 6, Lp6 & 0xff) ^ BOX (s_Cl, 7, Lp7 & 0xff); } stateh[0] = Lh[0] ^ Kh[0]; statel[0] = Ll[0] ^ Kl[0]; stateh[1] = Lh[1] ^ Kh[1]; statel[1] = Ll[1] ^ Kl[1]; stateh[2] = Lh[2] ^ Kh[2]; statel[2] = Ll[2] ^ Kl[2]; stateh[3] = Lh[3] ^ Kh[3]; statel[3] = Ll[3] ^ Kl[3]; stateh[4] = Lh[4] ^ Kh[4]; statel[4] = Ll[4] ^ Kl[4]; stateh[5] = Lh[5] ^ Kh[5]; statel[5] = Ll[5] ^ Kl[5]; stateh[6] = Lh[6] ^ Kh[6]; statel[6] = Ll[6] ^ Kl[6]; stateh[7] = Lh[7] ^ Kh[7]; statel[7] = Ll[7] ^ Kl[7]; } dgst[ 0] = stateh[0] ^ w[ 0]; dgst[ 1] = statel[0] ^ w[ 1]; dgst[ 2] = stateh[1] ^ w[ 2]; dgst[ 3] = statel[1] ^ w[ 3]; dgst[ 4] = stateh[2] ^ w[ 4]; dgst[ 5] = statel[2] ^ w[ 5]; dgst[ 6] = stateh[3] ^ w[ 6]; dgst[ 7] = statel[3] ^ w[ 7]; dgst[ 8] = stateh[4] ^ w[ 8]; dgst[ 9] = statel[4] ^ w[ 9]; dgst[10] = stateh[5] ^ w[10]; dgst[11] = statel[5] ^ w[11]; dgst[12] = stateh[6] ^ w[12]; dgst[13] = statel[6] ^ w[13]; dgst[14] = stateh[7] ^ w[14]; dgst[15] = statel[7] ^ w[15]; } static void m06100m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 s_Cl[8][256], __local u32 s_Ch[8][256]) { /** * modifier */ const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); /** * loop */ u32 w0l = w0[0]; for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++) { const u32 w0r = bfs_buf[il_pos].i; w0[0] = w0l | w0r; u32 wl[16]; wl[ 0] = w0[0]; wl[ 1] = w0[1]; wl[ 2] = w0[2]; wl[ 3] = w0[3]; wl[ 4] = w1[0]; wl[ 5] = w1[1]; wl[ 6] = w1[2]; wl[ 7] = w1[3]; wl[ 8] = w2[0]; wl[ 9] = w2[1]; wl[10] = w2[2]; wl[11] = w2[3]; wl[12] = w3[0]; wl[13] = w3[1]; wl[14] = 0; wl[15] = pw_len * 8; u32 dgst[16]; whirlpool_transform (wl, dgst, s_Ch, s_Cl); const u32 r0 = dgst[0]; const u32 r1 = dgst[1]; const u32 r2 = dgst[2]; const u32 r3 = dgst[3]; #include COMPARE_M } } static void m06100s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 s_Cl[8][256], __local u32 s_Ch[8][256]) { /** * modifier */ const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); /** * digest */ const u32 search[4] = { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R2], digests_buf[digests_offset].digest_buf[DGST_R3] }; /** * loop */ u32 w0l = w0[0]; for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++) { const u32 w0r = bfs_buf[il_pos].i; w0[0] = w0l | w0r; u32 wl[16]; wl[ 0] = w0[0]; wl[ 1] = w0[1]; wl[ 2] = w0[2]; wl[ 3] = w0[3]; wl[ 4] = w1[0]; wl[ 5] = w1[1]; wl[ 6] = w1[2]; wl[ 7] = w1[3]; wl[ 8] = w2[0]; wl[ 9] = w2[1]; wl[10] = w2[2]; wl[11] = w2[3]; wl[12] = w3[0]; wl[13] = w3[1]; wl[14] = 0; wl[15] = pw_len * 8; u32 dgst[16]; whirlpool_transform (wl, dgst, s_Ch, s_Cl); const u32 r0 = dgst[0]; const u32 r1 = dgst[1]; const u32 r2 = dgst[2]; const u32 r3 = dgst[3]; #include COMPARE_S } } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier */ const u32 lid = get_local_id (0); /** * base */ const u32 gid = get_global_id (0); u32 w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32 w1[4]; w1[0] = 0; w1[1] = 0; w1[2] = 0; w1[3] = 0; u32 w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32 w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * modifier */ __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; const u32 lid4 = lid * 4; for (u32 i = 0; i < 8; i++) { s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; } barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; /** * main */ m06100m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_Cl, s_Ch); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier */ const u32 lid = get_local_id (0); /** * base */ const u32 gid = get_global_id (0); u32 w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32 w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32 w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32 w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * modifier */ __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; const u32 lid4 = lid * 4; for (u32 i = 0; i < 8; i++) { s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; } barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; /** * main */ m06100m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_Cl, s_Ch); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier */ const u32 lid = get_local_id (0); /** * base */ const u32 gid = get_global_id (0); u32 w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32 w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32 w2[4]; w2[0] = pws[gid].i[ 8]; w2[1] = pws[gid].i[ 9]; w2[2] = pws[gid].i[10]; w2[3] = pws[gid].i[11]; u32 w3[4]; w3[0] = pws[gid].i[12]; w3[1] = pws[gid].i[13]; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * modifier */ __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; const u32 lid4 = lid * 4; for (u32 i = 0; i < 8; i++) { s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; } barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; /** * main */ m06100m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_Cl, s_Ch); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier */ const u32 lid = get_local_id (0); /** * base */ const u32 gid = get_global_id (0); u32 w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32 w1[4]; w1[0] = 0; w1[1] = 0; w1[2] = 0; w1[3] = 0; u32 w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32 w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * modifier */ __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; const u32 lid4 = lid * 4; for (u32 i = 0; i < 8; i++) { s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; } barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; /** * main */ m06100s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_Cl, s_Ch); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier */ const u32 lid = get_local_id (0); /** * base */ const u32 gid = get_global_id (0); u32 w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32 w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32 w2[4]; w2[0] = 0; w2[1] = 0; w2[2] = 0; w2[3] = 0; u32 w3[4]; w3[0] = 0; w3[1] = 0; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * modifier */ __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; const u32 lid4 = lid * 4; for (u32 i = 0; i < 8; i++) { s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; } barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; /** * main */ m06100s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_Cl, s_Ch); } __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06100_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier */ const u32 lid = get_local_id (0); /** * base */ const u32 gid = get_global_id (0); u32 w0[4]; w0[0] = pws[gid].i[ 0]; w0[1] = pws[gid].i[ 1]; w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; u32 w1[4]; w1[0] = pws[gid].i[ 4]; w1[1] = pws[gid].i[ 5]; w1[2] = pws[gid].i[ 6]; w1[3] = pws[gid].i[ 7]; u32 w2[4]; w2[0] = pws[gid].i[ 8]; w2[1] = pws[gid].i[ 9]; w2[2] = pws[gid].i[10]; w2[3] = pws[gid].i[11]; u32 w3[4]; w3[0] = pws[gid].i[12]; w3[1] = pws[gid].i[13]; w3[2] = 0; w3[3] = 0; const u32 pw_len = pws[gid].pw_len; /** * modifier */ __local u32 s_Ch[8][256]; __local u32 s_Cl[8][256]; const u32 lid4 = lid * 4; for (u32 i = 0; i < 8; i++) { s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; } barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; /** * main */ m06100s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset, s_Cl, s_Ch); }