arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-04 22:08:40 +00:00
Code Issues Releases Wiki Activity
a7ab9f6ca9
hashcat/OpenCL/m06231.cl

2242 lines
84 KiB
Common Lisp
Raw Normal View History

Initial commit
2015-12-04 14:47:52 +00:00
/**
Added docs/credits.txt Added docs/team.txt
2016-09-11 20:20:15 +00:00
* Author......: See docs/credits.txt
Initial commit
2015-12-04 14:47:52 +00:00
* License.....: MIT
*/
#define _WHIRLPOOL_
Move files from include/ to OpenCL/ if they are used within kernels Rename includes in OpenCL so that it's easier to recognize them as such
2016-05-25 21:04:26 +00:00
#include "inc_vendor.cl"
Move macros DGST_R0 - DGST_R3 to host, define dgst_size for opencl kernel from host; both at runtime
2016-06-26 21:39:42 +00:00
#include "inc_hash_constants.h"
Move files from include/ to OpenCL/ if they are used within kernels Rename includes in OpenCL so that it's easier to recognize them as such
2016-05-25 21:04:26 +00:00
#include "inc_hash_functions.cl"
#include "inc_types.cl"
#include "inc_common.cl"
Initial commit
2015-12-04 14:47:52 +00:00
Move files from include/ to OpenCL/ if they are used within kernels Rename includes in OpenCL so that it's easier to recognize them as such
2016-05-25 21:04:26 +00:00
#include "inc_cipher_aes256.cl"
#include "inc_cipher_twofish256.cl"
#include "inc_cipher_serpent256.cl"
Initial commit
2015-12-04 14:47:52 +00:00
#define R 10
__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,
},
};
More kernel fixes for function calls and vector datatypes
2015-12-15 16:46:00 +00:00
#define BOX(S,n,i) (S)[(n)][(i)]
Initial commit
2015-12-04 14:47:52 +00:00
- Added inline declaration to functions from simd.c, common.c, rp.c and types_ocl.c to increase performance - Dropped static declaration from functions in all kernel to achieve OpenCL 1.1 compatibility - Added -cl-std=CL1.1 to all kernel build options - Created environment variable to inform NVidia OpenCL runtime to not create its own kernel cache - Created environment variable to inform pocl OpenCL runtime to not create its own kernel cache
2016-05-01 21:15:26 +00:00
void whirlpool_transform_last (u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
{
const u32 rch[R + 1] =
{
0x00000000,
0x1823c6e8,
0x36a6d2f5,
0x60bc9b8e,
0x1de0d7c2,
0x157737e5,
0x58c9290a,
0xbd5d10f4,
0xe427418b,
0xfbee7c66,
0xca2dbf07,
};
const u32 rcl[R + 1] =
{
0x00000000,
0x87b8014f,
0x796f9152,
0xa30c7b35,
0x2e4bfe57,
0x9ff04ada,
0xb1a06b85,
0xcb3e0567,
0xa77d95d8,
0xdd17479e,
0xad5a8333,
};
u32 Kh[8];
u32 Kl[8];
Kh[0] = dgst[ 0];
Kl[0] = dgst[ 1];
Kh[1] = dgst[ 2];
Kl[1] = dgst[ 3];
Kh[2] = dgst[ 4];
Kl[2] = dgst[ 5];
Kh[3] = dgst[ 6];
Kl[3] = dgst[ 7];
Kh[4] = dgst[ 8];
Kl[4] = dgst[ 9];
Kh[5] = dgst[10];
Kl[5] = dgst[11];
Kh[6] = dgst[12];
Kl[6] = dgst[13];
Kh[7] = dgst[14];
Kl[7] = dgst[15];
u32 stateh[8];
u32 statel[8];
#define LAST_W00 0x80000000
#define LAST_W15 ((64 + 64) * 8)
stateh[0] = Kh[0] ^ LAST_W00;
statel[0] = Kl[0];
stateh[1] = Kh[1];
statel[1] = Kl[1];
stateh[2] = Kh[2];
statel[2] = Kl[2];
stateh[3] = Kh[3];
statel[3] = Kl[3];
stateh[4] = Kh[4];
statel[4] = Kl[4];
stateh[5] = Kh[5];
statel[5] = Kl[5];
stateh[6] = Kh[6];
statel[6] = Kl[6];
stateh[7] = Kh[7];
statel[7] = Kl[7] ^ LAST_W15;
u32 r;
for (r = 1; r <= R; r++)
{
u32 Lh[8];
u32 Ll[8];
u32 i;
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#ifdef _unroll
Fix Truecrypt Whirlpool speed
2016-02-16 09:47:40 +00:00
#pragma unroll
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#endif
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
for (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];
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#ifdef _unroll
Fix Truecrypt Whirlpool speed
2016-02-16 09:47:40 +00:00
#pragma unroll
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#endif
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
for (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] ^ LAST_W00;
dgst[ 1] ^= statel[0];
dgst[ 2] ^= stateh[1];
dgst[ 3] ^= statel[1];
dgst[ 4] ^= stateh[2];
dgst[ 5] ^= statel[2];
dgst[ 6] ^= stateh[3];
dgst[ 7] ^= statel[3];
dgst[ 8] ^= stateh[4];
dgst[ 9] ^= statel[4];
dgst[10] ^= stateh[5];
dgst[11] ^= statel[5];
dgst[12] ^= stateh[6];
dgst[13] ^= statel[6];
dgst[14] ^= stateh[7];
dgst[15] ^= statel[7] ^ LAST_W15;
}
- Added inline declaration to functions from simd.c, common.c, rp.c and types_ocl.c to increase performance - Dropped static declaration from functions in all kernel to achieve OpenCL 1.1 compatibility - Added -cl-std=CL1.1 to all kernel build options - Created environment variable to inform NVidia OpenCL runtime to not create its own kernel cache - Created environment variable to inform pocl OpenCL runtime to not create its own kernel cache
2016-05-01 21:15:26 +00:00
void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
Initial commit
2015-12-04 14:47:52 +00:00
{
const u32 rch[R + 1] =
{
0x00000000,
0x1823c6e8,
0x36a6d2f5,
0x60bc9b8e,
0x1de0d7c2,
0x157737e5,
0x58c9290a,
0xbd5d10f4,
0xe427418b,
0xfbee7c66,
0xca2dbf07,
};
const u32 rcl[R + 1] =
{
0x00000000,
0x87b8014f,
0x796f9152,
0xa30c7b35,
0x2e4bfe57,
0x9ff04ada,
0xb1a06b85,
0xcb3e0567,
0xa77d95d8,
0xdd17479e,
0xad5a8333,
};
u32 Kh[8];
u32 Kl[8];
Kh[0] = dgst[ 0];
Kl[0] = dgst[ 1];
Kh[1] = dgst[ 2];
Kl[1] = dgst[ 3];
Kh[2] = dgst[ 4];
Kl[2] = dgst[ 5];
Kh[3] = dgst[ 6];
Kl[3] = dgst[ 7];
Kh[4] = dgst[ 8];
Kl[4] = dgst[ 9];
Kh[5] = dgst[10];
Kl[5] = dgst[11];
Kh[6] = dgst[12];
Kl[6] = dgst[13];
Kh[7] = dgst[14];
Kl[7] = dgst[15];
u32 stateh[8];
u32 statel[8];
stateh[0] = w[ 0] ^ Kh[0];
statel[0] = w[ 1] ^ Kl[0];
stateh[1] = w[ 2] ^ Kh[1];
statel[1] = w[ 3] ^ Kl[1];
stateh[2] = w[ 4] ^ Kh[2];
statel[2] = w[ 5] ^ Kl[2];
stateh[3] = w[ 6] ^ Kh[3];
statel[3] = w[ 7] ^ Kl[3];
stateh[4] = w[ 8] ^ Kh[4];
statel[4] = w[ 9] ^ Kl[4];
stateh[5] = w[10] ^ Kh[5];
statel[5] = w[11] ^ Kl[5];
stateh[6] = w[12] ^ Kh[6];
statel[6] = w[13] ^ Kl[6];
stateh[7] = w[14] ^ Kh[7];
statel[7] = w[15] ^ Kl[7];
u32 r;
for (r = 1; r <= R; r++)
{
u32 Lh[8];
u32 Ll[8];
u32 i;
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#ifdef _unroll
Fix Truecrypt Whirlpool speed
2016-02-16 09:47:40 +00:00
#pragma unroll
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#endif
Initial commit
2015-12-04 14:47:52 +00:00
for (i = 0; i < 8; i++)
{
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
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;
Initial commit
2015-12-04 14:47:52 +00:00
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];
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#ifdef _unroll
Fix Truecrypt Whirlpool speed
2016-02-16 09:47:40 +00:00
#pragma unroll
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#endif
Initial commit
2015-12-04 14:47:52 +00:00
for (i = 0; i < 8; i++)
{
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
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;
Initial commit
2015-12-04 14:47:52 +00:00
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];
}
- Added inline declaration to functions from simd.c, common.c, rp.c and types_ocl.c to increase performance - Dropped static declaration from functions in all kernel to achieve OpenCL 1.1 compatibility - Added -cl-std=CL1.1 to all kernel build options - Created environment variable to inform NVidia OpenCL runtime to not create its own kernel cache - Created environment variable to inform pocl OpenCL runtime to not create its own kernel cache
2016-05-01 21:15:26 +00:00
void hmac_run2a (const u32 w1[16], const u32 w2[16], const u32 ipad[16], const u32 opad[16], u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
Initial commit
2015-12-04 14:47:52 +00:00
{
dgst[ 0] = ipad[ 0];
dgst[ 1] = ipad[ 1];
dgst[ 2] = ipad[ 2];
dgst[ 3] = ipad[ 3];
dgst[ 4] = ipad[ 4];
dgst[ 5] = ipad[ 5];
dgst[ 6] = ipad[ 6];
dgst[ 7] = ipad[ 7];
dgst[ 8] = ipad[ 8];
dgst[ 9] = ipad[ 9];
dgst[10] = ipad[10];
dgst[11] = ipad[11];
dgst[12] = ipad[12];
dgst[13] = ipad[13];
dgst[14] = ipad[14];
dgst[15] = ipad[15];
whirlpool_transform (w1, dgst, s_Ch, s_Cl);
whirlpool_transform (w2, dgst, s_Ch, s_Cl);
u32 w[16];
w[ 0] = dgst[ 0];
w[ 1] = dgst[ 1];
w[ 2] = dgst[ 2];
w[ 3] = dgst[ 3];
w[ 4] = dgst[ 4];
w[ 5] = dgst[ 5];
w[ 6] = dgst[ 6];
w[ 7] = dgst[ 7];
w[ 8] = dgst[ 8];
w[ 9] = dgst[ 9];
w[10] = dgst[10];
w[11] = dgst[11];
w[12] = dgst[12];
w[13] = dgst[13];
w[14] = dgst[14];
w[15] = dgst[15];
dgst[ 0] = opad[ 0];
dgst[ 1] = opad[ 1];
dgst[ 2] = opad[ 2];
dgst[ 3] = opad[ 3];
dgst[ 4] = opad[ 4];
dgst[ 5] = opad[ 5];
dgst[ 6] = opad[ 6];
dgst[ 7] = opad[ 7];
dgst[ 8] = opad[ 8];
dgst[ 9] = opad[ 9];
dgst[10] = opad[10];
dgst[11] = opad[11];
dgst[12] = opad[12];
dgst[13] = opad[13];
dgst[14] = opad[14];
dgst[15] = opad[15];
whirlpool_transform (w, dgst, s_Ch, s_Cl);
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
whirlpool_transform_last (dgst, s_Ch, s_Cl);
}
- Added inline declaration to functions from simd.c, common.c, rp.c and types_ocl.c to increase performance - Dropped static declaration from functions in all kernel to achieve OpenCL 1.1 compatibility - Added -cl-std=CL1.1 to all kernel build options - Created environment variable to inform NVidia OpenCL runtime to not create its own kernel cache - Created environment variable to inform pocl OpenCL runtime to not create its own kernel cache
2016-05-01 21:15:26 +00:00
void hmac_run2b (const u32 w1[16], const u32 ipad[16], const u32 opad[16], u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
{
dgst[ 0] = ipad[ 0];
dgst[ 1] = ipad[ 1];
dgst[ 2] = ipad[ 2];
dgst[ 3] = ipad[ 3];
dgst[ 4] = ipad[ 4];
dgst[ 5] = ipad[ 5];
dgst[ 6] = ipad[ 6];
dgst[ 7] = ipad[ 7];
dgst[ 8] = ipad[ 8];
dgst[ 9] = ipad[ 9];
dgst[10] = ipad[10];
dgst[11] = ipad[11];
dgst[12] = ipad[12];
dgst[13] = ipad[13];
dgst[14] = ipad[14];
dgst[15] = ipad[15];
whirlpool_transform (w1, dgst, s_Ch, s_Cl);
whirlpool_transform_last (dgst, s_Ch, s_Cl);
u32 w[16];
w[ 0] = dgst[ 0];
w[ 1] = dgst[ 1];
w[ 2] = dgst[ 2];
w[ 3] = dgst[ 3];
w[ 4] = dgst[ 4];
w[ 5] = dgst[ 5];
w[ 6] = dgst[ 6];
w[ 7] = dgst[ 7];
w[ 8] = dgst[ 8];
w[ 9] = dgst[ 9];
w[10] = dgst[10];
w[11] = dgst[11];
w[12] = dgst[12];
w[13] = dgst[13];
w[14] = dgst[14];
w[15] = dgst[15];
dgst[ 0] = opad[ 0];
dgst[ 1] = opad[ 1];
dgst[ 2] = opad[ 2];
dgst[ 3] = opad[ 3];
dgst[ 4] = opad[ 4];
dgst[ 5] = opad[ 5];
dgst[ 6] = opad[ 6];
dgst[ 7] = opad[ 7];
dgst[ 8] = opad[ 8];
dgst[ 9] = opad[ 9];
dgst[10] = opad[10];
dgst[11] = opad[11];
dgst[12] = opad[12];
dgst[13] = opad[13];
dgst[14] = opad[14];
dgst[15] = opad[15];
Initial commit
2015-12-04 14:47:52 +00:00
whirlpool_transform (w, dgst, s_Ch, s_Cl);
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
whirlpool_transform_last (dgst, s_Ch, s_Cl);
Initial commit
2015-12-04 14:47:52 +00:00
}
- Added inline declaration to functions from simd.c, common.c, rp.c and types_ocl.c to increase performance - Dropped static declaration from functions in all kernel to achieve OpenCL 1.1 compatibility - Added -cl-std=CL1.1 to all kernel build options - Created environment variable to inform NVidia OpenCL runtime to not create its own kernel cache - Created environment variable to inform pocl OpenCL runtime to not create its own kernel cache
2016-05-01 21:15:26 +00:00
void hmac_init (u32 w[16], u32 ipad[16], u32 opad[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
Initial commit
2015-12-04 14:47:52 +00:00
{
w[ 0] ^= 0x36363636;
w[ 1] ^= 0x36363636;
w[ 2] ^= 0x36363636;
w[ 3] ^= 0x36363636;
w[ 4] ^= 0x36363636;
w[ 5] ^= 0x36363636;
w[ 6] ^= 0x36363636;
w[ 7] ^= 0x36363636;
w[ 8] ^= 0x36363636;
w[ 9] ^= 0x36363636;
w[10] ^= 0x36363636;
w[11] ^= 0x36363636;
w[12] ^= 0x36363636;
w[13] ^= 0x36363636;
w[14] ^= 0x36363636;
w[15] ^= 0x36363636;
ipad[ 0] = 0;
ipad[ 1] = 0;
ipad[ 2] = 0;
ipad[ 3] = 0;
ipad[ 4] = 0;
ipad[ 5] = 0;
ipad[ 6] = 0;
ipad[ 7] = 0;
ipad[ 8] = 0;
ipad[ 9] = 0;
ipad[10] = 0;
ipad[11] = 0;
ipad[12] = 0;
ipad[13] = 0;
ipad[14] = 0;
ipad[15] = 0;
whirlpool_transform (w, ipad, s_Ch, s_Cl);
w[ 0] ^= 0x6a6a6a6a;
w[ 1] ^= 0x6a6a6a6a;
w[ 2] ^= 0x6a6a6a6a;
w[ 3] ^= 0x6a6a6a6a;
w[ 4] ^= 0x6a6a6a6a;
w[ 5] ^= 0x6a6a6a6a;
w[ 6] ^= 0x6a6a6a6a;
w[ 7] ^= 0x6a6a6a6a;
w[ 8] ^= 0x6a6a6a6a;
w[ 9] ^= 0x6a6a6a6a;
w[10] ^= 0x6a6a6a6a;
w[11] ^= 0x6a6a6a6a;
w[12] ^= 0x6a6a6a6a;
w[13] ^= 0x6a6a6a6a;
w[14] ^= 0x6a6a6a6a;
w[15] ^= 0x6a6a6a6a;
opad[ 0] = 0;
opad[ 1] = 0;
opad[ 2] = 0;
opad[ 3] = 0;
opad[ 4] = 0;
opad[ 5] = 0;
opad[ 6] = 0;
opad[ 7] = 0;
opad[ 8] = 0;
opad[ 9] = 0;
opad[10] = 0;
opad[11] = 0;
opad[12] = 0;
opad[13] = 0;
opad[14] = 0;
opad[15] = 0;
whirlpool_transform (w, opad, s_Ch, s_Cl);
}
- Added inline declaration to functions from simd.c, common.c, rp.c and types_ocl.c to increase performance - Dropped static declaration from functions in all kernel to achieve OpenCL 1.1 compatibility - Added -cl-std=CL1.1 to all kernel build options - Created environment variable to inform NVidia OpenCL runtime to not create its own kernel cache - Created environment variable to inform pocl OpenCL runtime to not create its own kernel cache
2016-05-01 21:15:26 +00:00
u32 u8add (const u32 a, const u32 b)
Initial commit
2015-12-04 14:47:52 +00:00
{
const u32 a1 = (a >> 0) & 0xff;
const u32 a2 = (a >> 8) & 0xff;
const u32 a3 = (a >> 16) & 0xff;
const u32 a4 = (a >> 24) & 0xff;
const u32 b1 = (b >> 0) & 0xff;
const u32 b2 = (b >> 8) & 0xff;
const u32 b3 = (b >> 16) & 0xff;
const u32 b4 = (b >> 24) & 0xff;
const u32 r1 = (a1 + b1) & 0xff;
const u32 r2 = (a2 + b2) & 0xff;
const u32 r3 = (a3 + b3) & 0xff;
const u32 r4 = (a4 + b4) & 0xff;
const u32 r = r1 << 0
| r2 << 8
| r3 << 16
| r4 << 24;
return r;
}
Workaround memory allocation limit from OpenCL by using multiple buffers for scrypt
2016-06-28 09:03:04 +00:00
__kernel void m06231_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_t *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)
Initial commit
2015-12-04 14:47:52 +00:00
{
/**
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
* modifier
Initial commit
2015-12-04 14:47:52 +00:00
*/
const u32 gid = get_global_id (0);
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
const u32 lid = get_local_id (0);
const u32 lsz = get_local_size (0);
/**
* shared
*/
__local u32 s_Ch[8][256];
__local u32 s_Cl[8][256];
for (u32 i = lid; i < 256; i += lsz)
{
s_Ch[0][i] = Ch[0][i];
s_Ch[1][i] = Ch[1][i];
s_Ch[2][i] = Ch[2][i];
s_Ch[3][i] = Ch[3][i];
s_Ch[4][i] = Ch[4][i];
s_Ch[5][i] = Ch[5][i];
s_Ch[6][i] = Ch[6][i];
s_Ch[7][i] = Ch[7][i];
s_Cl[0][i] = Cl[0][i];
s_Cl[1][i] = Cl[1][i];
s_Cl[2][i] = Cl[2][i];
s_Cl[3][i] = Cl[3][i];
s_Cl[4][i] = Cl[4][i];
s_Cl[5][i] = Cl[5][i];
s_Cl[6][i] = Cl[6][i];
s_Cl[7][i] = Cl[7][i];
}
barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
/**
* base
*/
Initial commit
2015-12-04 14:47:52 +00:00
- Dropped all vector code since new GPU's are all scalar, makes the code much easier - Some performance on low-end GPU may drop because of that, but only for a few hash-modes - Dropped scalar code (aka warp) since we do not have any vector datatypes anymore - Renamed C++ overloading functions memcat32_9 -> memcat_c32_w4x4_a3x4 - Still need to fix kernels to new function names, needs to be done manually - Temperature Management needs to be rewritten partially because of conflicting datatypes names - Added code to create different codepaths for NV on AMD in runtime in host (see data.vendor_id) - Added code to create different codepaths for NV on AMD in runtime in kernels (see IS_NV and IS_AMD) - First tests working for -m 0, for example - Great performance increases in general for NV so far - Tested amp_* and markov_* kernel - Migrated special NV optimizations for rule processor
2015-12-15 11:04:22 +00:00
u32 w0[4];
Initial commit
2015-12-04 14:47:52 +00:00
w0[0] = pws[gid].i[ 0];
w0[1] = pws[gid].i[ 1];
w0[2] = pws[gid].i[ 2];
w0[3] = pws[gid].i[ 3];
- Dropped all vector code since new GPU's are all scalar, makes the code much easier - Some performance on low-end GPU may drop because of that, but only for a few hash-modes - Dropped scalar code (aka warp) since we do not have any vector datatypes anymore - Renamed C++ overloading functions memcat32_9 -> memcat_c32_w4x4_a3x4 - Still need to fix kernels to new function names, needs to be done manually - Temperature Management needs to be rewritten partially because of conflicting datatypes names - Added code to create different codepaths for NV on AMD in runtime in host (see data.vendor_id) - Added code to create different codepaths for NV on AMD in runtime in kernels (see IS_NV and IS_AMD) - First tests working for -m 0, for example - Great performance increases in general for NV so far - Tested amp_* and markov_* kernel - Migrated special NV optimizations for rule processor
2015-12-15 11:04:22 +00:00
u32 w1[4];
Initial commit
2015-12-04 14:47:52 +00:00
w1[0] = pws[gid].i[ 4];
w1[1] = pws[gid].i[ 5];
w1[2] = pws[gid].i[ 6];
w1[3] = pws[gid].i[ 7];
- Dropped all vector code since new GPU's are all scalar, makes the code much easier - Some performance on low-end GPU may drop because of that, but only for a few hash-modes - Dropped scalar code (aka warp) since we do not have any vector datatypes anymore - Renamed C++ overloading functions memcat32_9 -> memcat_c32_w4x4_a3x4 - Still need to fix kernels to new function names, needs to be done manually - Temperature Management needs to be rewritten partially because of conflicting datatypes names - Added code to create different codepaths for NV on AMD in runtime in host (see data.vendor_id) - Added code to create different codepaths for NV on AMD in runtime in kernels (see IS_NV and IS_AMD) - First tests working for -m 0, for example - Great performance increases in general for NV so far - Tested amp_* and markov_* kernel - Migrated special NV optimizations for rule processor
2015-12-15 11:04:22 +00:00
u32 w2[4];
Initial commit
2015-12-04 14:47:52 +00:00
w2[0] = pws[gid].i[ 8];
w2[1] = pws[gid].i[ 9];
w2[2] = pws[gid].i[10];
w2[3] = pws[gid].i[11];
- Dropped all vector code since new GPU's are all scalar, makes the code much easier - Some performance on low-end GPU may drop because of that, but only for a few hash-modes - Dropped scalar code (aka warp) since we do not have any vector datatypes anymore - Renamed C++ overloading functions memcat32_9 -> memcat_c32_w4x4_a3x4 - Still need to fix kernels to new function names, needs to be done manually - Temperature Management needs to be rewritten partially because of conflicting datatypes names - Added code to create different codepaths for NV on AMD in runtime in host (see data.vendor_id) - Added code to create different codepaths for NV on AMD in runtime in kernels (see IS_NV and IS_AMD) - First tests working for -m 0, for example - Great performance increases in general for NV so far - Tested amp_* and markov_* kernel - Migrated special NV optimizations for rule processor
2015-12-15 11:04:22 +00:00
u32 w3[4];
Initial commit
2015-12-04 14:47:52 +00:00
w3[0] = pws[gid].i[12];
w3[1] = pws[gid].i[13];
w3[2] = pws[gid].i[14];
w3[3] = pws[gid].i[15];
/**
* keyfile
*/
w0[0] = u8add (w0[0], esalt_bufs[salt_pos].keyfile_buf[ 0]);
w0[1] = u8add (w0[1], esalt_bufs[salt_pos].keyfile_buf[ 1]);
w0[2] = u8add (w0[2], esalt_bufs[salt_pos].keyfile_buf[ 2]);
w0[3] = u8add (w0[3], esalt_bufs[salt_pos].keyfile_buf[ 3]);
w1[0] = u8add (w1[0], esalt_bufs[salt_pos].keyfile_buf[ 4]);
w1[1] = u8add (w1[1], esalt_bufs[salt_pos].keyfile_buf[ 5]);
w1[2] = u8add (w1[2], esalt_bufs[salt_pos].keyfile_buf[ 6]);
w1[3] = u8add (w1[3], esalt_bufs[salt_pos].keyfile_buf[ 7]);
w2[0] = u8add (w2[0], esalt_bufs[salt_pos].keyfile_buf[ 8]);
w2[1] = u8add (w2[1], esalt_bufs[salt_pos].keyfile_buf[ 9]);
w2[2] = u8add (w2[2], esalt_bufs[salt_pos].keyfile_buf[10]);
w2[3] = u8add (w2[3], esalt_bufs[salt_pos].keyfile_buf[11]);
w3[0] = u8add (w3[0], esalt_bufs[salt_pos].keyfile_buf[12]);
w3[1] = u8add (w3[1], esalt_bufs[salt_pos].keyfile_buf[13]);
w3[2] = u8add (w3[2], esalt_bufs[salt_pos].keyfile_buf[14]);
w3[3] = u8add (w3[3], esalt_bufs[salt_pos].keyfile_buf[15]);
/**
* salt
*/
u32 salt_buf1[16];
Fix more append_* functions in kernels
2015-12-15 15:50:21 +00:00
salt_buf1[ 0] = swap32 (esalt_bufs[salt_pos].salt_buf[ 0]);
salt_buf1[ 1] = swap32 (esalt_bufs[salt_pos].salt_buf[ 1]);
salt_buf1[ 2] = swap32 (esalt_bufs[salt_pos].salt_buf[ 2]);
salt_buf1[ 3] = swap32 (esalt_bufs[salt_pos].salt_buf[ 3]);
salt_buf1[ 4] = swap32 (esalt_bufs[salt_pos].salt_buf[ 4]);
salt_buf1[ 5] = swap32 (esalt_bufs[salt_pos].salt_buf[ 5]);
salt_buf1[ 6] = swap32 (esalt_bufs[salt_pos].salt_buf[ 6]);
salt_buf1[ 7] = swap32 (esalt_bufs[salt_pos].salt_buf[ 7]);
salt_buf1[ 8] = swap32 (esalt_bufs[salt_pos].salt_buf[ 8]);
salt_buf1[ 9] = swap32 (esalt_bufs[salt_pos].salt_buf[ 9]);
salt_buf1[10] = swap32 (esalt_bufs[salt_pos].salt_buf[10]);
salt_buf1[11] = swap32 (esalt_bufs[salt_pos].salt_buf[11]);
salt_buf1[12] = swap32 (esalt_bufs[salt_pos].salt_buf[12]);
salt_buf1[13] = swap32 (esalt_bufs[salt_pos].salt_buf[13]);
salt_buf1[14] = swap32 (esalt_bufs[salt_pos].salt_buf[14]);
salt_buf1[15] = swap32 (esalt_bufs[salt_pos].salt_buf[15]);
Initial commit
2015-12-04 14:47:52 +00:00
u32 salt_buf2[16];
salt_buf2[ 0] = 0;
salt_buf2[ 1] = 0x80000000;
salt_buf2[ 2] = 0;
salt_buf2[ 3] = 0;
salt_buf2[ 4] = 0;
salt_buf2[ 5] = 0;
salt_buf2[ 6] = 0;
salt_buf2[ 7] = 0;
salt_buf2[ 8] = 0;
salt_buf2[ 9] = 0;
salt_buf2[10] = 0;
salt_buf2[11] = 0;
salt_buf2[12] = 0;
salt_buf2[13] = 0;
salt_buf2[14] = 0;
salt_buf2[15] = (64 + 64 + 4) * 8;
const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
u32 w[16];
Fix more append_* functions in kernels
2015-12-15 15:50:21 +00:00
w[ 0] = swap32 (w0[0]);
w[ 1] = swap32 (w0[1]);
w[ 2] = swap32 (w0[2]);
w[ 3] = swap32 (w0[3]);
w[ 4] = swap32 (w1[0]);
w[ 5] = swap32 (w1[1]);
w[ 6] = swap32 (w1[2]);
w[ 7] = swap32 (w1[3]);
w[ 8] = swap32 (w2[0]);
w[ 9] = swap32 (w2[1]);
w[10] = swap32 (w2[2]);
w[11] = swap32 (w2[3]);
w[12] = swap32 (w3[0]);
w[13] = swap32 (w3[1]);
w[14] = swap32 (w3[2]);
w[15] = swap32 (w3[3]);
Initial commit
2015-12-04 14:47:52 +00:00
u32 ipad[16];
u32 opad[16];
hmac_init (w, ipad, opad, s_Ch, s_Cl);
tmps[gid].ipad[ 0] = ipad[ 0];
tmps[gid].ipad[ 1] = ipad[ 1];
tmps[gid].ipad[ 2] = ipad[ 2];
tmps[gid].ipad[ 3] = ipad[ 3];
tmps[gid].ipad[ 4] = ipad[ 4];
tmps[gid].ipad[ 5] = ipad[ 5];
tmps[gid].ipad[ 6] = ipad[ 6];
tmps[gid].ipad[ 7] = ipad[ 7];
tmps[gid].ipad[ 8] = ipad[ 8];
tmps[gid].ipad[ 9] = ipad[ 9];
tmps[gid].ipad[10] = ipad[10];
tmps[gid].ipad[11] = ipad[11];
tmps[gid].ipad[12] = ipad[12];
tmps[gid].ipad[13] = ipad[13];
tmps[gid].ipad[14] = ipad[14];
tmps[gid].ipad[15] = ipad[15];
tmps[gid].opad[ 0] = opad[ 0];
tmps[gid].opad[ 1] = opad[ 1];
tmps[gid].opad[ 2] = opad[ 2];
tmps[gid].opad[ 3] = opad[ 3];
tmps[gid].opad[ 4] = opad[ 4];
tmps[gid].opad[ 5] = opad[ 5];
tmps[gid].opad[ 6] = opad[ 6];
tmps[gid].opad[ 7] = opad[ 7];
tmps[gid].opad[ 8] = opad[ 8];
tmps[gid].opad[ 9] = opad[ 9];
tmps[gid].opad[10] = opad[10];
tmps[gid].opad[11] = opad[11];
tmps[gid].opad[12] = opad[12];
tmps[gid].opad[13] = opad[13];
tmps[gid].opad[14] = opad[14];
tmps[gid].opad[15] = opad[15];
for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 16, j += 1)
{
salt_buf2[0] = j;
u32 dgst[16];
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
hmac_run2a (salt_buf1, salt_buf2, ipad, opad, dgst, s_Ch, s_Cl);
Initial commit
2015-12-04 14:47:52 +00:00
tmps[gid].dgst[i + 0] = dgst[ 0];
tmps[gid].dgst[i + 1] = dgst[ 1];
tmps[gid].dgst[i + 2] = dgst[ 2];
tmps[gid].dgst[i + 3] = dgst[ 3];
tmps[gid].dgst[i + 4] = dgst[ 4];
tmps[gid].dgst[i + 5] = dgst[ 5];
tmps[gid].dgst[i + 6] = dgst[ 6];
tmps[gid].dgst[i + 7] = dgst[ 7];
tmps[gid].dgst[i + 8] = dgst[ 8];
tmps[gid].dgst[i + 9] = dgst[ 9];
tmps[gid].dgst[i + 10] = dgst[10];
tmps[gid].dgst[i + 11] = dgst[11];
tmps[gid].dgst[i + 12] = dgst[12];
tmps[gid].dgst[i + 13] = dgst[13];
tmps[gid].dgst[i + 14] = dgst[14];
tmps[gid].dgst[i + 15] = dgst[15];
tmps[gid].out[i + 0] = dgst[ 0];
tmps[gid].out[i + 1] = dgst[ 1];
tmps[gid].out[i + 2] = dgst[ 2];
tmps[gid].out[i + 3] = dgst[ 3];
tmps[gid].out[i + 4] = dgst[ 4];
tmps[gid].out[i + 5] = dgst[ 5];
tmps[gid].out[i + 6] = dgst[ 6];
tmps[gid].out[i + 7] = dgst[ 7];
tmps[gid].out[i + 8] = dgst[ 8];
tmps[gid].out[i + 9] = dgst[ 9];
tmps[gid].out[i + 10] = dgst[10];
tmps[gid].out[i + 11] = dgst[11];
tmps[gid].out[i + 12] = dgst[12];
tmps[gid].out[i + 13] = dgst[13];
tmps[gid].out[i + 14] = dgst[14];
tmps[gid].out[i + 15] = dgst[15];
}
}
Workaround memory allocation limit from OpenCL by using multiple buffers for scrypt
2016-06-28 09:03:04 +00:00
__kernel void m06231_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_t *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)
Initial commit
2015-12-04 14:47:52 +00:00
{
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
/**
* modifier
*/
Initial commit
2015-12-04 14:47:52 +00:00
const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0);
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
const u32 lsz = get_local_size (0);
/**
* shared
*/
Initial commit
2015-12-04 14:47:52 +00:00
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
__local u32 s_Ch[8][256];
__local u32 s_Cl[8][256];
Initial commit
2015-12-04 14:47:52 +00:00
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
for (u32 i = lid; i < 256; i += lsz)
Initial commit
2015-12-04 14:47:52 +00:00
{
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
s_Ch[0][i] = Ch[0][i];
s_Ch[1][i] = Ch[1][i];
s_Ch[2][i] = Ch[2][i];
s_Ch[3][i] = Ch[3][i];
s_Ch[4][i] = Ch[4][i];
s_Ch[5][i] = Ch[5][i];
s_Ch[6][i] = Ch[6][i];
s_Ch[7][i] = Ch[7][i];
s_Cl[0][i] = Cl[0][i];
s_Cl[1][i] = Cl[1][i];
s_Cl[2][i] = Cl[2][i];
s_Cl[3][i] = Cl[3][i];
s_Cl[4][i] = Cl[4][i];
s_Cl[5][i] = Cl[5][i];
s_Cl[6][i] = Cl[6][i];
s_Cl[7][i] = Cl[7][i];
Initial commit
2015-12-04 14:47:52 +00:00
}
barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
Initial commit
2015-12-04 14:47:52 +00:00
u32 ipad[16];
ipad[ 0] = tmps[gid].ipad[ 0];
ipad[ 1] = tmps[gid].ipad[ 1],
ipad[ 2] = tmps[gid].ipad[ 2];
ipad[ 3] = tmps[gid].ipad[ 3];
ipad[ 4] = tmps[gid].ipad[ 4];
ipad[ 5] = tmps[gid].ipad[ 5];
ipad[ 6] = tmps[gid].ipad[ 6],
ipad[ 7] = tmps[gid].ipad[ 7];
ipad[ 8] = tmps[gid].ipad[ 8];
ipad[ 9] = tmps[gid].ipad[ 9];
ipad[10] = tmps[gid].ipad[10];
ipad[11] = tmps[gid].ipad[11],
ipad[12] = tmps[gid].ipad[12];
ipad[13] = tmps[gid].ipad[13];
ipad[14] = tmps[gid].ipad[14];
ipad[15] = tmps[gid].ipad[15];
u32 opad[16];
opad[ 0] = tmps[gid].opad[ 0];
opad[ 1] = tmps[gid].opad[ 1],
opad[ 2] = tmps[gid].opad[ 2];
opad[ 3] = tmps[gid].opad[ 3];
opad[ 4] = tmps[gid].opad[ 4];
opad[ 5] = tmps[gid].opad[ 5];
opad[ 6] = tmps[gid].opad[ 6],
opad[ 7] = tmps[gid].opad[ 7];
opad[ 8] = tmps[gid].opad[ 8];
opad[ 9] = tmps[gid].opad[ 9];
opad[10] = tmps[gid].opad[10];
opad[11] = tmps[gid].opad[11],
opad[12] = tmps[gid].opad[12];
opad[13] = tmps[gid].opad[13];
opad[14] = tmps[gid].opad[14];
opad[15] = tmps[gid].opad[15];
for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 16)
{
u32 dgst[16];
dgst[ 0] = tmps[gid].dgst[i + 0];
dgst[ 1] = tmps[gid].dgst[i + 1];
dgst[ 2] = tmps[gid].dgst[i + 2];
dgst[ 3] = tmps[gid].dgst[i + 3];
dgst[ 4] = tmps[gid].dgst[i + 4];
dgst[ 5] = tmps[gid].dgst[i + 5];
dgst[ 6] = tmps[gid].dgst[i + 6];
dgst[ 7] = tmps[gid].dgst[i + 7];
dgst[ 8] = tmps[gid].dgst[i + 8];
dgst[ 9] = tmps[gid].dgst[i + 9];
dgst[10] = tmps[gid].dgst[i + 10];
dgst[11] = tmps[gid].dgst[i + 11];
dgst[12] = tmps[gid].dgst[i + 12];
dgst[13] = tmps[gid].dgst[i + 13];
dgst[14] = tmps[gid].dgst[i + 14];
dgst[15] = tmps[gid].dgst[i + 15];
u32 out[16];
out[ 0] = tmps[gid].out[i + 0];
out[ 1] = tmps[gid].out[i + 1];
out[ 2] = tmps[gid].out[i + 2];
out[ 3] = tmps[gid].out[i + 3];
out[ 4] = tmps[gid].out[i + 4];
out[ 5] = tmps[gid].out[i + 5];
out[ 6] = tmps[gid].out[i + 6];
out[ 7] = tmps[gid].out[i + 7];
out[ 8] = tmps[gid].out[i + 8];
out[ 9] = tmps[gid].out[i + 9];
out[10] = tmps[gid].out[i + 10];
out[11] = tmps[gid].out[i + 11];
out[12] = tmps[gid].out[i + 12];
out[13] = tmps[gid].out[i + 13];
out[14] = tmps[gid].out[i + 14];
out[15] = tmps[gid].out[i + 15];
for (u32 j = 0; j < loop_cnt; j++)
{
u32 w1[16];
w1[ 0] = dgst[ 0];
w1[ 1] = dgst[ 1];
w1[ 2] = dgst[ 2];
w1[ 3] = dgst[ 3];
w1[ 4] = dgst[ 4];
w1[ 5] = dgst[ 5];
w1[ 6] = dgst[ 6];
w1[ 7] = dgst[ 7];
w1[ 8] = dgst[ 8];
w1[ 9] = dgst[ 9];
w1[10] = dgst[10];
w1[11] = dgst[11];
w1[12] = dgst[12];
w1[13] = dgst[13];
w1[14] = dgst[14];
w1[15] = dgst[15];
Fix TC whirlpool speed for NV
2015-12-28 21:07:09 +00:00
hmac_run2b (w1, ipad, opad, dgst, s_Ch, s_Cl);
Initial commit
2015-12-04 14:47:52 +00:00
out[ 0] ^= dgst[ 0];
out[ 1] ^= dgst[ 1];
out[ 2] ^= dgst[ 2];
out[ 3] ^= dgst[ 3];
out[ 4] ^= dgst[ 4];
out[ 5] ^= dgst[ 5];
out[ 6] ^= dgst[ 6];
out[ 7] ^= dgst[ 7];
out[ 8] ^= dgst[ 8];
out[ 9] ^= dgst[ 9];
out[10] ^= dgst[10];
out[11] ^= dgst[11];
out[12] ^= dgst[12];
out[13] ^= dgst[13];
out[14] ^= dgst[14];
out[15] ^= dgst[15];
}
tmps[gid].dgst[i + 0] = dgst[ 0];
tmps[gid].dgst[i + 1] = dgst[ 1];
tmps[gid].dgst[i + 2] = dgst[ 2];
tmps[gid].dgst[i + 3] = dgst[ 3];
tmps[gid].dgst[i + 4] = dgst[ 4];
tmps[gid].dgst[i + 5] = dgst[ 5];
tmps[gid].dgst[i + 6] = dgst[ 6];
tmps[gid].dgst[i + 7] = dgst[ 7];
tmps[gid].dgst[i + 8] = dgst[ 8];
tmps[gid].dgst[i + 9] = dgst[ 9];
tmps[gid].dgst[i + 10] = dgst[10];
tmps[gid].dgst[i + 11] = dgst[11];
tmps[gid].dgst[i + 12] = dgst[12];
tmps[gid].dgst[i + 13] = dgst[13];
tmps[gid].dgst[i + 14] = dgst[14];
tmps[gid].dgst[i + 15] = dgst[15];
tmps[gid].out[i + 0] = out[ 0];
tmps[gid].out[i + 1] = out[ 1];
tmps[gid].out[i + 2] = out[ 2];
tmps[gid].out[i + 3] = out[ 3];
tmps[gid].out[i + 4] = out[ 4];
tmps[gid].out[i + 5] = out[ 5];
tmps[gid].out[i + 6] = out[ 6];
tmps[gid].out[i + 7] = out[ 7];
tmps[gid].out[i + 8] = out[ 8];
tmps[gid].out[i + 9] = out[ 9];
tmps[gid].out[i + 10] = out[10];
tmps[gid].out[i + 11] = out[11];
tmps[gid].out[i + 12] = out[12];
tmps[gid].out[i + 13] = out[13];
tmps[gid].out[i + 14] = out[14];
tmps[gid].out[i + 15] = out[15];
}
}
Workaround memory allocation limit from OpenCL by using multiple buffers for scrypt
2016-06-28 09:03:04 +00:00
__kernel void m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_t *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)
Initial commit
2015-12-04 14:47:52 +00:00
{
/**
* base
*/
const u32 gid = get_global_id (0);
const u32 lid = get_local_id (0);
if (gid >= gid_max) return;
u32 ukey1[8];
Fix more append_* functions in kernels
2015-12-15 15:50:21 +00:00
ukey1[0] = swap32 (tmps[gid].out[ 0]);
ukey1[1] = swap32 (tmps[gid].out[ 1]);
ukey1[2] = swap32 (tmps[gid].out[ 2]);
ukey1[3] = swap32 (tmps[gid].out[ 3]);
ukey1[4] = swap32 (tmps[gid].out[ 4]);
ukey1[5] = swap32 (tmps[gid].out[ 5]);
ukey1[6] = swap32 (tmps[gid].out[ 6]);
ukey1[7] = swap32 (tmps[gid].out[ 7]);
Initial commit
2015-12-04 14:47:52 +00:00
u32 ukey2[8];
Fix more append_* functions in kernels
2015-12-15 15:50:21 +00:00
ukey2[0] = swap32 (tmps[gid].out[ 8]);
ukey2[1] = swap32 (tmps[gid].out[ 9]);
ukey2[2] = swap32 (tmps[gid].out[10]);
ukey2[3] = swap32 (tmps[gid].out[11]);
ukey2[4] = swap32 (tmps[gid].out[12]);
ukey2[5] = swap32 (tmps[gid].out[13]);
ukey2[6] = swap32 (tmps[gid].out[14]);
ukey2[7] = swap32 (tmps[gid].out[15]);
Initial commit
2015-12-04 14:47:52 +00:00
u32 data[4];
data[0] = esalt_bufs[0].data_buf[0];
data[1] = esalt_bufs[0].data_buf[1];
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
Prepare for VeraCrypt integration
2016-05-19 18:53:17 +00:00
const u32 signature = esalt_bufs[0].signature;
Initial commit
2015-12-04 14:47:52 +00:00
u32 tmp[4];
{
tmp[0] = data[0];
tmp[1] = data[1];
tmp[2] = data[2];
tmp[3] = data[3];
aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
Prepare for VeraCrypt integration
2016-05-19 18:53:17 +00:00
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
Initial commit
2015-12-04 14:47:52 +00:00
{
Some fixes for rare kernels
2016-05-20 17:05:54 +00:00
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
Initial commit
2015-12-04 14:47:52 +00:00
}
}
{
tmp[0] = data[0];
tmp[1] = data[1];
tmp[2] = data[2];
tmp[3] = data[3];
serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
Prepare for VeraCrypt integration
2016-05-19 18:53:17 +00:00
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
Initial commit
2015-12-04 14:47:52 +00:00
{
Some fixes for rare kernels
2016-05-20 17:05:54 +00:00
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
Initial commit
2015-12-04 14:47:52 +00:00
}
}
{
tmp[0] = data[0];
tmp[1] = data[1];
tmp[2] = data[2];
tmp[3] = data[3];
twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
Prepare for VeraCrypt integration
2016-05-19 18:53:17 +00:00
if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5)))
Initial commit
2015-12-04 14:47:52 +00:00
{
Some fixes for rare kernels
2016-05-20 17:05:54 +00:00
mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0);
Initial commit
2015-12-04 14:47:52 +00:00
}
}
}
Reference in New Issue Copy Permalink