mirror of
https://github.com/GOSTSec/gostoclminer
synced 2025-01-29 16:04:28 +00:00
initial commit of gostcoin miner
This commit is contained in:
orignal
parent
c1bbf987d6
commit
2a115d3692
1299
CL/cl_platform.h
Normal file
1299
CL/cl_platform.h
Normal file
File diff suppressed because it is too large
Load Diff
16
Makefile
16
Makefile
@ -1,21 +1,23 @@
|
||||
INCLUDEPATHS=-I${ATISTREAMSDKROOT}/include/ -L${ATISTREAMSDKROOT}/
|
||||
INCLUDEPATHS=-I.
|
||||
|
||||
LIBS=-lOpenCL -ljansson -lcurl
|
||||
LIBS=-lOpenCL -ljansson -lcurl -lpthread
|
||||
# LDFLAGS = -L$(AMDAPPSDKROOT)/lib/x86_64
|
||||
LDFLAGS = -L/usr/local/cuda/lib64
|
||||
|
||||
DEFS=
|
||||
DEBUGFLAGS=-g
|
||||
CFLAGS=-O3 -Wformat $(DEBUGFLAGS) $(DEFS) $(INCLUDEPATHS)
|
||||
HEADERS=
|
||||
|
||||
OBJS=miner.o ocl.o findnonce.o util.o
|
||||
OBJS=miner.o ocl.o util.o streebog.o
|
||||
|
||||
all: oclminer
|
||||
all: gostoclminer
|
||||
|
||||
%.o: %.c $(HEADERS)
|
||||
gcc -c $(CFLAGS) -o $@ $<
|
||||
|
||||
oclminer: $(OBJS)
|
||||
gcc $(CFLAGS) -o $@ $^ $(LIBS)
|
||||
gostoclminer: $(OBJS)
|
||||
gcc $(LDFLAGS) -o $@ $^ $(LIBS)
|
||||
|
||||
clean:
|
||||
-rm *.o oclminer
|
||||
-rm *.o gostoclminer
|
||||
|
||||
939
gostminer.cl
Normal file
939
gostminer.cl
Normal file
@ -0,0 +1,939 @@
|
||||
#define WGS __attribute__((reqd_work_group_size(256, 1, 1)))
|
||||
|
||||
typedef unsigned int sph_u32;
|
||||
typedef int sph_s32;
|
||||
#ifndef __OPENCL_VERSION__
|
||||
typedef unsigned long long sph_u64;
|
||||
typedef long long sph_s64;
|
||||
#else
|
||||
typedef unsigned long sph_u64;
|
||||
typedef long sph_s64;
|
||||
#endif
|
||||
|
||||
#define SWAP4(x) as_uint(as_uchar4(x).wzyx)
|
||||
#define SWAP8(x) as_ulong(as_uchar8(x).s76543210)
|
||||
|
||||
inline uint LODWORD(const sph_u64 x)
|
||||
{
|
||||
return (uint)(x & 0xFFFFFFFF);
|
||||
}
|
||||
|
||||
inline uint HIDWORD(const sph_u64 x)
|
||||
{
|
||||
return (uint)(x >> 32);
|
||||
}
|
||||
|
||||
// Tables for function F
|
||||
__constant static sph_u64 T0[256] = {
|
||||
0xE6F87E5C5B711FD0, 0x258377800924FA16, 0xC849E07E852EA4A8, 0x5B4686A18F06C16A,
|
||||
0x0B32E9A2D77B416E, 0xABDA37A467815C66, 0xF61796A81A686676, 0xF5DC0B706391954B,
|
||||
0x4862F38DB7E64BF1, 0xFF5C629A68BD85C5, 0xCB827DA6FCD75795, 0x66D36DAF69B9F089,
|
||||
0x356C9F74483D83B0, 0x7CBCECB1238C99A1, 0x36A702AC31C4708D, 0x9EB6A8D02FBCDFD6,
|
||||
0x8B19FA51E5B3AE37, 0x9CCFB5408A127D0B, 0xBC0C78B508208F5A, 0xE533E3842288ECED,
|
||||
0xCEC2C7D377C15FD2, 0xEC7817B6505D0F5E, 0xB94CC2C08336871D, 0x8C205DB4CB0B04AD,
|
||||
0x763C855B28A0892F, 0x588D1B79F6FF3257, 0x3FECF69E4311933E, 0x0FC0D39F803A18C9,
|
||||
0xEE010A26F5F3AD83, 0x10EFE8F4411979A6, 0x5DCDA10C7DE93A10, 0x4A1BEE1D1248E92C,
|
||||
0x53BFF2DB21847339, 0xB4F50CCFA6A23D09, 0x5FB4BC9CD84798CD, 0xE88A2D8B071C56F9,
|
||||
0x7F7771695A756A9C, 0xC5F02E71A0BA1EBC, 0xA663F9AB4215E672, 0x2EB19E22DE5FBB78,
|
||||
0x0DB9CE0F2594BA14, 0x82520E6397664D84, 0x2F031E6A0208EA98, 0x5C7F2144A1BE6BF0,
|
||||
0x7A37CB1CD16362DB, 0x83E08E2B4B311C64, 0xCF70479BAB960E32, 0x856BA986B9DEE71E,
|
||||
0xB5478C877AF56CE9, 0xB8FE42885F61D6FD, 0x1BDD0156966238C8, 0x622157923EF8A92E,
|
||||
0xFC97FF42114476F8, 0x9D7D350856452CEB, 0x4C90C9B0E0A71256, 0x2308502DFBCB016C,
|
||||
0x2D7A03FAA7A64845, 0xF46E8B38BFC6C4AB, 0xBDBEF8FDD477DEBA, 0x3AAC4CEBC8079B79,
|
||||
0xF09CB105E8879D0C, 0x27FA6A10AC8A58CB, 0x8960E7C1401D0CEA, 0x1A6F811E4A356928,
|
||||
0x90C4FB0773D196FF, 0x43501A2F609D0A9F, 0xF7A516E0C63F3796, 0x1CE4A6B3B8DA9252,
|
||||
0x1324752C38E08A9B, 0xA5A864733BEC154F, 0x2BF124575549B33F, 0xD766DB15440DC5C7,
|
||||
0xA7D179E39E42B792, 0xDADF151A61997FD3, 0x86A0345EC0271423, 0x38D5517B6DA939A4,
|
||||
0x6518F077104003B4, 0x02791D90A5AEA2DD, 0x88D267899C4A5D0A, 0x930F66DF0A2865C2,
|
||||
0x4EE9D4204509B08B, 0x325538916685292A, 0x412907BFC533A842, 0xB27E2B62544DC673,
|
||||
0x6C5304456295E007, 0x5AF406E95351908A, 0x1F2F3B6BC123616F, 0xC37B09DC5255E5C6,
|
||||
0x3967D133B1FE6844, 0x298839C7F0E711E2, 0x409B87F71964F9A2, 0xE938ADC3DB4B0719,
|
||||
0x0C0B4E47F9C3EBF4, 0x5534D576D36B8843, 0x4610A05AEB8B02D8, 0x20C3CDF58232F251,
|
||||
0x6DE1840DBEC2B1E7, 0xA0E8DE06B0FA1D08, 0x7B854B540D34333B, 0x42E29A67BCCA5B7F,
|
||||
0xD8A6088AC437DD0E, 0xC63BB3A9D943ED81, 0x21714DBD5E65A3B1, 0x6761EDE7B5EEA169,
|
||||
0x2431F7C8D573ABF6, 0xD51FC685E1A3671A, 0x5E063CD40410C92D, 0x283AB98F2CB04002,
|
||||
0x8FEBC06CB2F2F790, 0x17D64F116FA1D33C, 0xE07359F1A99EE4AA, 0x784ED68C74CDC006,
|
||||
0x6E2A19D5C73B42DA, 0x8712B4161C7045C3, 0x371582E4ED93216D, 0xACE390414939F6FC,
|
||||
0x7EC5F12186223B7C, 0xC0B094042BAC16FB, 0xF9D745379A527EBF, 0x737C3F2EA3B68168,
|
||||
0x33E7B8D9BAD278CA, 0xA9A32A34C22FFEBB, 0xE48163CCFEDFBD0D, 0x8E5940246EA5A670,
|
||||
0x51C6EF4B842AD1E4, 0x22BAD065279C508C, 0xD91488C218608CEE, 0x319EA5491F7CDA17,
|
||||
0xD394E128134C9C60, 0x094BF43272D5E3B3, 0x9BF612A5A4AAD791, 0xCCBBDA43D26FFD0F,
|
||||
0x34DE1F3C946AD250, 0x4F5B5468995EE16B, 0xDF9FAF6FEA8F7794, 0x2648EA5870DD092B,
|
||||
0xBFC7E56D71D97C67, 0xDDE6B2FF4F21D549, 0x3C276B463AE86003, 0x91767B4FAF86C71F,
|
||||
0x68A13E7835D4B9A0, 0xB68C115F030C9FD4, 0x141DD2C916582001, 0x983D8F7DDD5324AC,
|
||||
0x64AA703FCC175254, 0xC2C989948E02B426, 0x3E5E76D69F46C2DE, 0x50746F03587D8004,
|
||||
0x45DB3D829272F1E5, 0x60584A029B560BF3, 0xFBAE58A73FFCDC62, 0xA15A5E4E6CAD4CE8,
|
||||
0x4BA96E55CE1FB8CC, 0x08F9747AAE82B253, 0xC102144CF7FB471B, 0x9F042898F3EB8E36,
|
||||
0x068B27ADF2EFFB7A, 0xEDCA97FE8C0A5EBE, 0x778E0513F4F7D8CF, 0x302C2501C32B8BF7,
|
||||
0x8D92DDFC175C554D, 0xF865C57F46052F5F, 0xEAF3301BA2B2F424, 0xAA68B7ECBBD60D86,
|
||||
0x998F0F350104754C, 0x0000000000000000, 0xF12E314D34D0CCEC, 0x710522BE061823B5,
|
||||
0xAF280D9930C005C1, 0x97FD5CE25D693C65, 0x19A41CC633CC9A15, 0x95844172F8C79EB8,
|
||||
0xDC5432B7937684A9, 0x9436C13A2490CF58, 0x802B13F332C8EF59, 0xC442AE397CED4F5C,
|
||||
0xFA1CD8EFE3AB8D82, 0xF2E5AC954D293FD1, 0x6AD823E8907A1B7D, 0x4D2249F83CF043B6,
|
||||
0x03CB9DD879F9F33D, 0xDE2D2F2736D82674, 0x2A43A41F891EE2DF, 0x6F98999D1B6C133A,
|
||||
0xD4AD46CD3DF436FA, 0xBB35DF50269825C0, 0x964FDCAA813E6D85, 0xEB41B0537EE5A5C4,
|
||||
0x0540BA758B160847, 0xA41AE43BE7BB44AF, 0xE3B8C429D0671797, 0x819993BBEE9FBEB9,
|
||||
0xAE9A8DD1EC975421, 0xF3572CDD917E6E31, 0x6393D7DAE2AFF8CE, 0x47A2201237DC5338,
|
||||
0xA32343DEC903EE35, 0x79FC56C4A89A91E6, 0x01B28048DC5751E0, 0x1296F564E4B7DB7B,
|
||||
0x75F7188351597A12, 0xDB6D9552BDCE2E33, 0x1E9DBB231D74308F, 0x520D7293FDD322D9,
|
||||
0xE20A44610C304677, 0xFEEEE2D2B4EAD425, 0xCA30FDEE20800675, 0x61EACA4A47015A13,
|
||||
0xE74AFE1487264E30, 0x2CC883B27BF119A5, 0x1664CF59B3F682DC, 0xA811AA7C1E78AF5B,
|
||||
0x1D5626FB648DC3B2, 0xB73E9117DF5BCE34, 0xD05F7CF06AB56F5D, 0xFD257F0ACD132718,
|
||||
0x574DC8E676C52A9E, 0x0739A7E52EB8AA9A, 0x5486553E0F3CD9A3, 0x56FF48AEAA927B7E,
|
||||
0xBE756525AD8E2D87, 0x7D0E6CF9FFDBC841, 0x3B1ECCA31450CA99, 0x6913BE30E983E840,
|
||||
0xAD511009956EA71C, 0xB1B5B6BA2DB4354E, 0x4469BDCA4E25A005, 0x15AF5281CA0F71E1,
|
||||
0x744598CB8D0E2BF2, 0x593F9B312AA863B7, 0xEFB38A6E29A4FC63, 0x6B6AA3A04C2D4A9D,
|
||||
0x3D95EB0EE6BF31E3, 0xA291C3961554BFD5, 0x18169C8EEF9BCBF5, 0x115D68BC9D4E2846,
|
||||
0xBA875F18FACF7420, 0xD1EDFCB8B6E23EBD, 0xB00736F2F1E364AE, 0x84D929CE6589B6FE,
|
||||
0x70B7A2F6DA4F7255, 0x0E7253D75C6D4929, 0x04F23A3D574159A7, 0x0A8069EA0B2C108E,
|
||||
0x49D073C56BB11A11, 0x8AAB7A1939E4FFD7, 0xCD095A0B0E38ACEF, 0xC9FB60365979F548,
|
||||
0x92BDE697D67F3422, 0xC78933E10514BC61, 0xE1C1D9B975C9B54A, 0xD2266160CF1BCD80,
|
||||
0x9A4492ED78FD8671, 0xB3CCAB2A881A9793, 0x72CEBF667FE1D088, 0xD6D45B5D985A9427
|
||||
};
|
||||
__constant static sph_u64 T1[256] = {
|
||||
0xC811A8058C3F55DE, 0x65F5B43196B50619, 0xF74F96B1D6706E43, 0x859D1E8BCB43D336,
|
||||
0x5AAB8A85CCFA3D84, 0xF9C7BF99C295FCFD, 0xA21FD5A1DE4B630F, 0xCDB3EF763B8B456D,
|
||||
0x803F59F87CF7C385, 0xB27C73BE5F31913C, 0x98E3AC6633B04821, 0xBF61674C26B8F818,
|
||||
0x0FFBC995C4C130C8, 0xAAA0862010761A98, 0x6057F342210116AA, 0xF63C760C0654CC35,
|
||||
0x2DDB45CC667D9042, 0xBCF45A964BD40382, 0x68E8A0C3EF3C6F3D, 0xA7BD92D269FF73BC,
|
||||
0x290AE20201ED2287, 0xB7DE34CDE885818F, 0xD901EEA7DD61059B, 0xD6FA273219A03553,
|
||||
0xD56F1AE874CCCEC9, 0xEA31245C2E83F554, 0x7034555DA07BE499, 0xCE26D2AC56E7BEF7,
|
||||
0xFD161857A5054E38, 0x6A0E7DA4527436D1, 0x5BD86A381CDE9FF2, 0xCAF7756231770C32,
|
||||
0xB09AAED9E279C8D0, 0x5DEF1091C60674DB, 0x111046A2515E5045, 0x23536CE4729802FC,
|
||||
0xC50CBCF7F5B63CFA, 0x73A16887CD171F03, 0x7D2941AFD9F28DBD, 0x3F5E3EB45A4F3B9D,
|
||||
0x84EEFE361B677140, 0x3DB8E3D3E7076271, 0x1A3A28F9F20FD248, 0x7EBC7C75B49E7627,
|
||||
0x74E5F293C7EB565C, 0x18DCF59E4F478BA4, 0x0C6EF44FA9ADCB52, 0xC699812D98DAC760,
|
||||
0x788B06DC6E469D0E, 0xFC65F8EA7521EC4E, 0x30A5F7219E8E0B55, 0x2BEC3F65BCA57B6B,
|
||||
0xDDD04969BAF1B75E, 0x99904CDBE394EA57, 0x14B201D1E6EA40F6, 0xBBB0C08241284ADD,
|
||||
0x50F20463BF8F1DFF, 0xE8D7F93B93CBACB8, 0x4D8CB68E477C86E8, 0xC1DD1B3992268E3F,
|
||||
0x7C5AA11209D62FCB, 0x2F3D98ABDB35C9AE, 0x671369562BFD5FF5, 0x15C1E16C36CEE280,
|
||||
0x1D7EB2EDF8F39B17, 0xDA94D37DB00DFE01, 0x877BC3EC760B8ADA, 0xCB8495DFE153AE44,
|
||||
0x05A24773B7B410B3, 0x12857B783C32ABDF, 0x8EB770D06812513B, 0x536739B9D2E3E665,
|
||||
0x584D57E271B26468, 0xD789C78FC9849725, 0xA935BBFA7D1AE102, 0x8B1537A3DFA64188,
|
||||
0xD0CD5D9BC378DE7A, 0x4AC82C9A4D80CFB7, 0x42777F1B83BDB620, 0x72D2883A1D33BD75,
|
||||
0x5E7A2D4BAB6A8F41, 0xF4DAAB6BBB1C95D9, 0x905CFFE7FD8D31B6, 0x83AA6422119B381F,
|
||||
0xC0AEFB8442022C49, 0xA0F908C663033AE3, 0xA428AF0804938826, 0xADE41C341A8A53C7,
|
||||
0xAE7121EE77E6A85D, 0xC47F5C4A25929E8C, 0xB538E9AA55CDD863, 0x06377AA9DAD8EB29,
|
||||
0xA18AE87BB3279895, 0x6EDFDA6A35E48414, 0x6B7D9D19825094A7, 0xD41CFA55A4E86CBF,
|
||||
0xE5CAEDC9EA42C59C, 0xA36C351C0E6FC179, 0x5181E4DE6FABBF89, 0xFFF0C530184D17D4,
|
||||
0x9D41EB1584045892, 0x1C0D525028D73961, 0xF178EC180CA8856A, 0x9A0571018EF811CD,
|
||||
0x4091A27C3EF5EFCC, 0x19AF15239F6329D2, 0x347450EFF91EB990, 0xE11B4A078DD27759,
|
||||
0xB9561DE5FC601331, 0x912F1F5A2DA993C0, 0x1654DCB65BA2191A, 0x3E2DDE098A6B99EB,
|
||||
0x8A66D71E0F82E3FE, 0x8C51ADB7D55A08D7, 0x4533E50F8941FF7F, 0x02E6DD67BD4859EC,
|
||||
0xE068AABA5DF6D52F, 0xC24826E3FF4A75A5, 0x6C39070D88ACDDF8, 0x6486548C4691A46F,
|
||||
0xD1BEBD26135C7C0C, 0xB30F93038F15334A, 0x82D9849FC1BF9A69, 0x9C320BA85420FAE4,
|
||||
0xFA528243AFF90767, 0x9ED4D6CFE968A308, 0xB825FD582C44B147, 0x9B7691BC5EDCB3BB,
|
||||
0xC7EA619048FE6516, 0x1063A61F817AF233, 0x47D538683409A693, 0x63C2CE984C6DED30,
|
||||
0x2A9FDFD86C81D91D, 0x7B1E3B06032A6694, 0x666089EBFBD9FD83, 0x0A598EE67375207B,
|
||||
0x07449A140AFC495F, 0x2CA8A571B6593234, 0x1F986F8A45BBC2FB, 0x381AA4A050B372C2,
|
||||
0x5423A3ADD81FAF3A, 0x17273C0B8B86BB6C, 0xFE83258DC869B5A2, 0x287902BFD1C980F1,
|
||||
0xF5A94BD66B3837AF, 0x88800A79B2CABA12, 0x55504310083B0D4C, 0xDF36940E07B9EEB2,
|
||||
0x04D1A7CE6790B2C5, 0x612413FFF125B4DC, 0x26F12B97C52C124F, 0x86082351A62F28AC,
|
||||
0xEF93632F9937E5E7, 0x3507B052293A1BE6, 0xE72C30AE570A9C70, 0xD3586041AE1425E0,
|
||||
0xDE4574B3D79D4CC4, 0x92BA228040C5685A, 0xF00B0CA5DC8C271C, 0xBE1287F1F69C5A6E,
|
||||
0xF39E317FB1E0DC86, 0x495D114020EC342D, 0x699B407E3F18CD4B, 0xDCA3A9D46AD51528,
|
||||
0x0D1D14F279896924, 0x0000000000000000, 0x593EB75FA196C61E, 0x2E4E78160B116BD8,
|
||||
0x6D4AE7B058887F8E, 0xE65FD013872E3E06, 0x7A6DDBBBD30EC4E2, 0xAC97FC89CAAEF1B1,
|
||||
0x09CCB33C1E19DBE1, 0x89F3EAC462EE1864, 0x7770CF49AA87ADC6, 0x56C57ECA6557F6D6,
|
||||
0x03953DDA6D6CFB9A, 0x36928D884456E07C, 0x1EEB8F37959F608D, 0x31D6179C4EAAA923,
|
||||
0x6FAC3AD7E5C02662, 0x43049FA653991456, 0xABD3669DC052B8EE, 0xAF02C153A7C20A2B,
|
||||
0x3CCB036E3723C007, 0x93C9C23D90E1CA2C, 0xC33BC65E2F6ED7D3, 0x4CFF56339758249E,
|
||||
0xB1E94E64325D6AA6, 0x37E16D359472420A, 0x79F8E661BE623F78, 0x5214D90402C74413,
|
||||
0x482EF1FDF0C8965B, 0x13F69BC5EC1609A9, 0x0E88292814E592BE, 0x4E198B542A107D72,
|
||||
0xCCC00FCBEBAFE71B, 0x1B49C844222B703E, 0x2564164DA840E9D5, 0x20C6513E1FF4F966,
|
||||
0xBAC3203F910CE8AB, 0xF2EDD1C261C47EF0, 0x814CB945ACD361F3, 0x95FEB8944A392105,
|
||||
0x5C9CF02C1622D6AD, 0x971865F3F77178E9, 0xBD87BA2B9BF0A1F4, 0x444005B259655D09,
|
||||
0xED75BE48247FBC0B, 0x7596122E17CFF42A, 0xB44B091785E97A15, 0x966B854E2755DA9F,
|
||||
0xEEE0839249134791, 0x32432A4623C652B9, 0xA8465B47AD3E4374, 0xF8B45F2412B15E8B,
|
||||
0x2417F6F078644BA3, 0xFB2162FE7FDDA511, 0x4BBBCC279DA46DC1, 0x0173E0BDD024A276,
|
||||
0x22208C59A2BCA08A, 0x8FC4906DB836F34D, 0xE4B90D743A6667EA, 0x7147B5E0705F46EF,
|
||||
0x2782CB2A1508B039, 0xEC065EF5F45B1E7D, 0x21B5B183CFD05B10, 0xDBE733C060295C77,
|
||||
0x9FA73672394C017E, 0xCF55321186C31C81, 0xD8720E1A0D45A7ED, 0x3B8F997A3DDF8958,
|
||||
0x3AFC79C7EDFB2B2E, 0xE9A4198643EF0ECE, 0x5F09CDF67B4E2D37, 0x4F6A6BE9FA34DF04,
|
||||
0xB6ADD47038A123F9, 0x8D224D0A057EAAA1, 0xC96248B85C1BF7A8, 0xE3FD9760309A2EB5,
|
||||
0x0B2A6E5BA351820D, 0xEB42C4E1FEA75722, 0x948D58299A1D8373, 0x7FCF9CC864BAD451,
|
||||
0xA55B4FB5D4B72A50, 0x08BF5381CE3D7997, 0x46A6D8D5E42D04E5, 0xD22B80FC7E308796,
|
||||
0x57B69E77B57354A0, 0x3969441D8097D0B4, 0x3330CAFBF3E2F0CF, 0xE28E77DDE0BE8CC3,
|
||||
0x62B12E259C494F46, 0xA6CE726FB9DBD1CA, 0x41E242C1EED14DBA, 0x76032FF47AA30FB0
|
||||
};
|
||||
__constant static sph_u64 T2[256] = {
|
||||
0x45B268A93ACDE4CC, 0xAF7F0BE884549D08, 0x048354B3C1468263, 0x925435C2C80EFED2,
|
||||
0xEE4E37F27FDFFBA7, 0x167A33920C60F14D, 0xFB123B52EA03E584, 0x4A0CAB53FDBB9007,
|
||||
0x9DEAF6380F788A19, 0xCB48EC558F0CB32A, 0xB59DC4B2D6FEF7E0, 0xDCDBCA22F4F3ECB6,
|
||||
0x11DF5813549A9C40, 0xE33FDEDF568ACED3, 0xA0C1C8124322E9C3, 0x07A56B8158FA6D0D,
|
||||
0x77279579B1E1F3DD, 0xD9B18B74422AC004, 0xB8EC2D9FFFABC294, 0xF4ACF8A82D75914F,
|
||||
0x7BBF69B1EF2B6878, 0xC4F62FAF487AC7E1, 0x76CE809CC67E5D0C, 0x6711D88F92E4C14C,
|
||||
0x627B99D9243DEDFE, 0x234AA5C3DFB68B51, 0x909B1F15262DBF6D, 0x4F66EA054B62BCB5,
|
||||
0x1AE2CF5A52AA6AE8, 0xBEA053FBD0CE0148, 0xED6808C0E66314C9, 0x43FE16CD15A82710,
|
||||
0xCD049231A06970F6, 0xE7BC8A6C97CC4CB0, 0x337CE835FCB3B9C0, 0x65DEF2587CC780F3,
|
||||
0x52214EDE4132BB50, 0x95F15E4390F493DF, 0x870839625DD2E0F1, 0x41313C1AFB8B66AF,
|
||||
0x91720AF051B211BC, 0x477D427ED4EEA573, 0x2E3B4CEEF6E3BE25, 0x82627834EB0BCC43,
|
||||
0x9C03E3DD78E724C8, 0x2877328AD9867DF9, 0x14B51945E243B0F2, 0x574B0F88F7EB97E2,
|
||||
0x88B6FA989AA4943A, 0x19C4F068CB168586, 0x50EE6409AF11FAEF, 0x7DF317D5C04EABA4,
|
||||
0x7A567C5498B4C6A9, 0xB6BBFB804F42188E, 0x3CC22BCF3BC5CD0B, 0xD04336EAAA397713,
|
||||
0xF02FAC1BEC33132C, 0x2506DBA7F0D3488D, 0xD7E65D6BF2C31A1E, 0x5EB9B2161FF820F5,
|
||||
0x842E0650C46E0F9F, 0x716BEB1D9E843001, 0xA933758CAB315ED4, 0x3FE414FDA2792265,
|
||||
0x27C9F1701EF00932, 0x73A4C1CA70A771BE, 0x94184BA6E76B3D0E, 0x40D829FF8C14C87E,
|
||||
0x0FBEC3FAC77674CB, 0x3616A9634A6A9572, 0x8F139119C25EF937, 0xF545ED4D5AEA3F9E,
|
||||
0xE802499650BA387B, 0x6437E7BD0B582E22, 0xE6559F89E053E261, 0x80AD52E305288DFC,
|
||||
0x6DC55A23E34B9935, 0xDE14E0F51AD0AD09, 0xC6390578A659865E, 0x96D7617109487CB1,
|
||||
0xE2D6CB3A21156002, 0x01E915E5779FAED1, 0xADB0213F6A77DCB7, 0x9880B76EB9A1A6AB,
|
||||
0x5D9F8D248644CF9B, 0xFD5E4536C5662658, 0xF1C6B9FE9BACBDFD, 0xEACD6341BE9979C4,
|
||||
0xEFA7221708405576, 0x510771ECD88E543E, 0xC2BA51CB671F043D, 0x0AD482AC71AF5879,
|
||||
0xFE787A045CDAC936, 0xB238AF338E049AED, 0xBD866CC94972EE26, 0x615DA6EBBD810290,
|
||||
0x3295FDD08B2C1711, 0xF834046073BF0AEA, 0xF3099329758FFC42, 0x1CAEB13E7DCFA934,
|
||||
0xBA2307481188832B, 0x24EFCE42874CE65C, 0x0E57D61FB0E9DA1A, 0xB3D1BAD6F99B343C,
|
||||
0xC0757B1C893C4582, 0x2B510DB8403A9297, 0x5C7698C1F1DB614A, 0x3E0D0118D5E68CB4,
|
||||
0xD60F488E855CB4CF, 0xAE961E0DF3CB33D9, 0x3A8E55AB14A00ED7, 0x42170328623789C1,
|
||||
0x838B6DD19C946292, 0x895FEF7DED3B3AEB, 0xCFCBB8E64E4A3149, 0x064C7E642F65C3DC,
|
||||
0x3D2B3E2A4C5A63DA, 0x5BD3F340A9210C47, 0xB474D157A1615931, 0xAC5934DA1DE87266,
|
||||
0x6EE365117AF7765B, 0xC86ED36716B05C44, 0x9BA6885C201D49C5, 0xB905387A88346C45,
|
||||
0x131072C4BAB9DDFF, 0xBF49461EA751AF99, 0xD52977BC1CE05BA1, 0xB0F785E46027DB52,
|
||||
0x546D30BA6E57788C, 0x305AD707650F56AE, 0xC987C682612FF295, 0xA5AB8944F5FBC571,
|
||||
0x7ED528E759F244CA, 0x8DDCBBCE2C7DB888, 0xAA154ABE328DB1BA, 0x1E619BE993ECE88B,
|
||||
0x09F2BD9EE813B717, 0x7401AA4B285D1CB3, 0x21858F143195CAEE, 0x48C381841398D1B8,
|
||||
0xFCB750D3B2F98889, 0x39A86A998D1CE1B9, 0x1F888E0CE473465A, 0x7899568376978716,
|
||||
0x02CF2AD7EE2341BF, 0x85C713B5B3F1A14E, 0xFF916FE12B4567E7, 0x7C1A0230B7D10575,
|
||||
0x0C98FCC85ECA9BA5, 0xA3E7F720DA9E06AD, 0x6A6031A2BBB1F438, 0x973E74947ED7D260,
|
||||
0x2CF4663918C0FF9A, 0x5F50A7F368678E24, 0x34D983B4A449D4CD, 0x68AF1B755592B587,
|
||||
0x7F3C3D022E6DEA1B, 0xABFC5F5B45121F6B, 0x0D71E92D29553574, 0xDFFDF5106D4F03D8,
|
||||
0x081BA87B9F8C19C6, 0xDB7EA1A3AC0981BB, 0xBBCA12AD66172DFA, 0x79704366010829C7,
|
||||
0x179326777BFF5F9C, 0x0000000000000000, 0xEB2476A4C906D715, 0x724DD42F0738DF6F,
|
||||
0xB752EE6538DDB65F, 0x37FFBC863DF53BA3, 0x8EFA84FCB5C157E6, 0xE9EB5C73272596AA,
|
||||
0x1B0BDABF2535C439, 0x86E12C872A4D4E20, 0x9969A28BCE3E087A, 0xFAFB2EB79D9C4B55,
|
||||
0x056A4156B6D92CB2, 0x5A3AE6A5DEBEA296, 0x22A3B026A8292580, 0x53C85B3B36AD1581,
|
||||
0xB11E900117B87583, 0xC51F3A4A3FE56930, 0xE019E1EDCF3621BD, 0xEC811D2591FCBA18,
|
||||
0x445B7D4C4D524A1D, 0xA8DA6069DCAEF005, 0x58F5CC72309DE329, 0xD4C062596B7FF570,
|
||||
0xCE22AD0339D59F98, 0x591CD99747024DF8, 0x8B90C5AA03187B54, 0xF663D27FC356D0F0,
|
||||
0xD8589E9135B56ED5, 0x35309651D3D67A1C, 0x12F96721CD26732E, 0xD28C1C3D441A36AC,
|
||||
0x492A946164077F69, 0x2D1D73DC6F5F514B, 0x6F0A70F40D68D88A, 0x60B4B30ECA1EAC41,
|
||||
0xD36509D83385987D, 0x0B3D97490630F6A8, 0x9ECCC90A96C46577, 0xA20EE2C5AD01A87C,
|
||||
0xE49AB55E0E70A3DE, 0xA4429CA182646BA0, 0xDA97B446DB962F6A, 0xCCED87D4D7F6DE27,
|
||||
0x2AB8185D37A53C46, 0x9F25DCEFE15BCBA6, 0xC19C6EF9FEA3EB53, 0xA764A3931BD884CE,
|
||||
0x2FD2590B817C10F4, 0x56A21A6D80743933, 0xE573A0BB79EF0D0F, 0x155C0CA095DC1E23,
|
||||
0x6C2C4FC694D437E4, 0x10364DF623053291, 0xDD32DFC7836C4267, 0x03263F3299BCEF6E,
|
||||
0x66F8CD6AE57B6F9D, 0x8C35AE2B5BE21659, 0x31B3C2E21290F87F, 0x93BD2027BF915003,
|
||||
0x69460E90220D1B56, 0x299E276FAE19D328, 0x63928C3C53A2432F, 0x7082FEF8E91B9ED0,
|
||||
0xBC6F792C3EED40F7, 0x4C40D537D2DE53DB, 0x75E8BFAE5FC2B262, 0x4DA9C0D2A541FD0A,
|
||||
0x4E8FFFE03CFD1264, 0x2620E495696FA7E3, 0xE1F0F408B8A98F6C, 0xD1AA230FDDA6D9C2,
|
||||
0xC7D0109DD1C6288F, 0x8A79D04F7487D585, 0x4694579BA3710BA2, 0x38417F7CFA834F68,
|
||||
0x1D47A4DB0A5007E5, 0x206C9AF1460A643F, 0xA128DDF734BD4712, 0x8144470672B7232D,
|
||||
0xF2E086CC02105293, 0x182DE58DBC892B57, 0xCAA1F9B0F8931DFB, 0x6B892447CC2E5AE9,
|
||||
0xF9DD11850420A43B, 0x4BE5BEB68A243ED6, 0x5584255F19C8D65D, 0x3B67404E633FA006,
|
||||
0xA68DB6766C472A1F, 0xF78AC79AB4C97E21, 0xC353442E1080AAEC, 0x9A4F9DB95782E714
|
||||
};
|
||||
__constant static sph_u64 T3[256] = {
|
||||
0x05BA7BC82C9B3220, 0x31A54665F8B65E4F, 0xB1B651F77547F4D4, 0x8BFA0D857BA46682,
|
||||
0x85A96C5AA16A98BB, 0x990FAEF908EB79C9, 0xA15E37A247F4A62D, 0x76857DCD5D27741E,
|
||||
0xF8C50B800A1820BC, 0xBE65DCB201F7A2B4, 0x666D1B986F9426E7, 0x4CC921BF53C4E648,
|
||||
0x95410A0F93D9CA42, 0x20CDCCAA647BA4EF, 0x429A4060890A1871, 0x0C4EA4F69B32B38B,
|
||||
0xCCDA362DDE354CD3, 0x96DC23BC7C5B2FA9, 0xC309BB68AA851AB3, 0xD26131A73648E013,
|
||||
0x021DC52941FC4DB2, 0xCD5ADAB7704BE48A, 0xA77965D984ED71E6, 0x32386FD61734BBA4,
|
||||
0xE82D6DD538AB7245, 0x5C2147EA6177B4B1, 0x5DA1AB70CF091CE8, 0xAC907FCE72B8BDFF,
|
||||
0x57C85DFD972278A8, 0xA4E44C6A6B6F940D, 0x3851995B4F1FDFE4, 0x62578CCAED71BC9E,
|
||||
0xD9882BB0C01D2C0A, 0x917B9D5D113C503B, 0xA2C31E11A87643C6, 0xE463C923A399C1CE,
|
||||
0xF71686C57EA876DC, 0x87B4A973E096D509, 0xAF0D567D9D3A5814, 0xB40C2A3F59DCC6F4,
|
||||
0x3602F88495D121DD, 0xD3E1DD3D9836484A, 0xF945E71AA46688E5, 0x7518547EB2A591F5,
|
||||
0x9366587450C01D89, 0x9EA81018658C065B, 0x4F54080CBC4603A3, 0x2D0384C65137BF3D,
|
||||
0xDC325078EC861E2A, 0xEA30A8FC79573FF7, 0x214D2030CA050CB6, 0x65F0322B8016C30C,
|
||||
0x69BE96DD1B247087, 0xDB95EE9981E161B8, 0xD1FC1814D9CA05F8, 0x820ED2BBCC0DE729,
|
||||
0x63D76050430F14C7, 0x3BCCB0E8A09D3A0F, 0x8E40764D573F54A2, 0x39D175C1E16177BD,
|
||||
0x12F5A37C734F1F4B, 0xAB37C12F1FDFC26D, 0x5648B167395CD0F1, 0x6C04ED1537BF42A7,
|
||||
0xED97161D14304065, 0x7D6C67DAAB72B807, 0xEC17FA87BA4EE83C, 0xDFAF79CB0304FBC1,
|
||||
0x733F060571BC463E, 0x78D61C1287E98A27, 0xD07CF48E77B4ADA1, 0xB9C262536C90DD26,
|
||||
0xE2449B5860801605, 0x8FC09AD7F941FCFB, 0xFAD8CEA94BE46D0E, 0xA343F28B0608EB9F,
|
||||
0x9B126BD04917347B, 0x9A92874AE7699C22, 0x1B017C42C4E69EE0, 0x3A4C5C720EE39256,
|
||||
0x4B6E9F5E3EA399DA, 0x6BA353F45AD83D35, 0xE7FEE0904C1B2425, 0x22D009832587E95D,
|
||||
0x842980C00F1430E2, 0xC6B3C0A0861E2893, 0x087433A419D729F2, 0x341F3DADD42D6C6F,
|
||||
0xEE0A3FAEFBB2A58E, 0x4AEE73C490DD3183, 0xAAB72DB5B1A16A34, 0xA92A04065E238FDF,
|
||||
0x7B4B35A1686B6FCC, 0x6A23BF6EF4A6956C, 0x191CB96B851AD352, 0x55D598D4D6DE351A,
|
||||
0xC9604DE5F2AE7EF3, 0x1CA6C2A3A981E172, 0xDE2F9551AD7A5398, 0x3025AAFF56C8F616,
|
||||
0x15521D9D1E2860D9, 0x506FE31CFA45073A, 0x189C55F12B647B0B, 0x0180EC9AAE7EA859,
|
||||
0x7CEC8B40050C105E, 0x2350E5198BF94104, 0xEF8AD33455CC0DD7, 0x07A7BEE16D677F92,
|
||||
0xE5E325B90DE76997, 0x5A061591A26E637A, 0xB611EF1618208B46, 0x09F4DF3EB7A981AB,
|
||||
0x1EBB078AE87DACC0, 0xB791038CB65E231F, 0x0FD38D4574B05660, 0x67EDF702C1EA8EBE,
|
||||
0xBA5F4BE0831238CD, 0xE3C477C2CEFEBE5C, 0x0DCE486C354C1BD2, 0x8C5DB36416C31910,
|
||||
0x26EA9ED1A7627324, 0x039D29B3EF82E5EB, 0x9F28FC82CBF2AE02, 0xA8AAE89CF05D2786,
|
||||
0x431AACFA2774B028, 0xCF471F9E31B7A938, 0x581BD0B8E3922EC8, 0xBC78199B400BEF06,
|
||||
0x90FB71C7BF42F862, 0x1F3BEB1046030499, 0x683E7A47B55AD8DE, 0x988F4263A695D190,
|
||||
0xD808C72A6E638453, 0x0627527BC319D7CB, 0xEBB04466D72997AE, 0xE67E0C0AE2658C7C,
|
||||
0x14D2F107B056C880, 0x7122C32C30400B8C, 0x8A7AE11FD5DACEDB, 0xA0DEDB38E98A0E74,
|
||||
0xAD109354DCC615A6, 0x0BE91A17F655CC19, 0x8DDD5FFEB8BDB149, 0xBFE53028AF890AED,
|
||||
0xD65BA6F5B4AD7A6A, 0x7956F0882997227E, 0x10E8665532B352F9, 0x0E5361DFDACEFE39,
|
||||
0xCEC7F3049FC90161, 0xFF62B561677F5F2E, 0x975CCF26D22587F0, 0x51EF0F86543BAF63,
|
||||
0x2F1E41EF10CBF28F, 0x52722635BBB94A88, 0xAE8DBAE73344F04D, 0x410769D36688FD9A,
|
||||
0xB3AB94DE34BBB966, 0x801317928DF1AA9B, 0xA564A0F0C5113C54, 0xF131D4BEBDB1A117,
|
||||
0x7F71A2F3EA8EF5B5, 0x40878549C8F655C3, 0x7EF14E6944F05DEC, 0xD44663DCF55137D8,
|
||||
0xF2ACFD0D523344FC, 0x0000000000000000, 0x5FBC6E598EF5515A, 0x16CF342EF1AA8532,
|
||||
0xB036BD6DDB395C8D, 0x13754FE6DD31B712, 0xBBDFA77A2D6C9094, 0x89E7C8AC3A582B30,
|
||||
0x3C6B0E09CDFA459D, 0xC4AE0589C7E26521, 0x49735A777F5FD468, 0xCAFD64561D2C9B18,
|
||||
0xDA1502032F9FC9E1, 0x8867243694268369, 0x3782141E3BAF8984, 0x9CB5D53124704BE9,
|
||||
0xD7DB4A6F1AD3D233, 0xA6F989432A93D9BF, 0x9D3539AB8A0EE3B0, 0x53F2CAAF15C7E2D1,
|
||||
0x6E19283C76430F15, 0x3DEBE2936384EDC4, 0x5E3C82C3208BF903, 0x33B8834CB94A13FD,
|
||||
0x6470DEB12E686B55, 0x359FD1377A53C436, 0x61CAA57902F35975, 0x043A975282E59A79,
|
||||
0xFD7F70482683129C, 0xC52EE913699CCD78, 0x28B9FF0E7DAC8D1D, 0x5455744E78A09D43,
|
||||
0xCB7D88CCB3523341, 0x44BD121B4A13CFBA, 0x4D49CD25FDBA4E11, 0x3E76CB208C06082F,
|
||||
0x3FF627BA2278A076, 0xC28957F204FBB2EA, 0x453DFE81E46D67E3, 0x94C1E6953DA7621B,
|
||||
0x2C83685CFF491764, 0xF32C1197FC4DECA5, 0x2B24D6BD922E68F6, 0xB22B78449AC5113F,
|
||||
0x48F3B6EDD1217C31, 0x2E9EAD75BEB55AD6, 0x174FD8B45FD42D6B, 0x4ED4E4961238ABFA,
|
||||
0x92E6B4EEFEBEB5D0, 0x46A0D7320BEF8208, 0x47203BA8A5912A51, 0x24F75BF8E69E3E96,
|
||||
0xF0B1382413CF094E, 0xFEE259FBC901F777, 0x276A724B091CDB7D, 0xBDF8F501EE75475F,
|
||||
0x599B3C224DEC8691, 0x6D84018F99C1EAFE, 0x7498B8E41CDB39AC, 0xE0595E71217C5BB7,
|
||||
0x2AA43A273C50C0AF, 0xF50B43EC3F543B6E, 0x838E3E2162734F70, 0xC09492DB4507FF58,
|
||||
0x72BFEA9FDFC2EE67, 0x11688ACF9CCDFAA0, 0x1A8190D86A9836B9, 0x7ACBD93BC615C795,
|
||||
0xC7332C3A286080CA, 0x863445E94EE87D50, 0xF6966A5FD0D6DE85, 0xE9AD814F96D5DA1C,
|
||||
0x70A22FB69E3EA3D5, 0x0A69F68D582B6440, 0xB8428EC9C2EE757F, 0x604A49E3AC8DF12C,
|
||||
0x5B86F90B0C10CB23, 0xE1D9B2EB8F02F3EE, 0x29391394D3D22544, 0xC8E0A17F5CD0D6AA,
|
||||
0xB58CC6A5F7A26EAD, 0x8193FB08238F02C2, 0xD5C68F465B2F9F81, 0xFCFF9CD288FDBAC5,
|
||||
0x77059157F359DC47, 0x1D262E3907FF492B, 0xFB582233E59AC557, 0xDDB2BCE242F8B673,
|
||||
0x2577B76248E096CF, 0x6F99C4A6D83DA74C, 0xC1147E41EB795701, 0xF48BAF76912A9337
|
||||
};
|
||||
__constant static sph_u64 T4[256] = {
|
||||
0x3EF29D249B2C0A19, 0xE9E16322B6F8622F, 0x5536994047757F7A, 0x9F4D56D5A47B0B33,
|
||||
0x822567466AA1174C, 0xB8F5057DEB082FB2, 0xCC48C10BF4475F53, 0x373088D4275DEC3A,
|
||||
0x968F4325180AED10, 0x173D232CF7016151, 0xAE4ED09F946FCC13, 0xFD4B4741C4539873,
|
||||
0x1B5B3F0DD9933765, 0x2FFCB0967B644052, 0xE02376D20A89840C, 0xA3AE3A70329B18D7,
|
||||
0x419CBD2335DE8526, 0xFAFEBF115B7C3199, 0x0397074F85AA9B0D, 0xC58AD4FB4836B970,
|
||||
0xBEC60BE3FC4104A8, 0x1EFF36DC4B708772, 0x131FDC33ED8453B6, 0x0844E33E341764D3,
|
||||
0x0FF11B6EAB38CD39, 0x64351F0A7761B85A, 0x3B5694F509CFBA0E, 0x30857084B87245D0,
|
||||
0x47AFB3BD2297AE3C, 0xF2BA5C2F6F6B554A, 0x74BDC4761F4F70E1, 0xCFDFC64471EDC45E,
|
||||
0xE610784C1DC0AF16, 0x7ACA29D63C113F28, 0x2DED411776A859AF, 0xAC5F211E99A3D5EE,
|
||||
0xD484F949A87EF33B, 0x3CE36CA596E013E4, 0xD120F0983A9D432C, 0x6BC40464DC597563,
|
||||
0x69D5F5E5D1956C9E, 0x9AE95F043698BB24, 0xC9ECC8DA66A4EF44, 0xD69508C8A5B2EAC6,
|
||||
0xC40C2235C0503B80, 0x38C193BA8C652103, 0x1CEEC75D46BC9E8F, 0xD331011937515AD1,
|
||||
0xD8E2E56886ECA50F, 0xB137108D5779C991, 0x709F3B6905CA4206, 0x4FEB50831680CAEF,
|
||||
0xEC456AF3241BD238, 0x58D673AFE181ABBE, 0x242F54E7CAD9BF8C, 0x0211F1810DCC19FD,
|
||||
0x90BC4DBB0F43C60A, 0x9518446A9DA0761D, 0xA1BFCBF13F57012A, 0x2BDE4F8961E172B5,
|
||||
0x27B853A84F732481, 0xB0B1E643DF1F4B61, 0x18CC38425C39AC68, 0xD2B7F7D7BF37D821,
|
||||
0x3103864A3014C720, 0x14AA246372ABFA5C, 0x6E600DB54EBAC574, 0x394765740403A3F3,
|
||||
0x09C215F0BC71E623, 0x2A58B947E987F045, 0x7B4CDF18B477BDD8, 0x9709B5EB906C6FE0,
|
||||
0x73083C268060D90B, 0xFEDC400E41F9037E, 0x284948C6E44BE9B8, 0x728ECAE808065BFB,
|
||||
0x06330E9E17492B1A, 0x5950856169E7294E, 0xBAE4F4FCE6C4364F, 0xCA7BCF95E30E7449,
|
||||
0x7D7FD186A33E96C2, 0x52836110D85AD690, 0x4DFAA1021B4CD312, 0x913ABB75872544FA,
|
||||
0xDD46ECB9140F1518, 0x3D659A6B1E869114, 0xC23F2CABD719109A, 0xD713FE062DD46836,
|
||||
0xD0A60656B2FBC1DC, 0x221C5A79DD909496, 0xEFD26DBCA1B14935, 0x0E77EDA0235E4FC9,
|
||||
0xCBFD395B6B68F6B9, 0x0DE0EAEFA6F4D4C4, 0x0422FF1F1A8532E7, 0xF969B85EDED6AA94,
|
||||
0x7F6E2007AEF28F3F, 0x3AD0623B81A938FE, 0x6624EE8B7AADA1A7, 0xB682E8DDC856607B,
|
||||
0xA78CC56F281E2A30, 0xC79B257A45FAA08D, 0x5B4174E0642B30B3, 0x5F638BFF7EAE0254,
|
||||
0x4BC9AF9C0C05F808, 0xCE59308AF98B46AE, 0x8FC58DA9CC55C388, 0x803496C7676D0EB1,
|
||||
0xF33CAAE1E70DD7BA, 0xBB6202326EA2B4BF, 0xD5020F87201871CB, 0x9D5CA754A9B712CE,
|
||||
0x841669D87DE83C56, 0x8A6184785EB6739F, 0x420BBA6CB0741E2B, 0xF12D5B60EAC1CE47,
|
||||
0x76AC35F71283691C, 0x2C6BB7D9FECEDB5F, 0xFCCDB18F4C351A83, 0x1F79C012C3160582,
|
||||
0xF0ABADAE62A74CB7, 0xE1A5801C82EF06FC, 0x67A21845F2CB2357, 0x5114665F5DF04D9D,
|
||||
0xBF40FD2D74278658, 0xA0393D3FB73183DA, 0x05A409D192E3B017, 0xA9FB28CF0B4065F9,
|
||||
0x25A9A22942BF3D7C, 0xDB75E22703463E02, 0xB326E10C5AB5D06C, 0xE7968E8295A62DE6,
|
||||
0xB973F3B3636EAD42, 0xDF571D3819C30CE5, 0xEE549B7229D7CBC5, 0x12992AFD65E2D146,
|
||||
0xF8EF4E9056B02864, 0xB7041E134030E28B, 0xC02EDD2ADAD50967, 0x932B4AF48AE95D07,
|
||||
0x6FE6FB7BC6DC4784, 0x239AACB755F61666, 0x401A4BEDBDB807D6, 0x485EA8D389AF6305,
|
||||
0xA41BC220ADB4B13D, 0x753B32B89729F211, 0x997E584BB3322029, 0x1D683193CEDA1C7F,
|
||||
0xFF5AB6C0C99F818E, 0x16BBD5E27F67E3A1, 0xA59D34EE25D233CD, 0x98F8AE853B54A2D9,
|
||||
0x6DF70AFACB105E79, 0x795D2E99B9BBA425, 0x8E437B6744334178, 0x0186F6CE886682F0,
|
||||
0xEBF092A3BB347BD2, 0xBCD7FA62F18D1D55, 0xADD9D7D011C5571E, 0x0BD3E471B1BDFFDE,
|
||||
0xAA6C2F808EEAFEF4, 0x5EE57D31F6C880A4, 0xF50FA47FF044FCA0, 0x1ADDC9C351F5B595,
|
||||
0xEA76646D3352F922, 0x0000000000000000, 0x85909F16F58EBEA6, 0x46294573AAF12CCC,
|
||||
0x0A5512BF39DB7D2E, 0x78DBD85731DD26D5, 0x29CFBE086C2D6B48, 0x218B5D36583A0F9B,
|
||||
0x152CD2ADFACD78AC, 0x83A39188E2C795BC, 0xC3B9DA655F7F926A, 0x9ECBA01B2C1D89C3,
|
||||
0x07B5F8509F2FA9EA, 0x7EE8D6C926940DCF, 0x36B67E1AAF3B6ECA, 0x86079859702425AB,
|
||||
0xFB7849DFD31AB369, 0x4C7C57CC932A51E2, 0xD96413A60E8A27FF, 0x263EA566C715A671,
|
||||
0x6C71FC344376DC89, 0x4A4F595284637AF8, 0xDAF314E98B20BCF2, 0x572768C14AB96687,
|
||||
0x1088DB7C682EC8BB, 0x887075F9537A6A62, 0x2E7A4658F302C2A2, 0x619116DBE582084D,
|
||||
0xA87DDE018326E709, 0xDCC01A779C6997E8, 0xEDC39C3DAC7D50C8, 0xA60A33A1A078A8C0,
|
||||
0xC1A82BE452B38B97, 0x3F746BEA134A88E9, 0xA228CCBEBAFD9A27, 0xABEAD94E068C7C04,
|
||||
0xF48952B178227E50, 0x5CF48CB0FB049959, 0x6017E0156DE48ABD, 0x4438B4F2A73D3531,
|
||||
0x8C528AE649FF5885, 0xB515EF924DFCFB76, 0x0C661C212E925634, 0xB493195CC59A7986,
|
||||
0x9CDA519A21D1903E, 0x32948105B5BE5C2D, 0x194ACE8CD45F2E98, 0x438D4CA238129CDB,
|
||||
0x9B6FA9CABEFE39D4, 0x81B26009EF0B8C41, 0xDED1EBF691A58E15, 0x4E6DA64D9EE6481F,
|
||||
0x54B06F8ECF13FD8A, 0x49D85E1D01C9E1F5, 0xAFC826511C094EE3, 0xF698A33075EE67AD,
|
||||
0x5AC7822EEC4DB243, 0x8DD47C28C199DA75, 0x89F68337DB1CE892, 0xCDCE37C57C21DDA3,
|
||||
0x530597DE503C5460, 0x6A42F2AA543FF793, 0x5D727A7E73621BA9, 0xE232875307459DF1,
|
||||
0x56A19E0FC2DFE477, 0xC61DD3B4CD9C227D, 0xE5877F03986A341B, 0x949EB2A415C6F4ED,
|
||||
0x6206119460289340, 0x6380E75AE84E11B0, 0x8BE772B6D6D0F16F, 0x50929091D596CF6D,
|
||||
0xE86795EC3E9EE0DF, 0x7CF927482B581432, 0xC86A3E14EEC26DB4, 0x7119CDA78DACC0F6,
|
||||
0xE40189CD100CB6EB, 0x92ADBC3A028FDFF7, 0xB2A017C2D2D3529C, 0x200DABF8D05C8D6B,
|
||||
0x34A78F9BA2F77737, 0xE3B4719D8F231F01, 0x45BE423C2F5BB7C1, 0xF71E55FEFD88E55D,
|
||||
0x6853032B59F3EE6E, 0x65B3E9C4FF073AAA, 0x772AC3399AE5EBEC, 0x87816E97F842A75B,
|
||||
0x110E2DB2E0484A4B, 0x331277CB3DD8DEDD, 0xBD510CAC79EB9FA5, 0x352179552A91F5C7
|
||||
};
|
||||
__constant static sph_u64 T5[256] = {
|
||||
0x8AB0A96846E06A6D, 0x43C7E80B4BF0B33A, 0x08C9B3546B161EE5, 0x39F1C235EBA990BE,
|
||||
0xC1BEF2376606C7B2, 0x2C209233614569AA, 0xEB01523B6FC3289A, 0x946953AB935ACEDD,
|
||||
0x272838F63E13340E, 0x8B0455ECA12BA052, 0x77A1B2C4978FF8A2, 0xA55122CA13E54086,
|
||||
0x2276135862D3F1CD, 0xDB8DDFDE08B76CFE, 0x5D1E12C89E4A178A, 0x0E56816B03969867,
|
||||
0xEE5F79953303ED59, 0xAFED748BAB78D71D, 0x6D929F2DF93E53EE, 0xF5D8A8F8BA798C2A,
|
||||
0xF619B1698E39CF6B, 0x95DDAF2F749104E2, 0xEC2A9C80E0886427, 0xCE5C8FD8825B95EA,
|
||||
0xC4E0D9993AC60271, 0x4699C3A5173076F9, 0x3D1B151F50A29F42, 0x9ED505EA2BC75946,
|
||||
0x34665ACFDC7F4B98, 0x61B1FB53292342F7, 0xC721C0080E864130, 0x8693CD1696FD7B74,
|
||||
0x872731927136B14B, 0xD3446C8A63A1721B, 0x669A35E8A6680E4A, 0xCAB658F239509A16,
|
||||
0xA4E5DE4EF42E8AB9, 0x37A7435EE83F08D9, 0x134E6239E26C7F96, 0x82791A3C2DF67488,
|
||||
0x3F6EF00A8329163C, 0x8E5A7E42FDEB6591, 0x5CAAEE4C7981DDB5, 0x19F234785AF1E80D,
|
||||
0x255DDDE3ED98BD70, 0x50898A32A99CCCAC, 0x28CA4519DA4E6656, 0xAE59880F4CB31D22,
|
||||
0x0D9798FA37D6DB26, 0x32F968F0B4FFCD1A, 0xA00F09644F258545, 0xFA3AD5175E24DE72,
|
||||
0xF46C547C5DB24615, 0x713E80FBFF0F7E20, 0x7843CF2B73D2AAFA, 0xBD17EA36AEDF62B4,
|
||||
0xFD111BACD16F92CF, 0x4ABAA7DBC72D67E0, 0xB3416B5DAD49FAD3, 0xBCA316B24914A88B,
|
||||
0x15D150068AECF914, 0xE27C1DEBE31EFC40, 0x4FE48C759BEDA223, 0x7EDCFD141B522C78,
|
||||
0x4E5070F17C26681C, 0xE696CAC15815F3BC, 0x35D2A64B3BB481A7, 0x800CFF29FE7DFDF6,
|
||||
0x1ED9FAC3D5BAA4B0, 0x6C2663A91EF599D1, 0x03C1199134404341, 0xF7AD4DED69F20554,
|
||||
0xCD9D9649B61BD6AB, 0xC8C3BDE7EADB1368, 0xD131899FB02AFB65, 0x1D18E352E1FAE7F1,
|
||||
0xDA39235AEF7CA6C1, 0xA1BBF5E0A8EE4F7A, 0x91377805CF9A0B1E, 0x3138716180BF8E5B,
|
||||
0xD9F83ACBDB3CE580, 0x0275E515D38B897E, 0x472D3F21F0FBBCC6, 0x2D946EB7868EA395,
|
||||
0xBA3C248D21942E09, 0xE7223645BFDE3983, 0xFF64FEB902E41BB1, 0xC97741630D10D957,
|
||||
0xC3CB1722B58D4ECC, 0xA27AEC719CAE0C3B, 0x99FECB51A48C15FB, 0x1465AC826D27332B,
|
||||
0xE1BD047AD75EBF01, 0x79F733AF941960C5, 0x672EC96C41A3C475, 0xC27FEBA6524684F3,
|
||||
0x64EFD0FD75E38734, 0xED9E60040743AE18, 0xFB8E2993B9EF144D, 0x38453EB10C625A81,
|
||||
0x6978480742355C12, 0x48CF42CE14A6EE9E, 0x1CAC1FD606312DCE, 0x7B82D6BA4792E9BB,
|
||||
0x9D141C7B1F871A07, 0x5616B80DC11C4A2E, 0xB849C198F21FA777, 0x7CA91801C8D9A506,
|
||||
0xB1348E487EC273AD, 0x41B20D1E987B3A44, 0x7460AB55A3CFBBE3, 0x84E628034576F20A,
|
||||
0x1B87D16D897A6173, 0x0FE27DEFE45D5258, 0x83CDE6B8CA3DBEB7, 0x0C23647ED01D1119,
|
||||
0x7A362A3EA0592384, 0xB61F40F3F1893F10, 0x75D457D1440471DC, 0x4558DA34237035B8,
|
||||
0xDCA6116587FC2043, 0x8D9B67D3C9AB26D0, 0x2B0B5C88EE0E2517, 0x6FE77A382AB5DA90,
|
||||
0x269CC472D9D8FE31, 0x63C41E46FAA8CB89, 0xB7ABBC771642F52F, 0x7D1DE4852F126F39,
|
||||
0xA8C6BA3024339BA0, 0x600507D7CEE888C8, 0x8FEE82C61A20AFAE, 0x57A2448926D78011,
|
||||
0xFCA5E72836A458F0, 0x072BCEBB8F4B4CBD, 0x497BBE4AF36D24A1, 0x3CAFE99BB769557D,
|
||||
0x12FA9EBD05A7B5A9, 0xE8C04BAA5B836BDB, 0x4273148FAC3B7905, 0x908384812851C121,
|
||||
0xE557D3506C55B0FD, 0x72FF996ACB4F3D61, 0x3EDA0C8E64E2DC03, 0xF0868356E6B949E9,
|
||||
0x04EAD72ABB0B0FFC, 0x17A4B5135967706A, 0xE3C8E16F04D5367F, 0xF84F30028DAF570C,
|
||||
0x1846C8FCBD3A2232, 0x5B8120F7F6CA9108, 0xD46FA231ECEA3EA6, 0x334D947453340725,
|
||||
0x58403966C28AD249, 0xBED6F3A79A9F21F5, 0x68CCB483A5FE962D, 0xD085751B57E1315A,
|
||||
0xFED0023DE52FD18E, 0x4B0E5B5F20E6ADDF, 0x1A332DE96EB1AB4C, 0xA3CE10F57B65C604,
|
||||
0x108F7BA8D62C3CD7, 0xAB07A3A11073D8E1, 0x6B0DAD1291BED56C, 0xF2F366433532C097,
|
||||
0x2E557726B2CEE0D4, 0x0000000000000000, 0xCB02A476DE9B5029, 0xE4E32FD48B9E7AC2,
|
||||
0x734B65EE2C84F75E, 0x6E5386BCCD7E10AF, 0x01B4FC84E7CBCA3F, 0xCFE8735C65905FD5,
|
||||
0x3613BFDA0FF4C2E6, 0x113B872C31E7F6E8, 0x2FE18BA255052AEB, 0xE974B72EBC48A1E4,
|
||||
0x0ABC5641B89D979B, 0xB46AA5E62202B66E, 0x44EC26B0C4BBFF87, 0xA6903B5B27A503C7,
|
||||
0x7F680190FC99E647, 0x97A84A3AA71A8D9C, 0xDD12EDE16037EA7C, 0xC554251DDD0DC84E,
|
||||
0x88C54C7D956BE313, 0x4D91696048662B5D, 0xB08072CC9909B992, 0xB5DE5962C5C97C51,
|
||||
0x81B803AD19B637C9, 0xB2F597D94A8230EC, 0x0B08AAC55F565DA4, 0xF1327FD2017283D6,
|
||||
0xAD98919E78F35E63, 0x6AB9519676751F53, 0x24E921670A53774F, 0xB9FD3D1C15D46D48,
|
||||
0x92F66194FBDA485F, 0x5A35DC7311015B37, 0xDED3F4705477A93D, 0xC00A0EB381CD0D8D,
|
||||
0xBB88D809C65FE436, 0x16104997BEACBA55, 0x21B70AC95693B28C, 0x59F4C5E225411876,
|
||||
0xD5DB5EB50B21F499, 0x55D7A19CF55C096F, 0xA97246B4C3F8519F, 0x8552D487A2BD3835,
|
||||
0x54635D181297C350, 0x23C2EFDC85183BF2, 0x9F61F96ECC0C9379, 0x534893A39DDC8FED,
|
||||
0x5EDF0B59AA0A54CB, 0xAC2C6D1A9F38945C, 0xD7AEBBA0D8AA7DE7, 0x2ABFA00C09C5EF28,
|
||||
0xD84CC64F3CF72FBF, 0x2003F64DB15878B3, 0xA724C7DFC06EC9F8, 0x069F323F68808682,
|
||||
0xCC296ACD51D01C94, 0x055E2BAE5CC0C5C3, 0x6270E2C21D6301B6, 0x3B842720382219C0,
|
||||
0xD2F0900E846AB824, 0x52FC6F277A1745D2, 0xC6953C8CE94D8B0F, 0xE009F8FE3095753E,
|
||||
0x655B2C7992284D0B, 0x984A37D54347DFC4, 0xEAB5AEBF8808E2A5, 0x9A3FD2C090CC56BA,
|
||||
0x9CA0E0FFF84CD038, 0x4C2595E4AFADE162, 0xDF6708F4B3BC6302, 0xBF620F237D54EBCA,
|
||||
0x93429D101C118260, 0x097D4FD08CDDD4DA, 0x8C2F9B572E60ECEF, 0x708A7C7F18C4B41F,
|
||||
0x3A30DBA4DFE9D3FF, 0x4006F19A7FB0F07B, 0x5F6BF7DD4DC19EF4, 0x1F6D064732716E8F,
|
||||
0xF9FBCC866A649D33, 0x308C8DE567744464, 0x8971B0F972A0292C, 0xD61A47243F61B7D8,
|
||||
0xEFEB8511D4C82766, 0x961CB6BE40D147A3, 0xAAB35F25F7B812DE, 0x76154E407044329D,
|
||||
0x513D76B64E570693, 0xF3479AC7D2F90AA8, 0x9B8B2E4477079C85, 0x297EB99D3D85AC69
|
||||
};
|
||||
__constant static sph_u64 T6[256] = {
|
||||
0x7E37E62DFC7D40C3, 0x776F25A4EE939E5B, 0xE045C850DD8FB5AD, 0x86ED5BA711FF1952,
|
||||
0xE91D0BD9CF616B35, 0x37E0AB256E408FFB, 0x9607F6C031025A7A, 0x0B02F5E116D23C9D,
|
||||
0xF3D8486BFB50650C, 0x621CFF27C40875F5, 0x7D40CB71FA5FD34A, 0x6DAA6616DAA29062,
|
||||
0x9F5F354923EC84E2, 0xEC847C3DC507C3B3, 0x025A3668043CE205, 0xA8BF9E6C4DAC0B19,
|
||||
0xFA808BE2E9BEBB94, 0xB5B99C5277C74FA3, 0x78D9BC95F0397BCC, 0xE332E50CDBAD2624,
|
||||
0xC74FCE129332797E, 0x1729ECEB2EA709AB, 0xC2D6B9F69954D1F8, 0x5D898CBFBAB8551A,
|
||||
0x859A76FB17DD8ADB, 0x1BE85886362F7FB5, 0xF6413F8FF136CD8A, 0xD3110FA5BBB7E35C,
|
||||
0x0A2FEED514CC4D11, 0xE83010EDCD7F1AB9, 0xA1E75DE55F42D581, 0xEEDE4A55C13B21B6,
|
||||
0xF2F5535FF94E1480, 0x0CC1B46D1888761E, 0xBCE15FDB6529913B, 0x2D25E8975A7181C2,
|
||||
0x71817F1CE2D7A554, 0x2E52C5CB5C53124B, 0xF9F7A6BEEF9C281D, 0x9E722E7D21F2F56E,
|
||||
0xCE170D9B81DCA7E6, 0x0E9B82051CB4941B, 0x1E712F623C49D733, 0x21E45CFA42F9F7DC,
|
||||
0xCB8E7A7F8BBA0F60, 0x8E98831A010FB646, 0x474CCF0D8E895B23, 0xA99285584FB27A95,
|
||||
0x8CC2B57205335443, 0x42D5B8E984EFF3A5, 0x012D1B34021E718C, 0x57A6626AAE74180B,
|
||||
0xFF19FC06E3D81312, 0x35BA9D4D6A7C6DFE, 0xC9D44C178F86ED65, 0x506523E6A02E5288,
|
||||
0x03772D5C06229389, 0x8B01F4FE0B691EC0, 0xF8DABD8AED825991, 0x4C4E3AEC985B67BE,
|
||||
0xB10DF0827FBF96A9, 0x6A69279AD4F8DAE1, 0xE78689DCD3D5FF2E, 0x812E1A2B1FA553D1,
|
||||
0xFBAD90D6EBA0CA18, 0x1AC543B234310E39, 0x1604F7DF2CB97827, 0xA6241C6951189F02,
|
||||
0x753513CCEAAF7C5E, 0x64F2A59FC84C4EFA, 0x247D2B1E489F5F5A, 0xDB64D718AB474C48,
|
||||
0x79F4A7A1F2270A40, 0x1573DA832A9BEBAE, 0x3497867968621C72, 0x514838D2A2302304,
|
||||
0xF0AF6537FD72F685, 0x1D06023E3A6B44BA, 0x678588C3CE6EDD73, 0x66A893F7CC70ACFF,
|
||||
0xD4D24E29B5EDA9DF, 0x3856321470EA6A6C, 0x07C3418C0E5A4A83, 0x2BCBB22F5635BACD,
|
||||
0x04B46CD00878D90A, 0x06EE5AB80C443B0F, 0x3B211F4876C8F9E5, 0x0958C38912EEDE98,
|
||||
0xD14B39CDBF8B0159, 0x397B292072F41BE0, 0x87C0409313E168DE, 0xAD26E98847CAA39F,
|
||||
0x4E140C849C6785BB, 0xD5FF551DB7F3D853, 0xA0CA46D15D5CA40D, 0xCD6020C787FE346F,
|
||||
0x84B76DCF15C3FB57, 0xDEFDA0FCA121E4CE, 0x4B8D7B6096012D3D, 0x9AC642AD298A2C64,
|
||||
0x0875D8BD10F0AF14, 0xB357C6EA7B8374AC, 0x4D6321D89A451632, 0xEDA96709C719B23F,
|
||||
0xF76C24BBF328BC06, 0xC662D526912C08F2, 0x3CE25EC47892B366, 0xB978283F6F4F39BD,
|
||||
0xC08C8F9E9D6833FD, 0x4F3917B09E79F437, 0x593DE06FB2C08C10, 0xD6887841B1D14BDA,
|
||||
0x19B26EEE32139DB0, 0xB494876675D93E2F, 0x825937771987C058, 0x90E9AC783D466175,
|
||||
0xF1827E03FF6C8709, 0x945DC0A8353EB87F, 0x4516F9658AB5B926, 0x3F9573987EB020EF,
|
||||
0xB855330B6D514831, 0x2AE6A91B542BCB41, 0x6331E413C6160479, 0x408F8E8180D311A0,
|
||||
0xEFF35161C325503A, 0xD06622F9BD9570D5, 0x8876D9A20D4B8D49, 0xA5533135573A0C8B,
|
||||
0xE168D364DF91C421, 0xF41B09E7F50A2F8F, 0x12B09B0F24C1A12D, 0xDA49CC2CA9593DC4,
|
||||
0x1F5C34563E57A6BF, 0x54D14F36A8568B82, 0xAF7CDFE043F6419A, 0xEA6A2685C943F8BC,
|
||||
0xE5DCBFB4D7E91D2B, 0xB27ADDDE799D0520, 0x6B443CAED6E6AB6D, 0x7BAE91C9F61BE845,
|
||||
0x3EB868AC7CAE5163, 0x11C7B65322E332A4, 0xD23C1491B9A992D0, 0x8FB5982E0311C7CA,
|
||||
0x70AC6428E0C9D4D8, 0x895BC2960F55FCC5, 0x76423E90EC8DEFD7, 0x6FF0507EDE9E7267,
|
||||
0x3DCF45F07A8CC2EA, 0x4AA06054941F5CB1, 0x5810FB5BB0DEFD9C, 0x5EFEA1E3BC9AC693,
|
||||
0x6EDD4B4ADC8003EB, 0x741808F8E8B10DD2, 0x145EC1B728859A22, 0x28BC9F7350172944,
|
||||
0x270A06424EBDCCD3, 0x972AEDF4331C2BF6, 0x059977E40A66A886, 0x2550302A4A812ED6,
|
||||
0xDD8A8DA0A7037747, 0xC515F87A970E9B7B, 0x3023EAA9601AC578, 0xB7E3AA3A73FBADA6,
|
||||
0x0FB699311EAAE597, 0x0000000000000000, 0x310EF19D6204B4F4, 0x229371A644DB6455,
|
||||
0x0DECAF591A960792, 0x5CA4978BB8A62496, 0x1C2B190A38753536, 0x41A295B582CD602C,
|
||||
0x3279DCC16426277D, 0xC1A194AA9F764271, 0x139D803B26DFD0A1, 0xAE51C4D441E83016,
|
||||
0xD813FA44AD65DFC1, 0xAC0BF2BC45D4D213, 0x23BE6A9246C515D9, 0x49D74D08923DCF38,
|
||||
0x9D05032127D066E7, 0x2F7FDEFF5E4D63C7, 0xA47E2A0155247D07, 0x99B16FF12FA8BFED,
|
||||
0x4661D4398C972AAF, 0xDFD0BBC8A33F9542, 0xDCA79694A51D06CB, 0xB020EBB67DA1E725,
|
||||
0xBA0F0563696DAA34, 0xE4F1A480D5F76CA7, 0xC438E34E9510EAF7, 0x939E81243B64F2FC,
|
||||
0x8DEFAE46072D25CF, 0x2C08F3A3586FF04E, 0xD7A56375B3CF3A56, 0x20C947CE40E78650,
|
||||
0x43F8A3DD86F18229, 0x568B795EAC6A6987, 0x8003011F1DBB225D, 0xF53612D3F7145E03,
|
||||
0x189F75DA300DEC3C, 0x9570DB9C3720C9F3, 0xBB221E576B73DBB8, 0x72F65240E4F536DD,
|
||||
0x443BE25188ABC8AA, 0xE21FFE38D9B357A8, 0xFD43CA6EE7E4F117, 0xCAA3614B89A47EEC,
|
||||
0xFE34E732E1C6629E, 0x83742C431B99B1D4, 0xCF3A16AF83C2D66A, 0xAAE5A8044990E91C,
|
||||
0x26271D764CA3BD5F, 0x91C4B74C3F5810F9, 0x7C6DD045F841A2C6, 0x7F1AFD19FE63314F,
|
||||
0xC8F957238D989CE9, 0xA709075D5306EE8E, 0x55FC5402AA48FA0E, 0x48FA563C9023BEB4,
|
||||
0x65DFBEABCA523F76, 0x6C877D22D8BCE1EE, 0xCC4D3BF385E045E3, 0xBEBB69B36115733E,
|
||||
0x10EAAD6720FD4328, 0xB6CEB10E71E5DC2A, 0xBDCC44EF6737E0B7, 0x523F158EA412B08D,
|
||||
0x989C74C52DB6CE61, 0x9BEB59992B945DE8, 0x8A2CEFCA09776F4C, 0xA3BD6B8D5B7E3784,
|
||||
0xEB473DB1CB5D8930, 0xC3FBA2C29B4AA074, 0x9C28181525CE176B, 0x683311F2D0C438E4,
|
||||
0x5FD3BAD7BE84B71F, 0xFC6ED15AE5FA809B, 0x36CDB0116C5EFE77, 0x29918447520958C8,
|
||||
0xA29070B959604608, 0x53120EBAA60CC101, 0x3A0C047C74D68869, 0x691E0AC6D2DA4968,
|
||||
0x73DB4974E6EB4751, 0x7A838AFDF40599C9, 0x5A4ACD33B4E21F99, 0x6046C94FC03497F0,
|
||||
0xE6AB92E8D1CB8EA2, 0x3354C7F5663856F1, 0xD93EE170AF7BAE4D, 0x616BD27BC22AE67C,
|
||||
0x92B39A10397A8370, 0xABC8B3304B8E9890, 0xBF967287630B02B2, 0x5B67D607B6FC6E15
|
||||
};
|
||||
__constant static sph_u64 T7[256] = {
|
||||
0xD031C397CE553FE6, 0x16BA5B01B006B525, 0xA89BADE6296E70C8, 0x6A1F525D77D3435B,
|
||||
0x6E103570573DFA0B, 0x660EFB2A17FC95AB, 0x76327A9E97634BF6, 0x4BAD9D6462458BF5,
|
||||
0xF1830CAEDBC3F748, 0xC5C8F542669131FF, 0x95044A1CDC48B0CB, 0x892962DF3CF8B866,
|
||||
0xB0B9E208E930C135, 0xA14FB3F0611A767C, 0x8D2605F21C160136, 0xD6B71922FECC549E,
|
||||
0x37089438A5907D8B, 0x0B5DA38E5803D49C, 0x5A5BCC9CEA6F3CBC, 0xEDAE246D3B73FFE5,
|
||||
0xD2B87E0FDE22EDCE, 0x5E54ABB1CA8185EC, 0x1DE7F88FE80561B9, 0xAD5E1A870135A08C,
|
||||
0x2F2ADBD665CECC76, 0x5780B5A782F58358, 0x3EDC8A2EEDE47B3F, 0xC9D95C3506BEE70F,
|
||||
0x83BE111D6C4E05EE, 0xA603B90959367410, 0x103C81B4809FDE5D, 0x2C69B6027D0C774A,
|
||||
0x399080D7D5C87953, 0x09D41E16487406B4, 0xCDD63B1826505E5F, 0xF99DC2F49B0298E8,
|
||||
0x9CD0540A943CB67F, 0xBCA84B7F891F17C5, 0x723D1DB3B78DF2A6, 0x78AA6E71E73B4F2E,
|
||||
0x1433E699A071670D, 0x84F21BE454620782, 0x98DF3327B4D20F2F, 0xF049DCE2D3769E5C,
|
||||
0xDB6C60199656EB7A, 0x648746B2078B4783, 0x32CD23598DCBADCF, 0x1EA4955BF0C7DA85,
|
||||
0xE9A143401B9D46B5, 0xFD92A5D9BBEC21B8, 0xC8138C790E0B8E1B, 0x2EE00B9A6D7BA562,
|
||||
0xF85712B893B7F1FC, 0xEB28FED80BEA949D, 0x564A65EB8A40EA4C, 0x6C9988E8474A2823,
|
||||
0x4535898B121D8F2D, 0xABD8C03231ACCBF4, 0xBA2E91CAB9867CBD, 0x7960BE3DEF8E263A,
|
||||
0x0C11A977602FD6F0, 0xCB50E1AD16C93527, 0xEAE22E94035FFD89, 0x2866D12F5DE2CE1A,
|
||||
0xFF1B1841AB9BF390, 0x9F9339DE8CFE0D43, 0x964727C8C48A0BF7, 0x524502C6AAAE531C,
|
||||
0x9B9C5EF3AC10B413, 0x4FA2FA4942AB32A5, 0x3F165A62E551122B, 0xC74148DA76E6E3D7,
|
||||
0x924840E5E464B2A7, 0xD372AE43D69784DA, 0x233B72A105E11A86, 0xA48A04914941A638,
|
||||
0xB4B68525C9DE7865, 0xDDEABAACA6CF8002, 0x0A9773C250B6BD88, 0xC284FFBB5EBD3393,
|
||||
0x8BA0DF472C8F6A4E, 0x2AEF6CB74D951C32, 0x427983722A318D41, 0x73F7CDFFBF389BB2,
|
||||
0x074C0AF9382C026C, 0x8A6A0F0B243A035A, 0x6FDAE53C5F88931F, 0xC68B98967E538AC3,
|
||||
0x44FF59C71AA8E639, 0xE2FCE0CE439E9229, 0xA20CDE2479D8CD40, 0x19E89FA2C8EBD8E9,
|
||||
0xF446BBCFF398270C, 0x43B3533E2284E455, 0xD82F0DCD8E945046, 0x51066F12B26CE820,
|
||||
0xE73957AF6BC5426D, 0x081ECE5A40C16FA0, 0x3B193D4FC5BFAB7B, 0x7FE66488DF174D42,
|
||||
0x0E9814EF705804D8, 0x8137AC857C39D7C6, 0xB1733244E185A821, 0x695C3F896F11F867,
|
||||
0xF6CF0657E3EFF524, 0x1AABF276D02963D5, 0x2DA3664E75B91E5E, 0x0289BD981077D228,
|
||||
0x90C1FD7DF413608F, 0x3C5537B6FD93A917, 0xAA12107E3919A2E0, 0x0686DAB530996B78,
|
||||
0xDAA6B0559EE3826E, 0xC34E2FF756085A87, 0x6D5358A44FFF4137, 0xFC587595B35948AC,
|
||||
0x7CA5095CC7D5F67E, 0xFB147F6C8B754AC0, 0xBFEB26AB91DDACF9, 0x6896EFC567A49173,
|
||||
0xCA9A31E11E7C5C33, 0xBBE44186B13315A9, 0x0DDB793B689ABFE4, 0x70B4A02BA7FA208E,
|
||||
0xE47A3A7B7307F951, 0x8CECD5BE14A36822, 0xEEED49B923B144D9, 0x17708B4DB8B3DC31,
|
||||
0x6088219F2765FED3, 0xB3FA8FDCF1F27A09, 0x910B2D31FCA6099B, 0x0F52C4A378ED6DCC,
|
||||
0x50CCBF5EBAD98134, 0x6BD582117F662A4F, 0x94CE9A50D4FDD9DF, 0x2B25BCFB45207526,
|
||||
0x67C42B661F49FCBF, 0x492420FC723259DD, 0x03436DD418C2BB3C, 0x1F6E4517F872B391,
|
||||
0xA08563BC69AF1F68, 0xD43EA4BAEEBB86B6, 0x01CAD04C08B56914, 0xAC94CACB0980C998,
|
||||
0x54C3D8739A373864, 0x26FEC5C02DBACAC2, 0xDEA9D778BE0D3B3E, 0x040F672D20EEB950,
|
||||
0xE5B0EA377BB29045, 0xF30AB136CBB42560, 0x62019C0737122CFB, 0xE86B930C13282FA1,
|
||||
0xCC1CEB542EE5374B, 0x538FD28AA21B3A08, 0x1B61223AD89C0AC1, 0x36C24474AD25149F,
|
||||
0x7A23D3E9F74C9D06, 0xBE21F6E79968C5ED, 0xCF5F868036278C77, 0xF705D61BEB5A9C30,
|
||||
0x4D2B47D152DCE08D, 0x5F9E7BFDC234ECF8, 0x247778583DCD18EA, 0x867BA67C4415D5AA,
|
||||
0x4CE1979D5A698999, 0x0000000000000000, 0xEC64F42133C696F1, 0xB57C5569C16B1171,
|
||||
0xC1C7926F467F88AF, 0x654D96FE0F3E2E97, 0x15F936D5A8C40E19, 0xB8A72C52A9F1AE95,
|
||||
0xA9517DAA21DB19DC, 0x58D27104FA18EE94, 0x5918A148F2AD8780, 0x5CDD1629DAF657C4,
|
||||
0x8274C15164FB6CFA, 0xD1FB13DBC6E056F2, 0x7D6FD910CF609F6A, 0xB63F38BDD9A9AA4D,
|
||||
0x3D9FE7FAF526C003, 0x74BBC706871499DE, 0xDF630734B6B8522A, 0x3AD3ED03CD0AC26F,
|
||||
0xFADEAF2083C023D4, 0xC00D42234ECAE1BB, 0x8538CBA85CD76E96, 0xC402250E6E2458EB,
|
||||
0x47BC3413026A5D05, 0xAFD7A71F114272A4, 0x978DF784CC3F62E3, 0xB96DFC1EA144C781,
|
||||
0x21B2CF391596C8AE, 0x318E4E8D950916F3, 0xCE9556CC3E92E563, 0x385A509BDD7D1047,
|
||||
0x358129A0B5E7AFA3, 0xE6F387E363702B79, 0xE0755D5653E94001, 0x7BE903A5FFF9F412,
|
||||
0x12B53C2C90E80C75, 0x3307F315857EC4DB, 0x8FAFB86A0C61D31E, 0xD9E5DD8186213952,
|
||||
0x77F8AAD29FD622E2, 0x25BDA814357871FE, 0x7571174A8FA1F0CA, 0x137FEC60985D6561,
|
||||
0x30449EC19DBC7FE7, 0xA540D4DD41F4CF2C, 0xDC206AE0AE7AE916, 0x5B911CD0E2DA55A8,
|
||||
0xB2305F90F947131D, 0x344BF9ECBD52C6B7, 0x5D17C665D2433ED0, 0x18224FEEC05EB1FD,
|
||||
0x9E59E992844B6457, 0x9A568EBFA4A5DD07, 0xA3C60E68716DA454, 0x7E2CB4C4D7A22456,
|
||||
0x87B176304CA0BCBE, 0x413AEEA632F3367D, 0x9915E36BBC67663B, 0x40F03EEA3A465F69,
|
||||
0x1C2D28C3E0B008AD, 0x4E682A054A1E5BB1, 0x05C5B761285BD044, 0xE1BF8D1A5B5C2915,
|
||||
0xF2C0617AC3014C74, 0xB7F5E8F1D11CC359, 0x63CB4C4B3FA745EF, 0x9D1A84469C89DF6B,
|
||||
0xE33630824B2BFB3D, 0xD5F474F6E60EEFA2, 0xF58C6B83FB2D4E18, 0x4676E45F0ADF3411,
|
||||
0x20781F751D23A1BA, 0xBD629B3381AA7ED1, 0xAE1D775319F71BB0, 0xFED1C80DA32E9A84,
|
||||
0x5509083F92825170, 0x29AC01635557A70E, 0xA7C9694551831D04, 0x8E65682604D4BA0A,
|
||||
0x11F651F8882AB749, 0xD77DC96EF6793D8A, 0xEF2799F52B042DCD, 0x48EEF0B07A8730C9,
|
||||
0x22F1A2ED0D547392, 0x6142F1D32FD097C7, 0x4A674D286AF0E2E1, 0x80FD7CC9748CBED2,
|
||||
0x717E7067AF4F499A, 0x938290A9ECD1DBB3, 0x88E3B293344DD172, 0x2734158C250FA3D6
|
||||
};
|
||||
|
||||
|
||||
// KeySchedule
|
||||
__constant static sph_u64 CC[12][8] = {{
|
||||
0xe9daca1eda5b08b1, 0x1f7c65c0812fcbeb, 0x16d0452e43766a2f, 0xfcc485758db84e71,
|
||||
0x0169679291e07c4b, 0x15d360a4082a42a2, 0x234d74cc36747605, 0x0745a6f2596580dd
|
||||
}, {
|
||||
0x1a2f9da98ab5a36f, 0xd7b5700f469de34f, 0x982b230a72eafef3, 0x3101b5160f5ed561,
|
||||
0x5899d6126b17b59a, 0xcaa70adbc261b55c, 0x56cdcbd71ba2dd55, 0xb79bb121700479e6
|
||||
}, {
|
||||
0xc72fce2bacdc74f5, 0x35843d6a28fc390a, 0x8b1f9c525f5ef106, 0x7b7b29b11475eaf2,
|
||||
0xb19e3590e40fe2d3, 0x09db6260373ac9c1, 0x31db7a8643f4b6c2, 0xb20aba0af5961e99
|
||||
}, {
|
||||
0xd26615e8b3df1fef, 0xdde4715da0e148f9, 0x7d3c5c337e858e48, 0x3f355e68ad1c729d,
|
||||
0x75d603ed822cd7a9, 0xbe0352933313b7d8, 0xf137e893a1ea5334, 0x2ed1e384bcbe0c22
|
||||
}, {
|
||||
0x994747adac6bea4b, 0x6323a96c0c413f9a, 0x4a1086161f1c157f, 0xbdff0f80d7359e35,
|
||||
0xa3f53a254717cdbf, 0x161a2723b700ffdf, 0xf563eaa97ea2567a, 0x57fe6c7cfd581760
|
||||
}, {
|
||||
0xd9d33a1daeae4fae, 0xc039307a3bc3a46f, 0x6ca44251f9c4662d, 0xc68ef09ab49a7f18,
|
||||
0xb4b79a1cb7a6facf, 0xb6c6bec2661ff20a, 0x354f903672c571bf, 0x6e7d64467a4068fa
|
||||
}, {
|
||||
0xecc5aaee160ec7f4, 0x540924bffe86ac51, 0xc987bfe6c7c69e39, 0xc9937a19333e47d3,
|
||||
0x372c822dc5ab9209, 0x04054a2883694706, 0xf34a3ca24c451735, 0x93d4143a4d568688
|
||||
}, {
|
||||
0xa7c9934d425b1f9b, 0x41416e0c02aae703, 0x1ede369c71f8b74e, 0x9ac4db4d3b44b489,
|
||||
0x90069b92cb2b89f4, 0x2fc4a5d12b8dd169, 0xd9a8515935c2ac36, 0x1ee702bfd40d7fa4
|
||||
}, {
|
||||
0x9b223116545a8f37, 0xde5f16ecd89a4c94, 0x244289251b3a7d3a, 0x84090de0b755d93c,
|
||||
0xb1ceb2db0b440a80, 0x549c07a69a8a2b7b, 0x602a1fcb92dc380e, 0xdb5a238351446172
|
||||
}, {
|
||||
0x526f0580a6debeab, 0xf3f3e4b248e52a38, 0xdb788aff1ce74189, 0x0361331b8ae1ff1f,
|
||||
0x4b3369af0267e79f, 0xf452763b306c1e7a, 0xc3b63b15d1fa9836, 0xed9c4598fbc7b474
|
||||
}, {
|
||||
0xfb89c8efd09ecd7b, 0x94fe5a63cdc60230, 0x6107abebbb6bfad8, 0x7966841421800120,
|
||||
0xcab948eaef711d8a, 0x986e477d1dcdbaef, 0x5dd86fc04a59a2de, 0x1b2df381cda4ca6b
|
||||
}, {
|
||||
0xba3116f167e78e37, 0x7ab14904b08013d2, 0x771ddfbc323ca4cd, 0x9b9f2130d41220f8,
|
||||
0x86cc91189def805d, 0x5228e188aaa41de7, 0x991bb2d9d517f4fa, 0x20d71bf14a92bc48
|
||||
}};
|
||||
|
||||
__constant static sph_u64 F0[8] = // GOST_F(0)
|
||||
{
|
||||
0x74a5d4ce2efc83b3, 0x74a5d4ce2efc83b3, 0x74a5d4ce2efc83b3, 0x74a5d4ce2efc83b3,
|
||||
0x74a5d4ce2efc83b3, 0x74a5d4ce2efc83b3, 0x74a5d4ce2efc83b3, 0x74a5d4ce2efc83b3
|
||||
};
|
||||
|
||||
__constant static sph_u64 CC_F0[12][8] =
|
||||
{
|
||||
{ 0x8FD72F640708B0D0, 0x0DE874C7EBC3F213, 0xE92EEF3AD202E9E0, 0xC1E9DA0708013DA7, 0x9727DAB2F014BE88, 0x103051A02BCD6935, 0x33EC7E1DBD28F736, 0x1ECF460CF78AD1F4 },
|
||||
{ 0x0B2D9F89C775449D, 0x6B6EEFC6DAB7E8B0, 0xF1A0D31667F6EC44, 0x2A71132D5E108166, 0x0E9357C2EC87931A, 0xC99F5C1B4A01612D, 0x7E60B16E637D4EE2, 0xA9FCB827F9BA6D81 },
|
||||
{ 0x231FECA5AB3D285C, 0x70C6E1483C838C3B, 0x9C21C3C40CE4E2DA, 0x2FA796BD5688E573, 0x04C0E3FF55809FDF, 0x5FF978BFB8E3CDC8, 0xC54A19D6A3D07033, 0x0FCA83FDDE872478 },
|
||||
{ 0xBDF9312726339F10, 0x51A5BA1793BC9C56, 0xC4428DA14F96D2D4, 0xEC925222374EAB1F, 0x79477893747DD92F, 0xC495E19A46886304, 0x9C23F893BA7CFA36, 0x0C47268881FC5FEB },
|
||||
{ 0xCF117966029B2CB3, 0x07179ABE77088A8F, 0x671EF4CC2650E257, 0x7474B8B170DAB5C6, 0x4224FEBECF35113E, 0x993D156C675C5537, 0x2DEE3A5782C39B45, 0xE7C586F2990DD385 },
|
||||
{ 0x8608FD95B1C1138A, 0x8BB0847D9E9849AC, 0x5E76623F4F0EB0C7, 0x34C2BDBAFC5060CE, 0xE9E814475907826C, 0x22C9ED94D6AAC7C9, 0xE6B75E28171EB0D6, 0xF1329E5534E60215 },
|
||||
{ 0x86BB4814B1C3CE52, 0xE8F226C9FBDDD017, 0xCEDED67991CB3087, 0x76C33E32FDBFACA5, 0xDBB13BE1A9F7474C, 0x3D0273470342C356, 0x8E7246C51CF07F61, 0xAC8C125DDEF8DF71 },
|
||||
{ 0x6D73E747795B8CF3, 0x4E4AA65EA0072050, 0xA14A1582CB43C2B9, 0x748EF2B7BB63B938, 0x126789534410D7D4, 0xD4D48FF40301D791, 0xC67DFBE315C41FC0, 0x35E7A1A1AF88601C },
|
||||
{ 0x9BD33EA0FAB34007, 0xF51B7CDBE3D67D25, 0xD3ABDA0CE4186E6B, 0x8E61DDADCBCE1706, 0x58994565B41BE6A5, 0x7A87ABC1240CD31D, 0xFAFE6C28487968D0, 0x15B368609FF9EEA7 },
|
||||
{ 0xAE33263CCF115818, 0x93B2DBE9CADFCFC8, 0x0A91952BF91B0147, 0x458E67CA5F1ED73A, 0x94C2E5F288F074E3, 0x377895E85C69E996, 0xF11A4456AAB37B10, 0x163131934816821A },
|
||||
{ 0xD07E4A2366BF469D, 0x5EF1A3D220213B6C, 0x3C5BB78971D8ED0F, 0x0DE05E6B9006F2D2, 0xC58CFB00B8EAA1C9, 0xEFCDB54D1F250B76, 0xFD135634FA527042, 0x4CEE791290516407 },
|
||||
{ 0xD800B9264010790F, 0x974C4823E2B668D7, 0xA605A4B385C5E361, 0x3F6C92DA5A56D8D2, 0x82B9D67C12EF8277, 0x0AB6B4582561BF90, 0x46954FD98FC2CBA3, 0x70BE45CB21B6760D }
|
||||
};
|
||||
|
||||
|
||||
__constant static sph_u64 F1[8] = // GOST_F(1)
|
||||
{
|
||||
0x155f7bb040eec523, 0x155f7bb040eec523, 0x155f7bb040eec523, 0x155f7bb040eec523,
|
||||
0x155f7bb040eec523, 0x155f7bb040eec523, 0x155f7bb040eec523, 0x155f7bb040eec523
|
||||
};
|
||||
|
||||
__constant static sph_u64 CC_F1[12][8] =
|
||||
{
|
||||
{ 0xeaebb276318fee18, 0xea4c693382cbd63b, 0xbf26be88df699734, 0x49a504a9b6fa1c45, 0xb1666aa693de22da, 0x113563ea5e6b7e9c, 0xcdbf01848cd611e6, 0xb95e4a9dc30c7d0c },
|
||||
{ 0x919565a231cfa4aa, 0x46fde791cec8ae57, 0xe3c56411e2de27bf, 0x1f9d9e511aba0b94, 0x57773e25f11309ce, 0x2ce14b67cd005091, 0x00fb26ba738ef6c7, 0x2d5f800141af74fd },
|
||||
{ 0xf57a17cc650afe61, 0x26d3deadafe23502, 0xf87b7436229a32a5, 0x85459ccaae2842a5, 0x0d3a74dda91e80cd, 0x330e2b60f01ed098, 0x56c16add5dfb6720, 0x8692832019310082 },
|
||||
{ 0x6f63d34f5f688399, 0xa826bf5fb7abd51f, 0x3ecb2eaa144393e2, 0x4e7d6cc0863c69e4, 0x61e175af40d59b16, 0xba60d963cd6a540a, 0x69bf99c14c3995d5, 0x5a3de79f30d5a599 },
|
||||
{ 0x25f0e72cae7257f0, 0xfdb8c6bc7f9a6c15, 0x326e9413d635e7f1, 0xeaff2028e5942992, 0x1a55b07e905d6162, 0x882060860a9970d1, 0xe2b0cd223cc898af, 0x56a1f7c0137c29be },
|
||||
{ 0x4e6e5462c344d15a, 0xb7fb298868e7b346, 0x33741921c3e95374, 0xacb5e26b0e8d2b0b, 0x59f16751b3b69ec8, 0xa659593ea405b0b7, 0x98408efc8cb1a951, 0x8dbbcf819b3df0fc },
|
||||
{ 0x8d0aa21b9aec6c6a, 0x2b3534b940a84fb6, 0x2a1230d58e638c51, 0xc9daefb8e02f3383, 0xc709f5a9e5878201, 0x6f42d5dc6a746c8d, 0x3fb7df9057ada0b0, 0xaa6d0139a591f1c1 },
|
||||
{ 0xb3a97a7336702199, 0x51bd05f743668d8a, 0xc50f8f941f5351f3, 0xbdd89dee5fa35fe3, 0x9c4e220a589d4cbb, 0xed49fc69200e2ed8, 0x38354437945f7d36, 0x0904ddf5a8b68f2b },
|
||||
{ 0x1afa89fcc0636790, 0xda9d9eecd88892e6, 0xfec3d6bfe830769a, 0xafae622e5dc303d7, 0x7f7a31a7805db3f0, 0x916752f22230f876, 0x7b33cb8f67df8fca, 0xd205cb3c39e54fd7 },
|
||||
{ 0x648e61636c99ce88, 0x8533e43ee0c8a504, 0xbb9189e6eee32a4e, 0x6edbda389dc2f3bf, 0xdf6ddca6e9daa1d6, 0xd3962f27af34ce52, 0xe1e63f4c628c9c15, 0xd5ad89fc0b5c693d },
|
||||
{ 0x0646bda91e280a3e, 0x3a6f57000155ec3e, 0x579182cf68a16a50, 0x382fa3cafc78b976, 0x45ca8299c7305fb5, 0x778479d865838e62, 0x2a119981c6495ae7, 0xdbf255760f5a7b1d },
|
||||
{ 0xeb1ab39e4073b2f0, 0x22216718aefb32e4, 0xf9926a2b4248c862, 0x838bd14eb5ba6c3f, 0xa33f1ec5ff1cb214, 0xdb6aef763e43ff19, 0xa17f903ce0f5f90e, 0x03bf0065a0ecf9fc }
|
||||
};
|
||||
|
||||
|
||||
#define GOST_FS(_state, _return_state) \
|
||||
({ \
|
||||
unsigned char *state=(unsigned char *)_state; \
|
||||
sph_u64 *return_state=(sph_u64 *)_return_state; \
|
||||
sph_u64 r = 0; \
|
||||
r ^= T0S[state[56]]; \
|
||||
r ^= T1S[state[48]]; \
|
||||
r ^= T2S[state[40]]; \
|
||||
r ^= T3S[state[32]]; \
|
||||
r ^= T4S[state[24]]; \
|
||||
r ^= T5S[state[16]]; \
|
||||
r ^= T6S[state[8]]; \
|
||||
r ^= T7S[state[0]]; \
|
||||
*return_state = r; \
|
||||
r = 0; \
|
||||
\
|
||||
r ^= T0S[state[57]]; \
|
||||
r ^= T1S[state[49]]; \
|
||||
r ^= T2S[state[41]]; \
|
||||
r ^= T3S[state[33]]; \
|
||||
r ^= T4S[state[25]]; \
|
||||
r ^= T5S[state[17]]; \
|
||||
r ^= T6S[state[9]]; \
|
||||
r ^= T7S[state[1]]; \
|
||||
*(return_state+1) = r; \
|
||||
r = 0; \
|
||||
\
|
||||
r ^= T0S[state[58]]; \
|
||||
r ^= T1S[state[50]]; \
|
||||
r ^= T2S[state[42]]; \
|
||||
r ^= T3S[state[34]]; \
|
||||
r ^= T4S[state[26]]; \
|
||||
r ^= T5S[state[18]]; \
|
||||
r ^= T6S[state[10]]; \
|
||||
r ^= T7S[state[2]]; \
|
||||
*(return_state+2) = r; \
|
||||
r = 0; \
|
||||
\
|
||||
r ^= T0S[state[59]]; \
|
||||
r ^= T1S[state[51]]; \
|
||||
r ^= T2S[state[43]]; \
|
||||
r ^= T3S[state[35]]; \
|
||||
r ^= T4S[state[27]]; \
|
||||
r ^= T5S[state[19]]; \
|
||||
r ^= T6S[state[11]]; \
|
||||
r ^= T7S[state[3]]; \
|
||||
*(return_state+3) = r; \
|
||||
r = 0; \
|
||||
\
|
||||
r ^= T0S[state[60]]; \
|
||||
r ^= T1S[state[52]]; \
|
||||
r ^= T2S[state[44]]; \
|
||||
r ^= T3S[state[36]]; \
|
||||
r ^= T4S[state[28]]; \
|
||||
r ^= T5S[state[20]]; \
|
||||
r ^= T6S[state[12]]; \
|
||||
r ^= T7S[state[4]]; \
|
||||
*(return_state+4) = r; \
|
||||
r = 0; \
|
||||
\
|
||||
r ^= T0S[state[61]]; \
|
||||
r ^= T1S[state[53]]; \
|
||||
r ^= T2S[state[45]]; \
|
||||
r ^= T3S[state[37]]; \
|
||||
r ^= T4S[state[29]]; \
|
||||
r ^= T5S[state[21]]; \
|
||||
r ^= T6S[state[13]]; \
|
||||
r ^= T7S[state[5]]; \
|
||||
*(return_state+5) = r; \
|
||||
r = 0; \
|
||||
\
|
||||
r ^= T0S[state[62]]; \
|
||||
r ^= T1S[state[54]]; \
|
||||
r ^= T2S[state[46]]; \
|
||||
r ^= T3S[state[38]]; \
|
||||
r ^= T4S[state[30]]; \
|
||||
r ^= T5S[state[22]]; \
|
||||
r ^= T6S[state[14]]; \
|
||||
r ^= T7S[state[6]]; \
|
||||
*(return_state+6) = r; \
|
||||
r = 0; \
|
||||
\
|
||||
r ^= T0S[state[63]]; \
|
||||
r ^= T1S[state[55]]; \
|
||||
r ^= T2S[state[47]]; \
|
||||
r ^= T3S[state[39]]; \
|
||||
r ^= T4S[state[31]]; \
|
||||
r ^= T5S[state[23]]; \
|
||||
r ^= T6S[state[15]]; \
|
||||
r ^= T7S[state[7]]; \
|
||||
*(return_state+7) = r; \
|
||||
})
|
||||
|
||||
inline void GOST_Copy512(sph_u64* dst, const sph_u64* src)
|
||||
{
|
||||
_Pragma("unroll") for (int i=0; i<8; i++)
|
||||
dst[i] = src[i];
|
||||
}
|
||||
|
||||
|
||||
#define GOST_F(state) \
|
||||
({ \
|
||||
sph_u64 t[8]; \
|
||||
GOST_FS(state, t); \
|
||||
GOST_Copy512(state, t); \
|
||||
})
|
||||
|
||||
inline void GOST_Xor512(sph_u64* C, const sph_u64* A, const sph_u64* B)
|
||||
{
|
||||
C[0] = A[0] ^ B[0];
|
||||
C[1] = A[1] ^ B[1];
|
||||
C[2] = A[2] ^ B[2];
|
||||
C[3] = A[3] ^ B[3];
|
||||
C[4] = A[4] ^ B[4];
|
||||
C[5] = A[5] ^ B[5];
|
||||
C[6] = A[6] ^ B[6];
|
||||
C[7] = A[7] ^ B[7];
|
||||
}
|
||||
|
||||
inline void GOST_Xor512_constant(sph_u64* C, const sph_u64* A, __constant const sph_u64* B)
|
||||
{
|
||||
C[0] = A[0] ^ B[0];
|
||||
C[1] = A[1] ^ B[1];
|
||||
C[2] = A[2] ^ B[2];
|
||||
C[3] = A[3] ^ B[3];
|
||||
C[4] = A[4] ^ B[4];
|
||||
C[5] = A[5] ^ B[5];
|
||||
C[6] = A[6] ^ B[6];
|
||||
C[7] = A[7] ^ B[7];
|
||||
}
|
||||
|
||||
inline void GOST_Xor512_3(sph_u64* C, const sph_u64* A, const sph_u64* B)
|
||||
{
|
||||
C[0] ^= A[0] ^ B[0];
|
||||
C[1] ^= A[1] ^ B[1];
|
||||
C[2] ^= A[2] ^ B[2];
|
||||
C[3] ^= A[3] ^ B[3];
|
||||
C[4] ^= A[4] ^ B[4];
|
||||
C[5] ^= A[5] ^ B[5];
|
||||
C[6] ^= A[6] ^ B[6];
|
||||
C[7] ^= A[7] ^ B[7];
|
||||
}
|
||||
|
||||
inline void GOST_Xor512_c(sph_u64* C, const sph_u64* A, const sph_u64* B, sph_u64 c)
|
||||
{
|
||||
C[0] = A[0] ^ B[0] ^ c;
|
||||
C[1] = A[1] ^ B[1] ^ c;
|
||||
C[2] = A[2] ^ B[2] ^ c;
|
||||
C[3] = A[3] ^ B[3] ^ c;
|
||||
C[4] = A[4] ^ B[4] ^ c;
|
||||
C[5] = A[5] ^ B[5] ^ c;
|
||||
C[6] = A[6] ^ B[6] ^ c;
|
||||
C[7] = A[7] ^ B[7] ^ c;
|
||||
}
|
||||
|
||||
#define GOST_E12(K, state) \
|
||||
({ \
|
||||
sph_u64 state1[8], K1[8]; \
|
||||
GOST_Xor512_constant(state1, K, CC[0]); \
|
||||
GOST_FS(state1, K1); \
|
||||
GOST_FS(state, state1); \
|
||||
GOST_Xor512(state, state1, K1); \
|
||||
_Pragma("unroll") for(int i=1; i<12; i++) \
|
||||
{ \
|
||||
GOST_Xor512_constant(state1, K1, CC[i]); \
|
||||
GOST_FS(state1, K1); \
|
||||
GOST_FS(state, state1); \
|
||||
GOST_Xor512(state, state1, K1); \
|
||||
} \
|
||||
})
|
||||
|
||||
#define GOST_E(K, m, state /* out only */) \
|
||||
({ \
|
||||
GOST_Xor512(state, m, K); /* state = m ^ K */ \
|
||||
GOST_E12(K, state); \
|
||||
})
|
||||
|
||||
#define GOST_E_F0(m, state /* out only */) \
|
||||
({ \
|
||||
GOST_Xor512_constant(state, m, F0); /* state = m ^ F0 */ \
|
||||
sph_u64 state1[8]; \
|
||||
_Pragma("unroll") for(int i=0; i<12; i++) \
|
||||
{ \
|
||||
GOST_FS(state, state1); \
|
||||
GOST_Xor512_constant(state, state1, CC_F0[i]); \
|
||||
} \
|
||||
})
|
||||
|
||||
#define GOST_E_F1(m, state /* out only */) \
|
||||
({ \
|
||||
GOST_Xor512_constant(state, m, F1); /* state = m ^ F0 */ \
|
||||
sph_u64 state1[8]; \
|
||||
_Pragma("unroll") for(int i=0; i<12; i++) \
|
||||
{ \
|
||||
GOST_FS(state, state1); \
|
||||
GOST_Xor512_constant(state, state1, CC_F1[i]); \
|
||||
} \
|
||||
})
|
||||
|
||||
#define GOST_g_N(h, M, N) \
|
||||
({ \
|
||||
sph_u64 K[8]; \
|
||||
GOST_Xor512(K, N, h); /* K = N ^ h */ \
|
||||
GOST_F(K); \
|
||||
sph_u64 t[8]; \
|
||||
GOST_E(K, M, t); \
|
||||
GOST_Xor512_3(h, t, M); /* h = h ^ t ^ M */ \
|
||||
})
|
||||
|
||||
#define GOST_g_0(h, M) \
|
||||
({ \
|
||||
sph_u64 K[8]; \
|
||||
GOST_FS(h, K); \
|
||||
sph_u64 t[8]; \
|
||||
GOST_E(K, M, t); \
|
||||
GOST_Xor512_3(h, t, M); /* h = h ^ t ^ M */ \
|
||||
})
|
||||
|
||||
#define GOST_g_0_0(h, M) /* input h assumed zero, for iv 512 */ \
|
||||
({ \
|
||||
GOST_E_F0 (M, h); \
|
||||
GOST_Xor512 (h, h, M); /* h = h ^ M */ \
|
||||
})
|
||||
|
||||
#define GOST_g_0_1(h, M) /* input h assumed all bytes one, for iv 256 */ \
|
||||
({ \
|
||||
GOST_E_F1 (M, h); \
|
||||
GOST_Xor512_c(h, h, M, 0x0101010101010101); /* h = h ^ M ^ 1 */ \
|
||||
})
|
||||
|
||||
|
||||
inline void GOST_Add128(void *x, void * const a, void * const b)
|
||||
{
|
||||
unsigned short t = 0;
|
||||
_Pragma("unroll") for(int i = 15; i >= 0; i--)
|
||||
{
|
||||
t = ((unsigned char *)a)[i] + ((unsigned char *)b)[i] + (t >> 8);
|
||||
((unsigned char *)x)[i] = t & 0xFF;
|
||||
}
|
||||
}
|
||||
|
||||
typedef struct {
|
||||
uint data[20];
|
||||
uint target[8];
|
||||
} dev_blk_ctx;
|
||||
|
||||
__kernel __attribute__((vec_type_hint(uint))) WGS void gostminer(
|
||||
__constant dev_blk_ctx *ctx, __global uint *output)
|
||||
{
|
||||
// move to shared memory
|
||||
__local sph_u64 T0S[256];
|
||||
__local sph_u64 T1S[256];
|
||||
__local sph_u64 T2S[256];
|
||||
__local sph_u64 T3S[256];
|
||||
__local sph_u64 T4S[256];
|
||||
__local sph_u64 T5S[256];
|
||||
__local sph_u64 T6S[256];
|
||||
__local sph_u64 T7S[256];
|
||||
|
||||
int x = get_local_id(0);
|
||||
T0S[x] = T0[x]; T1S[x] = T1[x]; T2S[x] = T2[x]; T3S[x] = T3[x];
|
||||
T4S[x] = T4[x]; T5S[x] = T5[x]; T6S[x] = T6[x]; T7S[x] = T7[x];
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
const uint myid = get_global_id(0);
|
||||
const uint nonce = ctx->data[19] + myid;
|
||||
sph_u64 N[8] = {0};
|
||||
// first hash (GOST 34.11-512 over 80 bytes)
|
||||
sph_u64 block2[8];
|
||||
for (int i = 0; i < 15; i++) ((uint *)block2)[i] = ctx->data[i+4];
|
||||
((uint *)block2)[15] = nonce; // change nonce
|
||||
sph_u64 hash1[8];
|
||||
// second block
|
||||
GOST_g_0_0 (hash1, block2); // zero iv for 512 assumed
|
||||
N[7] = 0x0002000000000000; // 512
|
||||
// first block
|
||||
sph_u64 block1[10] = { 0, 0, 0, 0, 0, 0, 0, 0x0100000000000000, 0, 0 };
|
||||
block1[8] = ((__constant sph_u64*)ctx->data)[0];
|
||||
block1[9] = ((__constant sph_u64*)ctx->data)[1];
|
||||
GOST_g_N(hash1, block1 + 2, N);
|
||||
N[7] |= 0x8000000000000000; // +128
|
||||
GOST_g_0(hash1, N);
|
||||
GOST_Add128(block2 + 6, block2 + 6, block1 + 8);
|
||||
block2[5] += 0x0100000000000000;
|
||||
GOST_g_0(hash1, block2);
|
||||
// second hash (GOST 34.11-256 over 64 bytes)
|
||||
sph_u64 hash[8];
|
||||
// second block
|
||||
GOST_g_0_1(hash, hash1); // // iv for 256 assumed (all bytes one)
|
||||
N[7] = 0x0002000000000000; // 512
|
||||
// first block
|
||||
GOST_g_N(hash, block1, N);
|
||||
GOST_g_0(hash, N);
|
||||
hash1[7] += 0x0100000000000000;
|
||||
GOST_g_0(hash, hash1);
|
||||
// result is first 32 bytes of hash
|
||||
output[myid] = (SWAP4 (LODWORD (hash[0])) <= ctx->target[7] && SWAP4 (HIDWORD (hash[0])) <= ctx->target[6]) ? nonce : 0;
|
||||
}
|
||||
115
miner.c
115
miner.c
@ -16,8 +16,14 @@
|
||||
#include <stdint.h>
|
||||
#include <unistd.h>
|
||||
#include <sys/time.h>
|
||||
|
||||
#ifdef _WIN32
|
||||
#include <winsock2.h>
|
||||
#else
|
||||
#include <sys/resource.h>
|
||||
#include <sys/select.h>
|
||||
#endif
|
||||
|
||||
#include <pthread.h>
|
||||
#include <getopt.h>
|
||||
#include <jansson.h>
|
||||
@ -25,13 +31,13 @@
|
||||
#include <math.h>
|
||||
#include "miner.h"
|
||||
|
||||
#include "findnonce.h"
|
||||
#include "ocl.h"
|
||||
#include "streebog.h"
|
||||
|
||||
#define VERSION "0.1"
|
||||
|
||||
#define PROGRAM_NAME "oclminer"
|
||||
#define DEF_RPC_URL "http://127.0.0.1:8332/"
|
||||
#define PROGRAM_NAME "gostoclminer"
|
||||
#define DEF_RPC_URL "http://127.0.0.1:9376/"
|
||||
#define DEF_RPC_USERPASS "rpcuser:rpcpass"
|
||||
|
||||
enum {
|
||||
@ -126,12 +132,6 @@ static bool jobj_binary(const json_t *obj, const char *key,
|
||||
|
||||
static bool work_decode(const json_t *val, struct work_t *work)
|
||||
{
|
||||
if (!jobj_binary(val, "midstate",
|
||||
work->midstate, sizeof(work->midstate))) {
|
||||
fprintf(stderr, "JSON inval midstate\n");
|
||||
goto err_out;
|
||||
}
|
||||
|
||||
if (!jobj_binary(val, "data", work->data, sizeof(work->data))) {
|
||||
fprintf(stderr, "JSON inval data\n");
|
||||
goto err_out;
|
||||
@ -185,8 +185,9 @@ static void submit_work(struct work_t *work)
|
||||
|
||||
res = json_object_get(val, "result");
|
||||
|
||||
printf("PROOF OF WORK RESULT: %s\n",
|
||||
json_is_true(res) ? "true (yay!!!)" : "false (booooo)");
|
||||
printf("PROOF OF WORK RESULT: %s\n", json_is_true(res) ? "true (yay!!!)" : "false (booooo)");
|
||||
if (!json_is_true(res))
|
||||
printf ("REASON: %s\n", json_string_value(json_object_get(val, "reject-reason")));
|
||||
|
||||
json_decref(val);
|
||||
|
||||
@ -320,9 +321,15 @@ static void *miner_thread(void *thr_id_int)
|
||||
continue;
|
||||
}
|
||||
|
||||
precalc_hash(&work[frame].blk, (uint32_t *)(work[frame].midstate), (uint32_t *)(work[frame].data + 64));
|
||||
memcpy (work[frame].blk.data, work[frame].data, 80);
|
||||
int k;
|
||||
for (k = 0; k < 19; k++) work[frame].blk.data[k] = swap32 (work[frame].blk.data[k]);
|
||||
memcpy (work[frame].blk.target, work[frame].target, 32);
|
||||
|
||||
work[frame].blk.nonce = 0;
|
||||
/*work[frame].blk.target[6] = 0xFFFFFFFF;
|
||||
work[frame].blk.target[7] = 0x000000FF;*/
|
||||
|
||||
work[frame].blk.data[19] = 0;
|
||||
work[frame].valid = true;
|
||||
work[frame].ready = 0;
|
||||
|
||||
@ -332,9 +339,9 @@ static void *miner_thread(void *thr_id_int)
|
||||
|
||||
gettimeofday(&tv_start, NULL);
|
||||
|
||||
int threads = 102400 * 4;
|
||||
int threads = 65536; // TODO:
|
||||
globalThreads[0] = threads;
|
||||
localThreads[0] = 128;
|
||||
localThreads[0] = 256;
|
||||
|
||||
status = clEnqueueWriteBuffer(clState->commandQueue, clState->inputBuffer, CL_TRUE, 0,
|
||||
sizeof(dev_blk_ctx), (void *)&work[frame].blk, 0, NULL, NULL);
|
||||
@ -348,56 +355,46 @@ static void *miner_thread(void *thr_id_int)
|
||||
|
||||
clFlush(clState->commandQueue);
|
||||
|
||||
hashes_done = 1024 * threads;
|
||||
hashes_done = 1024 * threads;
|
||||
|
||||
if (work[res_frame].ready) {
|
||||
rc = false;
|
||||
|
||||
uint32_t bestG = ~0;
|
||||
uint32_t nonce;
|
||||
if (work[res_frame].ready)
|
||||
{
|
||||
int j;
|
||||
for(j = 0; j < work[res_frame].ready; j++) {
|
||||
if(res[j]) {
|
||||
uint32_t start = (work[res_frame].res_nonce + j)<<10;
|
||||
uint32_t my_g, my_nonce;
|
||||
my_g = postcalc_hash(&work[res_frame].blk, &work[res_frame], start, start + 1026, &my_nonce, opt_pool, &h0count);
|
||||
|
||||
if (!opt_pool) {
|
||||
if (opt_debug)
|
||||
fprintf(stderr, "DEBUG: H0 within %u .. %u, best G = %08x, nonce = %08x\n", start, start + 1024, my_g, my_nonce);
|
||||
|
||||
if(my_g < bestG) {
|
||||
bestG = my_g;
|
||||
nonce = my_nonce;
|
||||
if (opt_debug)
|
||||
fprintf(stderr, "new best\n");
|
||||
}
|
||||
for(j = 0; j < work[res_frame].ready; j++)
|
||||
{
|
||||
if(res[j])
|
||||
{
|
||||
uint32_t hash[8];
|
||||
work[frame].blk.data[19] = res[j];
|
||||
gostd_hash (hash, work[frame].blk.data);
|
||||
work[frame].blk.data[19] = 0;
|
||||
int k;
|
||||
for (k = 0; k < 8; k++) printf ("%08x ", hash[k]);
|
||||
printf ("\n");
|
||||
if (swap32 (hash[0]) <= work[frame].blk.target[7])
|
||||
{
|
||||
uint32_t *target1 = (uint32_t *)(work[res_frame].target + 24);
|
||||
uint32_t *target2 = (uint32_t *)(work[res_frame].target + 28);
|
||||
printf("Found solution for %08x %08x: %08x\n", *target1, *target2, res[j]);
|
||||
submit_nonce(&work[res_frame], swap32 (res[j]));
|
||||
block++;
|
||||
h0count++;
|
||||
need_work = true;
|
||||
break;
|
||||
}
|
||||
|
||||
rc = true;
|
||||
else
|
||||
printf ("result for %08x does not validate on CPU!", res[j]);
|
||||
}
|
||||
}
|
||||
|
||||
work[res_frame].ready = false;
|
||||
|
||||
uint32_t *target = (uint32_t *)(work[res_frame].target + 24);
|
||||
|
||||
if(!opt_pool && rc && bestG <= *target) {
|
||||
printf("Found solution for %08x: %08x %u\n", *target, bestG, nonce);
|
||||
|
||||
submit_nonce(&work[res_frame], nonce);
|
||||
|
||||
block++;
|
||||
|
||||
need_work = true;
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
if (h0count != 0) {
|
||||
if (h0count != 0)
|
||||
{
|
||||
double secs;
|
||||
secs = time2secs(&tv_start0);
|
||||
hashrates[thr_id] = h0count * pow(2, 256) / (secs * 1e6 * (pow(2, 224) - 1.0));
|
||||
hashrates[thr_id] = h0count * pow (2, 256) / (secs * 1e6 * (pow (2, 224) - 1.0));
|
||||
}
|
||||
|
||||
status = clEnqueueReadBuffer(clState->commandQueue, clState->outputBuffer, CL_TRUE, 0,
|
||||
@ -406,11 +403,11 @@ static void *miner_thread(void *thr_id_int)
|
||||
|
||||
res_frame = frame;
|
||||
work[res_frame].ready = threads;
|
||||
work[res_frame].res_nonce = work[res_frame].blk.nonce;
|
||||
work[res_frame].res_nonce = work[res_frame].blk.data[19];
|
||||
|
||||
work[frame].blk.nonce += threads;
|
||||
work[frame].blk.data[19] += threads;
|
||||
|
||||
if (work[frame].blk.nonce > 4000000 - threads)
|
||||
if (work[frame].blk.data[19] > 4000000 - threads)
|
||||
need_work = true;
|
||||
|
||||
failures = 0;
|
||||
@ -423,8 +420,8 @@ static void show_usage(void)
|
||||
{
|
||||
int i;
|
||||
|
||||
printf("oclminer version %s\n\n", VERSION);
|
||||
printf("Usage:\tminerd [options]\n\nSupported options:\n");
|
||||
printf("gostoclminer version %s\n\n", VERSION);
|
||||
printf("Usage:\tgostoclminer [options]\n\nSupported options:\n");
|
||||
for (i = 0; i < ARRAY_SIZE(options_help); i++) {
|
||||
struct option_help *h;
|
||||
|
||||
|
||||
2
miner.h
2
miner.h
@ -1,5 +1,6 @@
|
||||
#ifndef __MINER_H__
|
||||
#define __MINER_H__
|
||||
#include <inttypes.h>
|
||||
#include "ocl.h"
|
||||
|
||||
#ifndef ARRAY_SIZE
|
||||
@ -14,5 +15,6 @@ extern char *bin2hex(unsigned char *p, size_t len);
|
||||
extern int hex2bin(unsigned char *p, const char *hexstr, size_t len);
|
||||
|
||||
extern int timeval_subtract (struct timeval *result, struct timeval *x, struct timeval *y);
|
||||
extern uint32_t swap32(uint32_t value);
|
||||
|
||||
#endif /* __MINER_H__ */
|
||||
|
||||
21
ocl.c
21
ocl.c
@ -3,14 +3,19 @@
|
||||
#include <string.h>
|
||||
#include <stdio.h>
|
||||
#include <sys/types.h>
|
||||
|
||||
#ifdef _WIN32
|
||||
#include <winsock2.h>
|
||||
#else
|
||||
#include <sys/socket.h>
|
||||
#include <netinet/in.h>
|
||||
#include <netdb.h>
|
||||
#include <netdb.h>
|
||||
#endif
|
||||
|
||||
#include <time.h>
|
||||
#include <sys/time.h>
|
||||
#include <pthread.h>
|
||||
|
||||
#include "findnonce.h"
|
||||
#include "ocl.h"
|
||||
|
||||
char *file_contents(const char *filename, int *length)
|
||||
@ -210,7 +215,7 @@ _clState *initCl(int gpu, char *name, size_t nameSize) {
|
||||
// Load CL file, build CL program object, create CL kernel object
|
||||
/////////////////////////////////////////////////////////////////
|
||||
//
|
||||
const char * filename = "oclminer.cl";
|
||||
const char * filename = "gostminer.cl";
|
||||
int pl;
|
||||
char *source = file_contents(filename, &pl);
|
||||
size_t sourceSize[] = {(size_t)pl};
|
||||
@ -223,7 +228,7 @@ _clState *initCl(int gpu, char *name, size_t nameSize) {
|
||||
}
|
||||
|
||||
/* create a cl program executable for all the devices specified */
|
||||
status = clBuildProgram(clState->program, 1, &devices[gpu], NULL, NULL, NULL);
|
||||
status = clBuildProgram(clState->program, 1, &devices[gpu], "", NULL, NULL);
|
||||
if(status != CL_SUCCESS)
|
||||
{
|
||||
printf("Error: Building Program (clBuildProgram)\n");
|
||||
@ -237,7 +242,7 @@ _clState *initCl(int gpu, char *name, size_t nameSize) {
|
||||
}
|
||||
|
||||
/* get a kernel object handle for a kernel with the given name */
|
||||
clState->kernel = clCreateKernel(clState->program, "oclminer", &status);
|
||||
clState->kernel = clCreateKernel(clState->program, "gostminer", &status);
|
||||
if(status != CL_SUCCESS)
|
||||
{
|
||||
printf("Error: Creating Kernel from program. (clCreateKernel)\n");
|
||||
@ -254,16 +259,16 @@ _clState *initCl(int gpu, char *name, size_t nameSize) {
|
||||
return NULL;
|
||||
}
|
||||
|
||||
clState->inputBuffer = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, sizeof(dev_blk_ctx), NULL, &status);
|
||||
clState->inputBuffer = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, sizeof (dev_blk_ctx), NULL, &status);
|
||||
if(status != CL_SUCCESS) {
|
||||
printf("Error: clCreateBuffer (inputBuffer)\n");
|
||||
return;
|
||||
return NULL;
|
||||
}
|
||||
|
||||
clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, sizeof(uint32_t) * MAXTHREADS, NULL, &status);
|
||||
if(status != CL_SUCCESS) {
|
||||
printf("Error: clCreateBuffer (outputBuffer)\n");
|
||||
return;
|
||||
return NULL;
|
||||
}
|
||||
|
||||
return clState;
|
||||
|
||||
22
ocl.h
22
ocl.h
@ -6,6 +6,8 @@
|
||||
#include <CL/cl.h>
|
||||
#endif
|
||||
|
||||
#define MAXTHREADS 256000
|
||||
|
||||
typedef struct {
|
||||
cl_context context;
|
||||
cl_kernel kernel;
|
||||
@ -15,6 +17,26 @@ typedef struct {
|
||||
cl_mem outputBuffer;
|
||||
} _clState;
|
||||
|
||||
typedef struct {
|
||||
uint32_t data[20];
|
||||
uint32_t target[8];
|
||||
} dev_blk_ctx;
|
||||
|
||||
|
||||
struct work_t {
|
||||
unsigned char data[128];
|
||||
unsigned char hash1[64];
|
||||
unsigned char midstate[32];
|
||||
unsigned char target[32];
|
||||
|
||||
unsigned char hash[32];
|
||||
uint32_t output[MAXTHREADS];
|
||||
uint32_t res_nonce;
|
||||
uint32_t valid;
|
||||
uint32_t ready;
|
||||
dev_blk_ctx blk;
|
||||
};
|
||||
|
||||
extern char *file_contents(const char *filename, int *length);
|
||||
extern int clDevicesNum();
|
||||
extern _clState *initCl(int gpu, char *name, size_t nameSize);
|
||||
|
||||
1976
sph_types.h
Normal file
1976
sph_types.h
Normal file
File diff suppressed because it is too large
Load Diff
1054
streebog.c
Normal file
1054
streebog.c
Normal file
File diff suppressed because it is too large
Load Diff
187
streebog.h
Normal file
187
streebog.h
Normal file
@ -0,0 +1,187 @@
|
||||
/* $Id: sph_gost.h 216 2010-06-08 09:46:57Z tp $ */
|
||||
/**
|
||||
* GOST interface. This is the interface for GOST R 12 with the
|
||||
* recommended parameters for SHA-3, with output lengths 256
|
||||
* and 512 bits.
|
||||
*
|
||||
* ==========================(LICENSE BEGIN)============================
|
||||
*
|
||||
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
|
||||
*
|
||||
* Permission is hereby granted, free of charge, to any person obtaining
|
||||
* a copy of this software and associated documentation files (the
|
||||
* "Software"), to deal in the Software without restriction, including
|
||||
* without limitation the rights to use, copy, modify, merge, publish,
|
||||
* distribute, sublicense, and/or sell copies of the Software, and to
|
||||
* permit persons to whom the Software is furnished to do so, subject to
|
||||
* the following conditions:
|
||||
*
|
||||
* The above copyright notice and this permission notice shall be
|
||||
* included in all copies or substantial portions of the Software.
|
||||
*
|
||||
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
|
||||
* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
|
||||
* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
|
||||
* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
|
||||
* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
|
||||
* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
|
||||
* SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
|
||||
*
|
||||
* ===========================(LICENSE END)=============================
|
||||
*
|
||||
* @file sph_gost.h
|
||||
* @author Mish <mish@btchouse.com>
|
||||
*/
|
||||
|
||||
#ifndef SPH_GOST_H__
|
||||
#define SPH_GOST_H__
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C"{
|
||||
#endif
|
||||
|
||||
#include <stddef.h>
|
||||
#include "sph_types.h"
|
||||
|
||||
/**
|
||||
* Output size (in bits) for GOST-256.
|
||||
*/
|
||||
#define SPH_SIZE_gost256 256
|
||||
|
||||
/**
|
||||
* Output size (in bits) for GOST-512.
|
||||
*/
|
||||
#define SPH_SIZE_gost512 512
|
||||
|
||||
/**
|
||||
* This structure is a context for Keccak computations: it contains the
|
||||
* intermediate values and some data from the last entered block. Once a
|
||||
* GOST computation has been performed, the context can be reused for
|
||||
* another computation.
|
||||
*
|
||||
* The contents of this structure are private. A running GOST computation
|
||||
* can be cloned by copying the context (e.g. with a simple
|
||||
* <code>memcpy()</code>).
|
||||
*/
|
||||
|
||||
/**
|
||||
* This structure is a context for Gost-256 computations.
|
||||
*/
|
||||
|
||||
typedef struct {
|
||||
#ifndef DOXYGEN_IGNORE
|
||||
unsigned char buf[32]; /* first field, for alignment */
|
||||
size_t ptr;
|
||||
sph_u32 V[3][8];
|
||||
#endif
|
||||
} sph_gost256_context;
|
||||
|
||||
/**
|
||||
* This structure is a context for Gost-512 computations.
|
||||
*/
|
||||
typedef struct {
|
||||
#ifndef DOXYGEN_IGNORE
|
||||
unsigned char buf[64]; /* first field, for alignment */
|
||||
size_t ptr;
|
||||
sph_u32 V[5][8];
|
||||
#endif
|
||||
} sph_gost512_context;
|
||||
|
||||
|
||||
/**
|
||||
* Initialize a GOST-256 context. This process performs no memory allocation.
|
||||
*
|
||||
* @param cc the GOST-256 context (pointer to a
|
||||
* <code>sph_gost256_context</code>)
|
||||
*/
|
||||
void sph_gost256_init(void *cc);
|
||||
|
||||
/**
|
||||
* Process some data bytes. It is acceptable that <code>len</code> is zero
|
||||
* (in which case this function does nothing).
|
||||
*
|
||||
* @param cc the Gost-256 context
|
||||
* @param data the input data
|
||||
* @param len the input data length (in bytes)
|
||||
*/
|
||||
void sph_gost256(void *cc, const void *data, size_t len);
|
||||
|
||||
/**
|
||||
* Terminate the current GOST-256 computation and output the result into
|
||||
* the provided buffer. The destination buffer must be wide enough to
|
||||
* accomodate the result (32 bytes). The context is automatically
|
||||
* reinitialized.
|
||||
*
|
||||
* @param cc the GOST-256 context
|
||||
* @param dst the destination buffer
|
||||
*/
|
||||
void sph_gost256_close(void *cc, void *dst);
|
||||
|
||||
/**
|
||||
* Add a few additional bits (0 to 7) to the current computation, then
|
||||
* terminate it and output the result in the provided buffer, which must
|
||||
* be wide enough to accomodate the result (32 bytes). If bit number i
|
||||
* in <code>ub</code> has value 2^i, then the extra bits are those
|
||||
* numbered 7 downto 8-n (this is the big-endian convention at the byte
|
||||
* level). The context is automatically reinitialized.
|
||||
*
|
||||
* @param cc the GOST-256 context
|
||||
* @param ub the extra bits
|
||||
* @param n the number of extra bits (0 to 7)
|
||||
* @param dst the destination buffer
|
||||
*/
|
||||
void sph_gost256_addbits_and_close(
|
||||
void *cc, unsigned ub, unsigned n, void *dst);
|
||||
|
||||
/**
|
||||
* Initialize a Gost-512 context. This process performs no memory allocation.
|
||||
*
|
||||
* @param cc the GOST-512 context (pointer to a
|
||||
* <code>sph_gost512_context</code>)
|
||||
*/
|
||||
void sph_gost512_init(void *cc);
|
||||
|
||||
/**
|
||||
* Process some data bytes. It is acceptable that <code>len</code> is zero
|
||||
* (in which case this function does nothing).
|
||||
*
|
||||
* @param cc the GOST-512 context
|
||||
* @param data the input data
|
||||
* @param len the input data length (in bytes)
|
||||
*/
|
||||
void sph_gost512(void *cc, const void *data, size_t len);
|
||||
|
||||
/**
|
||||
* Terminate the current GOST-512 computation and output the result into
|
||||
* the provided buffer. The destination buffer must be wide enough to
|
||||
* accomodate the result (64 bytes). The context is automatically
|
||||
* reinitialized.
|
||||
*
|
||||
* @param cc the GOST-512 context
|
||||
* @param dst the destination buffer
|
||||
*/
|
||||
void sph_gost512_close(void *cc, void *dst);
|
||||
|
||||
/**
|
||||
* Add a few additional bits (0 to 7) to the current computation, then
|
||||
* terminate it and output the result in the provided buffer, which must
|
||||
* be wide enough to accomodate the result (64 bytes). If bit number i
|
||||
* in <code>ub</code> has value 2^i, then the extra bits are those
|
||||
* numbered 7 downto 8-n (this is the big-endian convention at the byte
|
||||
* level). The context is automatically reinitialized.
|
||||
*
|
||||
* @param cc the GOST-512 context
|
||||
* @param ub the extra bits
|
||||
* @param n the number of extra bits (0 to 7)
|
||||
* @param dst the destination buffer
|
||||
*/
|
||||
void sph_gost512_addbits_and_close(
|
||||
void *cc, unsigned ub, unsigned n, void *dst);
|
||||
|
||||
void gostd_hash(void *output, const void *input);
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
#endif
|
||||
19
util.c
19
util.c
@ -129,7 +129,7 @@ json_t *json_rpc_call(const char *url, const char *userpass, const char *rpc_req
|
||||
if (rc)
|
||||
goto err_out;
|
||||
|
||||
val = json_loads(all_data.buf, &err);
|
||||
val = json_loads(all_data.buf, 0, &err);
|
||||
if (!val) {
|
||||
fprintf(stderr, "JSON failed(%d): %s\n", err.line, err.text);
|
||||
goto err_out;
|
||||
@ -226,3 +226,20 @@ timeval_subtract (
|
||||
return x->tv_sec < y->tv_sec;
|
||||
}
|
||||
|
||||
uint32_t swap32(uint32_t value)
|
||||
{
|
||||
#if defined(__x86_64__) || defined(__i386__)
|
||||
__asm__
|
||||
(
|
||||
"bswap %0"
|
||||
: "=r"(value)
|
||||
: "0"(value)
|
||||
:
|
||||
);
|
||||
return value;
|
||||
#else
|
||||
value = ((value & 0xFF00FF00) >> 8) | ((value & 0x00FF00FF) << 8);
|
||||
return (value<<16) | (value>>16);
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user