Bitcoin Forum
August 16, 2017, 07:42:23 PM *
News: ALL CLEAR: You can now use Bitcoin as you were previously. For more info, including how to claim your BCH (optional), see here.
 
   Home   Help Search Donate Login Register  
Pages: « 1 2 3 4 5 6 7 8 9 10 11 12 13 [14] 15 16 17 18 19 20 21 »  All
  Print  
Author Topic: [ANN][GRS][DMD][DGB] Pallas optimized groestl opencl kernels  (Read 56081 times)
This is a self-moderated topic. If you do not want to be moderated by the person who started this topic, create a new topic.
pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
February 25, 2015, 02:43:16 PM
 #261

Thx pallas! But im out of the game cos i have 280s only. Utahjohn too (i think so).

the new kernel should make no difference on tahiti cards, but we will eventually make some tests later anyway: the reason is compatibility with newer drivers.

Advertised sites are not endorsed by the Bitcoin Forum. They may be unsafe, untrustworthy, or illegal in your jurisdiction. Advertise here.
1502912543
Hero Member
*
Offline Offline

Posts: 1502912543

View Profile Personal Message (Offline)

Ignore
1502912543
Reply with quote  #2

1502912543
Report to moderator
1502912543
Hero Member
*
Offline Offline

Posts: 1502912543

View Profile Personal Message (Offline)

Ignore
1502912543
Reply with quote  #2

1502912543
Report to moderator
utahjohn
Hero Member
*****
Offline Offline

Activity: 616


View Profile WWW
February 25, 2015, 07:26:18 PM
 #262

Thx pallas! But im out of the game cos i have 280s only. Utahjohn too (i think so).
Indeed 280x only here, I am in talks with Pallas to work on this further Smiley

DMD: dUTjohnrXHGYkh7jELWrZkGJbMnE6mdsuh (Staking)
BTC: 1HANJQygp3jHuzutceBgMT7wfCgEug6h4L (Donation)
ETH: 0xba90d7c1ab2bb9d5c07d843476153d1722637250 Mine ETH for 0.5% http://donkeypool.com
M1ST3R
Member
**
Offline Offline

Activity: 84


View Profile
February 27, 2015, 03:44:29 AM
 #263

experimental new bin for Hawaii (r9 290/290X) only:

https://dl.dropboxusercontent.com/u/40353042/Diamond/diamondHawaiiw128l8.bin

use worksize 128.

this is my opencl kernel, tweaked for speed and compatibility.
please report hashrates and show your support!

Hi Pallas,

The bin file is not working on both the sgminer 4.1.0 from Diamond website and sgminer 5 from Wolf0.
After the ...kernel is experimental... display, both sgminer version either hanged or display black screen.
Maybe the sgminer needs the specific v2 diamond.cl file to function properly.

BTW, unlike v1, changing the name of your .bin file to match the one sgminer generated does not work either.
pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
February 27, 2015, 09:30:03 AM
 #264

experimental new bin for Hawaii (r9 290/290X) only:

https://dl.dropboxusercontent.com/u/40353042/Diamond/diamondHawaiiw128l8.bin

use worksize 128.

this is my opencl kernel, tweaked for speed and compatibility.
please report hashrates and show your support!

Hi Pallas,

The bin file is not working on both the sgminer 4.1.0 from Diamond website and sgminer 5 from Wolf0.
After the ...kernel is experimental... display, both sgminer version either hanged or display black screen.
Maybe the sgminer needs the specific v2 diamond.cl file to function properly.

BTW, unlike v1, changing the name of your .bin file to match the one sgminer generated does not work either.

the binary can work without the sources.
check that you are running a 64 bit miner (the official diamond miner is 32 bit), that you are using worksize 128 and that you are setting the correct bin file name.
and of course that you have a hawaii card! :-)

utahjohn
Hero Member
*****
Offline Offline

Activity: 616


View Profile WWW
February 27, 2015, 09:33:34 AM
 #265

experimental new bin for Hawaii (r9 290/290X) only:

https://dl.dropboxusercontent.com/u/40353042/Diamond/diamondHawaiiw128l8.bin

use worksize 128.

this is my opencl kernel, tweaked for speed and compatibility.
please report hashrates and show your support!

Hi Pallas,

The bin file is not working on both the sgminer 4.1.0 from Diamond website and sgminer 5 from Wolf0.
After the ...kernel is experimental... display, both sgminer version either hanged or display black screen.
Maybe the sgminer needs the specific v2 diamond.cl file to function properly.

BTW, unlike v1, changing the name of your .bin file to match the one sgminer generated does not work either.

the binary can work without the sources.
check that you are running a 64 bit miner (the official diamond miner is 32 bit), that you are using worksize 128 and that you are setting the correct bin file name.
and of course that you have a hawaii card! :-)
Good reason to have OCL source LOL, I am 64 bit OS but my miner is 32 bit ...
In my experience blackscreen or just hung miner indicates too much O/C or memclock not set as recommended (GPU crash before it can be reported) ...

Without the OCL source I can not test on Tahiti properly so I have no definte answer ... u did not specify OS, config etc so a bit hard to troubleshoot ...

DMD: dUTjohnrXHGYkh7jELWrZkGJbMnE6mdsuh (Staking)
BTC: 1HANJQygp3jHuzutceBgMT7wfCgEug6h4L (Donation)
ETH: 0xba90d7c1ab2bb9d5c07d843476153d1722637250 Mine ETH for 0.5% http://donkeypool.com
sammy007
Legendary
*
Offline Offline

Activity: 1134



View Profile
February 27, 2015, 04:35:20 PM
 #266

Very nice results with 290(X), any chance for 280x gain?

pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
February 27, 2015, 04:41:36 PM
 #267

Very nice results with 290(X), any chance for 280x gain?

In order to do the same optimizations on 280(x), the code would need to be almost completely rewritten to work on 32 bit numbers instead of 64 (because of lds usage), hoping for the vgprs count (which is mostly in compiler control and very difficult to reduce by modifying the opencl code) to be low enough to permit 2 wavefronts.
Or, maybe, a better compiler in the future could do it by itself.
As of now, I think the best is using the asm version for 280(x) and my last binary for 290(x).

utahjohn
Hero Member
*****
Offline Offline

Activity: 616


View Profile WWW
February 27, 2015, 05:00:08 PM
 #268

Very nice results with 290(X), any chance for 280x gain?

In order to do the same optimizations on 280(x), the code would need to be almost completely rewritten to work on 32 bit numbers instead of 64 (because of lds usage), hoping for the vgprs count (which is mostly in compiler control and very difficult to reduce by modifying the opencl code) to be low enough to permit 2 wavefronts.
Or, maybe, a better compiler in the future could do it by itself.
As of now, I think the best is using the asm version for 280(x) and my last binary for 290(x).
Why do you think this, ASM version is driver independent and relies on directly coding for GPU, there is very little difference between 280x and 290, 290 has more shaders, true.   Buts basic code optimize such as your first/last pass should work in ASM just as well or better considering that AMD lobotomized OCL compiler after 14.7

DMD: dUTjohnrXHGYkh7jELWrZkGJbMnE6mdsuh (Staking)
BTC: 1HANJQygp3jHuzutceBgMT7wfCgEug6h4L (Donation)
ETH: 0xba90d7c1ab2bb9d5c07d843476153d1722637250 Mine ETH for 0.5% http://donkeypool.com
pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
February 27, 2015, 09:57:54 PM
 #269

Very nice results with 290(X), any chance for 280x gain?

In order to do the same optimizations on 280(x), the code would need to be almost completely rewritten to work on 32 bit numbers instead of 64 (because of lds usage), hoping for the vgprs count (which is mostly in compiler control and very difficult to reduce by modifying the opencl code) to be low enough to permit 2 wavefronts.
Or, maybe, a better compiler in the future could do it by itself.
As of now, I think the best is using the asm version for 280(x) and my last binary for 290(x).
Why do you think this, ASM version is driver independent and relies on directly coding for GPU, there is very little difference between 280x and 290, 290 has more shaders, true.   Buts basic code optimize such as your first/last pass should work in ASM just as well or better considering that AMD lobotomized OCL compiler after 14.7

14.12 is the first version making Hawaii specific code which, in some cases, may bring sensible improvements.
On Tahiti, the compiler simply can't make code capable of running 2 wavefronts. Or maybe it can but I'm not able to make it do it, on Hawaii I can instead.
Hawaii is not just Tahiti with more shaders...

utahjohn
Hero Member
*****
Offline Offline

Activity: 616


View Profile WWW
February 27, 2015, 10:11:06 PM
 #270

Very nice results with 290(X), any chance for 280x gain?

In order to do the same optimizations on 280(x), the code would need to be almost completely rewritten to work on 32 bit numbers instead of 64 (because of lds usage), hoping for the vgprs count (which is mostly in compiler control and very difficult to reduce by modifying the opencl code) to be low enough to permit 2 wavefronts.
Or, maybe, a better compiler in the future could do it by itself.
As of now, I think the best is using the asm version for 280(x) and my last binary for 290(x).
Why do you think this, ASM version is driver independent and relies on directly coding for GPU, there is very little difference between 280x and 290, 290 has more shaders, true.   Buts basic code optimize such as your first/last pass should work in ASM just as well or better considering that AMD lobotomized OCL compiler after 14.7

14.12 is the first version making Hawaii specific code which, in some cases, may bring sensible improvements.
On Tahiti, the compiler simply can't make code capable of running 2 wavefronts. Or maybe it can but I'm not able to make it do it, on Hawaii I can instead.
Hawaii is not just Tahiti with more shaders...
What are the differences, 280x (Tahiti)  can do multiple gpu-threads on many other coins (up to 4 gpu-threads on x11) with great efficiency, I do not understand why groesl can not.
Forgive me for asking such questions, but like my question about neoscrypt (which performs best with only 1 gpu-thread)  WS being totally tuned by amount of shaders ...

DMD: dUTjohnrXHGYkh7jELWrZkGJbMnE6mdsuh (Staking)
BTC: 1HANJQygp3jHuzutceBgMT7wfCgEug6h4L (Donation)
ETH: 0xba90d7c1ab2bb9d5c07d843476153d1722637250 Mine ETH for 0.5% http://donkeypool.com
pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
February 27, 2015, 10:34:32 PM
 #271

I believe multiple threads help with algos which use gpu ram: groestl does not. Only WS and intensity matter. TC is a buffer in ram so not relevant as well.

utahjohn
Hero Member
*****
Offline Offline

Activity: 616


View Profile WWW
February 27, 2015, 10:41:24 PM
 #272

OK that makes sense Smiley
So I can still play with WS in your new OCL? (I think WS may be card specific tuning).

DMD: dUTjohnrXHGYkh7jELWrZkGJbMnE6mdsuh (Staking)
BTC: 1HANJQygp3jHuzutceBgMT7wfCgEug6h4L (Donation)
ETH: 0xba90d7c1ab2bb9d5c07d843476153d1722637250 Mine ETH for 0.5% http://donkeypool.com
HR
Legendary
*
Offline Offline

Activity: 1036


Transparency & Integrity


View Profile
February 28, 2015, 09:35:46 AM
 #273


Pallas,

Are you planning on adding myriad-groestl support in the future? If not, could you explain why not? Is it because your groestl kernel is already faster than the myriad-groestl?

Also, are you planning on putting your work on github? Again, if not, could you explain why not?

It seems to me that both are important ways to further your efforts and establish your reputation.

Best regards as always.

HR

If there is only 4 times more LTC than BTC, when is LTC going to be worth 1/4 of BTC?
pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
February 28, 2015, 09:51:37 AM
 #274

Myr-groestl: If there is interest and I can get some free time I'd love to do it :-)
Github: for a couple files it's not worth, IMHO

utahjohn
Hero Member
*****
Offline Offline

Activity: 616


View Profile WWW
February 28, 2015, 10:41:20 AM
 #275

Myr-groestl: If there is interest and I can get some free time I'd love to do it :-)
Github: for a couple files it's not worth, IMHO
There is interest, I am leaving DMD groestl as I am sick of crap there.  HR has convinced me to move to Digibyte which is myriad-groestl?  IDK yet I have to d/l wallet and blockchain ...

DMD: dUTjohnrXHGYkh7jELWrZkGJbMnE6mdsuh (Staking)
BTC: 1HANJQygp3jHuzutceBgMT7wfCgEug6h4L (Donation)
ETH: 0xba90d7c1ab2bb9d5c07d843476153d1722637250 Mine ETH for 0.5% http://donkeypool.com
pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
February 28, 2015, 11:29:29 AM
 #276

Myr-groestl: If there is interest and I can get some free time I'd love to do it :-)
Github: for a couple files it's not worth, IMHO
There is interest, I am leaving DMD groestl as I am sick of crap there.  HR has convinced me to move to Digibyte which is myriad-groestl?  IDK yet I have to d/l wallet and blockchain ...

I think it's multi-algo including myr-groestl, skein etc.

Heavyiron
Newbie
*
Offline Offline

Activity: 18


View Profile
March 01, 2015, 07:06:44 PM
 #277

Hello pallas and thank you for your kernel.
I have groestl code from smelter (first GPU miner for quark). May be it have some tricks for your work. It was rather fast on radeon HD 5xxx series. But it is not adapted for sgminer and I have no skills to do this job.

Code:
#define CONSTANT __constant
#define LOCAL __local
#define GLOBAL __global
#define RESTRICT restrict
#define GLOBALID (uint)(get_global_id(0))
#define LOCALID get_local_id(0)

#define EXT_BYTE32_0(n) ((uint)(as_uchar4((uint)(n)).x))
#define EXT_BYTE32_1(n) ((uint)(as_uchar4((uint)(n)).y))
#define EXT_BYTE32_2(n) ((uint)(as_uchar4((uint)(n)).z))
#define EXT_BYTE32_3(n) ((uint)(as_uchar4((uint)(n)).w))

#define groestl_EXT_BYTE_0(n) EXT_BYTE32_0(n)
#define groestl_EXT_BYTE_1(n) EXT_BYTE32_1(n)
#define groestl_EXT_BYTE_2(n) EXT_BYTE32_2(n)
#define groestl_EXT_BYTE_3(n) EXT_BYTE32_3(n)


#define groestl_PMIX(src, dst, r)\
src[ 0] ^= (r);\
src[ 2] ^= 0x00000010u^(r);\
src[ 4] ^= 0x00000020u^(r);\
src[ 6] ^= 0x00000030u^(r);\
src[ 8] ^= 0x00000040u^(r);\
src[10] ^= 0x00000050u^(r);\
src[12] ^= 0x00000060u^(r);\
src[14] ^= 0x00000070u^(r);\
src[16] ^= 0x00000080u^(r);\
src[18] ^= 0x00000090u^(r);\
src[20] ^= 0x000000a0u^(r);\
src[22] ^= 0x000000b0u^(r);\
src[24] ^= 0x000000c0u^(r);\
src[26] ^= 0x000000d0u^(r);\
src[28] ^= 0x000000e0u^(r);\
src[30] ^= 0x000000f0u^(r);\
dst[ 0]  = groestl_T0[groestl_EXT_BYTE_0(src[ 0])];\
dst[ 1]  = groestl_T0[groestl_EXT_BYTE_0(src[ 9])];\
dst[ 2]  = groestl_T0[groestl_EXT_BYTE_0(src[ 2])];\
dst[ 3]  = groestl_T0[groestl_EXT_BYTE_0(src[11])];\
dst[ 4]  = groestl_T0[groestl_EXT_BYTE_0(src[ 4])];\
dst[ 5]  = groestl_T0[groestl_EXT_BYTE_0(src[13])];\
dst[ 6]  = groestl_T0[groestl_EXT_BYTE_0(src[ 6])];\
dst[ 7]  = groestl_T0[groestl_EXT_BYTE_0(src[15])];\
dst[ 8]  = groestl_T0[groestl_EXT_BYTE_0(src[ 8])];\
dst[ 9]  = groestl_T0[groestl_EXT_BYTE_0(src[17])];\
dst[10]  = groestl_T0[groestl_EXT_BYTE_0(src[10])];\
dst[11]  = groestl_T0[groestl_EXT_BYTE_0(src[19])];\
dst[12]  = groestl_T0[groestl_EXT_BYTE_0(src[12])];\
dst[13]  = groestl_T0[groestl_EXT_BYTE_0(src[21])];\
dst[14]  = groestl_T0[groestl_EXT_BYTE_0(src[14])];\
dst[15]  = groestl_T0[groestl_EXT_BYTE_0(src[23])];\
dst[16]  = groestl_T0[groestl_EXT_BYTE_0(src[16])];\
dst[17]  = groestl_T0[groestl_EXT_BYTE_0(src[25])];\
dst[18]  = groestl_T0[groestl_EXT_BYTE_0(src[18])];\
dst[19]  = groestl_T0[groestl_EXT_BYTE_0(src[27])];\
dst[20]  = groestl_T0[groestl_EXT_BYTE_0(src[20])];\
dst[21]  = groestl_T0[groestl_EXT_BYTE_0(src[29])];\
dst[22]  = groestl_T0[groestl_EXT_BYTE_0(src[22])];\
dst[23]  = groestl_T0[groestl_EXT_BYTE_0(src[31])];\
dst[24]  = groestl_T0[groestl_EXT_BYTE_0(src[24])];\
dst[25]  = groestl_T0[groestl_EXT_BYTE_0(src[ 1])];\
dst[26]  = groestl_T0[groestl_EXT_BYTE_0(src[26])];\
dst[27]  = groestl_T0[groestl_EXT_BYTE_0(src[ 3])];\
dst[28]  = groestl_T0[groestl_EXT_BYTE_0(src[28])];\
dst[29]  = groestl_T0[groestl_EXT_BYTE_0(src[ 5])];\
dst[30]  = groestl_T0[groestl_EXT_BYTE_0(src[30])];\
dst[31]  = groestl_T0[groestl_EXT_BYTE_0(src[ 7])];\
dst[ 0] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 2])];\
dst[ 1] ^= groestl_T1[groestl_EXT_BYTE_1(src[11])];\
dst[ 2] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 4])];\
dst[ 3] ^= groestl_T1[groestl_EXT_BYTE_1(src[13])];\
dst[ 4] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 6])];\
dst[ 5] ^= groestl_T1[groestl_EXT_BYTE_1(src[15])];\
dst[ 6] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 8])];\
dst[ 7] ^= groestl_T1[groestl_EXT_BYTE_1(src[17])];\
dst[ 8] ^= groestl_T1[groestl_EXT_BYTE_1(src[10])];\
dst[ 9] ^= groestl_T1[groestl_EXT_BYTE_1(src[19])];\
dst[10] ^= groestl_T1[groestl_EXT_BYTE_1(src[12])];\
dst[11] ^= groestl_T1[groestl_EXT_BYTE_1(src[21])];\
dst[12] ^= groestl_T1[groestl_EXT_BYTE_1(src[14])];\
dst[13] ^= groestl_T1[groestl_EXT_BYTE_1(src[23])];\
dst[14] ^= groestl_T1[groestl_EXT_BYTE_1(src[16])];\
dst[15] ^= groestl_T1[groestl_EXT_BYTE_1(src[25])];\
dst[16] ^= groestl_T1[groestl_EXT_BYTE_1(src[18])];\
dst[17] ^= groestl_T1[groestl_EXT_BYTE_1(src[27])];\
dst[18] ^= groestl_T1[groestl_EXT_BYTE_1(src[20])];\
dst[19] ^= groestl_T1[groestl_EXT_BYTE_1(src[29])];\
dst[20] ^= groestl_T1[groestl_EXT_BYTE_1(src[22])];\
dst[21] ^= groestl_T1[groestl_EXT_BYTE_1(src[31])];\
dst[22] ^= groestl_T1[groestl_EXT_BYTE_1(src[24])];\
dst[23] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 1])];\
dst[24] ^= groestl_T1[groestl_EXT_BYTE_1(src[26])];\
dst[25] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 3])];\
dst[26] ^= groestl_T1[groestl_EXT_BYTE_1(src[28])];\
dst[27] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 5])];\
dst[28] ^= groestl_T1[groestl_EXT_BYTE_1(src[30])];\
dst[29] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 7])];\
dst[30] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 0])];\
dst[31] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 9])];\
dst[ 0] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 4])];\
dst[ 1] ^= groestl_T2[groestl_EXT_BYTE_2(src[13])];\
dst[ 2] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 6])];\
dst[ 3] ^= groestl_T2[groestl_EXT_BYTE_2(src[15])];\
dst[ 4] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 8])];\
dst[ 5] ^= groestl_T2[groestl_EXT_BYTE_2(src[17])];\
dst[ 6] ^= groestl_T2[groestl_EXT_BYTE_2(src[10])];\
dst[ 7] ^= groestl_T2[groestl_EXT_BYTE_2(src[19])];\
dst[ 8] ^= groestl_T2[groestl_EXT_BYTE_2(src[12])];\
dst[ 9] ^= groestl_T2[groestl_EXT_BYTE_2(src[21])];\
dst[10] ^= groestl_T2[groestl_EXT_BYTE_2(src[14])];\
dst[11] ^= groestl_T2[groestl_EXT_BYTE_2(src[23])];\
dst[12] ^= groestl_T2[groestl_EXT_BYTE_2(src[16])];\
dst[13] ^= groestl_T2[groestl_EXT_BYTE_2(src[25])];\
dst[14] ^= groestl_T2[groestl_EXT_BYTE_2(src[18])];\
dst[15] ^= groestl_T2[groestl_EXT_BYTE_2(src[27])];\
dst[16] ^= groestl_T2[groestl_EXT_BYTE_2(src[20])];\
dst[17] ^= groestl_T2[groestl_EXT_BYTE_2(src[29])];\
dst[18] ^= groestl_T2[groestl_EXT_BYTE_2(src[22])];\
dst[19] ^= groestl_T2[groestl_EXT_BYTE_2(src[31])];\
dst[20] ^= groestl_T2[groestl_EXT_BYTE_2(src[24])];\
dst[21] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 1])];\
dst[22] ^= groestl_T2[groestl_EXT_BYTE_2(src[26])];\
dst[23] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 3])];\
dst[24] ^= groestl_T2[groestl_EXT_BYTE_2(src[28])];\
dst[25] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 5])];\
dst[26] ^= groestl_T2[groestl_EXT_BYTE_2(src[30])];\
dst[27] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 7])];\
dst[28] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 0])];\
dst[29] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 9])];\
dst[30] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 2])];\
dst[31] ^= groestl_T2[groestl_EXT_BYTE_2(src[11])];\
dst[ 0] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 6])];\
dst[ 1] ^= groestl_T3[groestl_EXT_BYTE_3(src[23])];\
dst[ 2] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 8])];\
dst[ 3] ^= groestl_T3[groestl_EXT_BYTE_3(src[25])];\
dst[ 4] ^= groestl_T3[groestl_EXT_BYTE_3(src[10])];\
dst[ 5] ^= groestl_T3[groestl_EXT_BYTE_3(src[27])];\
dst[ 6] ^= groestl_T3[groestl_EXT_BYTE_3(src[12])];\
dst[ 7] ^= groestl_T3[groestl_EXT_BYTE_3(src[29])];\
dst[ 8] ^= groestl_T3[groestl_EXT_BYTE_3(src[14])];\
dst[ 9] ^= groestl_T3[groestl_EXT_BYTE_3(src[31])];\
dst[10] ^= groestl_T3[groestl_EXT_BYTE_3(src[16])];\
dst[11] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 1])];\
dst[12] ^= groestl_T3[groestl_EXT_BYTE_3(src[18])];\
dst[13] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 3])];\
dst[14] ^= groestl_T3[groestl_EXT_BYTE_3(src[20])];\
dst[15] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 5])];\
dst[16] ^= groestl_T3[groestl_EXT_BYTE_3(src[22])];\
dst[17] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 7])];\
dst[18] ^= groestl_T3[groestl_EXT_BYTE_3(src[24])];\
dst[19] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 9])];\
dst[20] ^= groestl_T3[groestl_EXT_BYTE_3(src[26])];\
dst[21] ^= groestl_T3[groestl_EXT_BYTE_3(src[11])];\
dst[22] ^= groestl_T3[groestl_EXT_BYTE_3(src[28])];\
dst[23] ^= groestl_T3[groestl_EXT_BYTE_3(src[13])];\
dst[24] ^= groestl_T3[groestl_EXT_BYTE_3(src[30])];\
dst[25] ^= groestl_T3[groestl_EXT_BYTE_3(src[15])];\
dst[26] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 0])];\
dst[27] ^= groestl_T3[groestl_EXT_BYTE_3(src[17])];\
dst[28] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 2])];\
dst[29] ^= groestl_T3[groestl_EXT_BYTE_3(src[19])];\
dst[30] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 4])];\
dst[31] ^= groestl_T3[groestl_EXT_BYTE_3(src[21])];\
dst[ 0] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 9])];\
dst[ 1] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 0])];\
dst[ 2] ^= groestl_T4[groestl_EXT_BYTE_0(src[11])];\
dst[ 3] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 2])];\
dst[ 4] ^= groestl_T4[groestl_EXT_BYTE_0(src[13])];\
dst[ 5] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 4])];\
dst[ 6] ^= groestl_T4[groestl_EXT_BYTE_0(src[15])];\
dst[ 7] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 6])];\
dst[ 8] ^= groestl_T4[groestl_EXT_BYTE_0(src[17])];\
dst[ 9] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 8])];\
dst[10] ^= groestl_T4[groestl_EXT_BYTE_0(src[19])];\
dst[11] ^= groestl_T4[groestl_EXT_BYTE_0(src[10])];\
dst[12] ^= groestl_T4[groestl_EXT_BYTE_0(src[21])];\
dst[13] ^= groestl_T4[groestl_EXT_BYTE_0(src[12])];\
dst[14] ^= groestl_T4[groestl_EXT_BYTE_0(src[23])];\
dst[15] ^= groestl_T4[groestl_EXT_BYTE_0(src[14])];\
dst[16] ^= groestl_T4[groestl_EXT_BYTE_0(src[25])];\
dst[17] ^= groestl_T4[groestl_EXT_BYTE_0(src[16])];\
dst[18] ^= groestl_T4[groestl_EXT_BYTE_0(src[27])];\
dst[19] ^= groestl_T4[groestl_EXT_BYTE_0(src[18])];\
dst[20] ^= groestl_T4[groestl_EXT_BYTE_0(src[29])];\
dst[21] ^= groestl_T4[groestl_EXT_BYTE_0(src[20])];\
dst[22] ^= groestl_T4[groestl_EXT_BYTE_0(src[31])];\
dst[23] ^= groestl_T4[groestl_EXT_BYTE_0(src[22])];\
dst[24] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 1])];\
dst[25] ^= groestl_T4[groestl_EXT_BYTE_0(src[24])];\
dst[26] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 3])];\
dst[27] ^= groestl_T4[groestl_EXT_BYTE_0(src[26])];\
dst[28] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 5])];\
dst[29] ^= groestl_T4[groestl_EXT_BYTE_0(src[28])];\
dst[30] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 7])];\
dst[31] ^= groestl_T4[groestl_EXT_BYTE_0(src[30])];\
dst[ 0] ^= groestl_T5[groestl_EXT_BYTE_1(src[11])];\
dst[ 1] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 2])];\
dst[ 2] ^= groestl_T5[groestl_EXT_BYTE_1(src[13])];\
dst[ 3] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 4])];\
dst[ 4] ^= groestl_T5[groestl_EXT_BYTE_1(src[15])];\
dst[ 5] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 6])];\
dst[ 6] ^= groestl_T5[groestl_EXT_BYTE_1(src[17])];\
dst[ 7] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 8])];\
dst[ 8] ^= groestl_T5[groestl_EXT_BYTE_1(src[19])];\
dst[ 9] ^= groestl_T5[groestl_EXT_BYTE_1(src[10])];\
dst[10] ^= groestl_T5[groestl_EXT_BYTE_1(src[21])];\
dst[11] ^= groestl_T5[groestl_EXT_BYTE_1(src[12])];\
dst[12] ^= groestl_T5[groestl_EXT_BYTE_1(src[23])];\
dst[13] ^= groestl_T5[groestl_EXT_BYTE_1(src[14])];\
dst[14] ^= groestl_T5[groestl_EXT_BYTE_1(src[25])];\
dst[15] ^= groestl_T5[groestl_EXT_BYTE_1(src[16])];\
dst[16] ^= groestl_T5[groestl_EXT_BYTE_1(src[27])];\
dst[17] ^= groestl_T5[groestl_EXT_BYTE_1(src[18])];\
dst[18] ^= groestl_T5[groestl_EXT_BYTE_1(src[29])];\
dst[19] ^= groestl_T5[groestl_EXT_BYTE_1(src[20])];\
dst[20] ^= groestl_T5[groestl_EXT_BYTE_1(src[31])];\
dst[21] ^= groestl_T5[groestl_EXT_BYTE_1(src[22])];\
dst[22] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 1])];\
dst[23] ^= groestl_T5[groestl_EXT_BYTE_1(src[24])];\
dst[24] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 3])];\
dst[25] ^= groestl_T5[groestl_EXT_BYTE_1(src[26])];\
dst[26] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 5])];\
dst[27] ^= groestl_T5[groestl_EXT_BYTE_1(src[28])];\
dst[28] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 7])];\
dst[29] ^= groestl_T5[groestl_EXT_BYTE_1(src[30])];\
dst[30] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 9])];\
dst[31] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 0])];\
dst[ 0] ^= groestl_T6[groestl_EXT_BYTE_2(src[13])];\
dst[ 1] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 4])];\
dst[ 2] ^= groestl_T6[groestl_EXT_BYTE_2(src[15])];\
dst[ 3] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 6])];\
dst[ 4] ^= groestl_T6[groestl_EXT_BYTE_2(src[17])];\
dst[ 5] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 8])];\
dst[ 6] ^= groestl_T6[groestl_EXT_BYTE_2(src[19])];\
dst[ 7] ^= groestl_T6[groestl_EXT_BYTE_2(src[10])];\
dst[ 8] ^= groestl_T6[groestl_EXT_BYTE_2(src[21])];\
dst[ 9] ^= groestl_T6[groestl_EXT_BYTE_2(src[12])];\
dst[10] ^= groestl_T6[groestl_EXT_BYTE_2(src[23])];\
dst[11] ^= groestl_T6[groestl_EXT_BYTE_2(src[14])];\
dst[12] ^= groestl_T6[groestl_EXT_BYTE_2(src[25])];\
dst[13] ^= groestl_T6[groestl_EXT_BYTE_2(src[16])];\
dst[14] ^= groestl_T6[groestl_EXT_BYTE_2(src[27])];\
dst[15] ^= groestl_T6[groestl_EXT_BYTE_2(src[18])];\
dst[16] ^= groestl_T6[groestl_EXT_BYTE_2(src[29])];\
dst[17] ^= groestl_T6[groestl_EXT_BYTE_2(src[20])];\
dst[18] ^= groestl_T6[groestl_EXT_BYTE_2(src[31])];\
dst[19] ^= groestl_T6[groestl_EXT_BYTE_2(src[22])];\
dst[20] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 1])];\
dst[21] ^= groestl_T6[groestl_EXT_BYTE_2(src[24])];\
dst[22] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 3])];\
dst[23] ^= groestl_T6[groestl_EXT_BYTE_2(src[26])];\
dst[24] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 5])];\
dst[25] ^= groestl_T6[groestl_EXT_BYTE_2(src[28])];\
dst[26] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 7])];\
dst[27] ^= groestl_T6[groestl_EXT_BYTE_2(src[30])];\
dst[28] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 9])];\
dst[29] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 0])];\
dst[30] ^= groestl_T6[groestl_EXT_BYTE_2(src[11])];\
dst[31] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 2])];\
dst[ 0] ^= groestl_T7[groestl_EXT_BYTE_3(src[23])];\
dst[ 1] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 6])];\
dst[ 2] ^= groestl_T7[groestl_EXT_BYTE_3(src[25])];\
dst[ 3] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 8])];\
dst[ 4] ^= groestl_T7[groestl_EXT_BYTE_3(src[27])];\
dst[ 5] ^= groestl_T7[groestl_EXT_BYTE_3(src[10])];\
dst[ 6] ^= groestl_T7[groestl_EXT_BYTE_3(src[29])];\
dst[ 7] ^= groestl_T7[groestl_EXT_BYTE_3(src[12])];\
dst[ 8] ^= groestl_T7[groestl_EXT_BYTE_3(src[31])];\
dst[ 9] ^= groestl_T7[groestl_EXT_BYTE_3(src[14])];\
dst[10] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 1])];\
dst[11] ^= groestl_T7[groestl_EXT_BYTE_3(src[16])];\
dst[12] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 3])];\
dst[13] ^= groestl_T7[groestl_EXT_BYTE_3(src[18])];\
dst[14] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 5])];\
dst[15] ^= groestl_T7[groestl_EXT_BYTE_3(src[20])];\
dst[16] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 7])];\
dst[17] ^= groestl_T7[groestl_EXT_BYTE_3(src[22])];\
dst[18] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 9])];\
dst[19] ^= groestl_T7[groestl_EXT_BYTE_3(src[24])];\
dst[20] ^= groestl_T7[groestl_EXT_BYTE_3(src[11])];\
dst[21] ^= groestl_T7[groestl_EXT_BYTE_3(src[26])];\
dst[22] ^= groestl_T7[groestl_EXT_BYTE_3(src[13])];\
dst[23] ^= groestl_T7[groestl_EXT_BYTE_3(src[28])];\
dst[24] ^= groestl_T7[groestl_EXT_BYTE_3(src[15])];\
dst[25] ^= groestl_T7[groestl_EXT_BYTE_3(src[30])];\
dst[26] ^= groestl_T7[groestl_EXT_BYTE_3(src[17])];\
dst[27] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 0])];\
dst[28] ^= groestl_T7[groestl_EXT_BYTE_3(src[19])];\
dst[29] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 2])];\
dst[30] ^= groestl_T7[groestl_EXT_BYTE_3(src[21])];\
dst[31] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 4])];

#define groestl_QMIX(src, dst, r)\
src[ 0] = ~src[ 0];\
src[ 1] ^= ~(r);\
src[ 2] = ~src[ 2];\
src[ 3] ^= 0xefffffffu^(r);\
src[ 4] = ~src[ 4];\
src[ 5] ^= 0xdfffffffu^(r);\
src[ 6] = ~src[ 6];\
src[ 7] ^= 0xcfffffffu^(r);\
src[ 8] = ~src[ 8];\
src[ 9] ^= 0xbfffffffu^(r);\
src[10] = ~src[10];\
src[11] ^= 0xafffffffu^(r);\
src[12] = ~src[12];\
src[13] ^= 0x9fffffffu^(r);\
src[14] = ~src[14];\
src[15] ^= 0x8fffffffu^(r);\
src[16] = ~src[16];\
src[17] ^= 0x7fffffffu^(r);\
src[18] = ~src[18];\
src[19] ^= 0x6fffffffu^(r);\
src[20] = ~src[20];\
src[21] ^= 0x5fffffffu^(r);\
src[22] = ~src[22];\
src[23] ^= 0x4fffffffu^(r);\
src[24] = ~src[24];\
src[25] ^= 0x3fffffffu^(r);\
src[26] = ~src[26];\
src[27] ^= 0x2fffffffu^(r);\
src[28] = ~src[28];\
src[29] ^= 0x1fffffffu^(r);\
src[30] = ~src[30];\
src[31] ^= 0x0fffffffu^(r);\
dst[ 0]  = groestl_T0[groestl_EXT_BYTE_0(src[ 2])];\
dst[ 1]  = groestl_T0[groestl_EXT_BYTE_0(src[ 1])];\
dst[ 2]  = groestl_T0[groestl_EXT_BYTE_0(src[ 4])];\
dst[ 3]  = groestl_T0[groestl_EXT_BYTE_0(src[ 3])];\
dst[ 4]  = groestl_T0[groestl_EXT_BYTE_0(src[ 6])];\
dst[ 5]  = groestl_T0[groestl_EXT_BYTE_0(src[ 5])];\
dst[ 6]  = groestl_T0[groestl_EXT_BYTE_0(src[ 8])];\
dst[ 7]  = groestl_T0[groestl_EXT_BYTE_0(src[ 7])];\
dst[ 8]  = groestl_T0[groestl_EXT_BYTE_0(src[10])];\
dst[ 9]  = groestl_T0[groestl_EXT_BYTE_0(src[ 9])];\
dst[10]  = groestl_T0[groestl_EXT_BYTE_0(src[12])];\
dst[11]  = groestl_T0[groestl_EXT_BYTE_0(src[11])];\
dst[12]  = groestl_T0[groestl_EXT_BYTE_0(src[14])];\
dst[13]  = groestl_T0[groestl_EXT_BYTE_0(src[13])];\
dst[14]  = groestl_T0[groestl_EXT_BYTE_0(src[16])];\
dst[15]  = groestl_T0[groestl_EXT_BYTE_0(src[15])];\
dst[16]  = groestl_T0[groestl_EXT_BYTE_0(src[18])];\
dst[17]  = groestl_T0[groestl_EXT_BYTE_0(src[17])];\
dst[18]  = groestl_T0[groestl_EXT_BYTE_0(src[20])];\
dst[19]  = groestl_T0[groestl_EXT_BYTE_0(src[19])];\
dst[20]  = groestl_T0[groestl_EXT_BYTE_0(src[22])];\
dst[21]  = groestl_T0[groestl_EXT_BYTE_0(src[21])];\
dst[22]  = groestl_T0[groestl_EXT_BYTE_0(src[24])];\
dst[23]  = groestl_T0[groestl_EXT_BYTE_0(src[23])];\
dst[24]  = groestl_T0[groestl_EXT_BYTE_0(src[26])];\
dst[25]  = groestl_T0[groestl_EXT_BYTE_0(src[25])];\
dst[26]  = groestl_T0[groestl_EXT_BYTE_0(src[28])];\
dst[27]  = groestl_T0[groestl_EXT_BYTE_0(src[27])];\
dst[28]  = groestl_T0[groestl_EXT_BYTE_0(src[30])];\
dst[29]  = groestl_T0[groestl_EXT_BYTE_0(src[29])];\
dst[30]  = groestl_T0[groestl_EXT_BYTE_0(src[ 0])];\
dst[31]  = groestl_T0[groestl_EXT_BYTE_0(src[31])];\
dst[ 0] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 6])];\
dst[ 1] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 5])];\
dst[ 2] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 8])];\
dst[ 3] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 7])];\
dst[ 4] ^= groestl_T1[groestl_EXT_BYTE_1(src[10])];\
dst[ 5] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 9])];\
dst[ 6] ^= groestl_T1[groestl_EXT_BYTE_1(src[12])];\
dst[ 7] ^= groestl_T1[groestl_EXT_BYTE_1(src[11])];\
dst[ 8] ^= groestl_T1[groestl_EXT_BYTE_1(src[14])];\
dst[ 9] ^= groestl_T1[groestl_EXT_BYTE_1(src[13])];\
dst[10] ^= groestl_T1[groestl_EXT_BYTE_1(src[16])];\
dst[11] ^= groestl_T1[groestl_EXT_BYTE_1(src[15])];\
dst[12] ^= groestl_T1[groestl_EXT_BYTE_1(src[18])];\
dst[13] ^= groestl_T1[groestl_EXT_BYTE_1(src[17])];\
dst[14] ^= groestl_T1[groestl_EXT_BYTE_1(src[20])];\
dst[15] ^= groestl_T1[groestl_EXT_BYTE_1(src[19])];\
dst[16] ^= groestl_T1[groestl_EXT_BYTE_1(src[22])];\
dst[17] ^= groestl_T1[groestl_EXT_BYTE_1(src[21])];\
dst[18] ^= groestl_T1[groestl_EXT_BYTE_1(src[24])];\
dst[19] ^= groestl_T1[groestl_EXT_BYTE_1(src[23])];\
dst[20] ^= groestl_T1[groestl_EXT_BYTE_1(src[26])];\
dst[21] ^= groestl_T1[groestl_EXT_BYTE_1(src[25])];\
dst[22] ^= groestl_T1[groestl_EXT_BYTE_1(src[28])];\
dst[23] ^= groestl_T1[groestl_EXT_BYTE_1(src[27])];\
dst[24] ^= groestl_T1[groestl_EXT_BYTE_1(src[30])];\
dst[25] ^= groestl_T1[groestl_EXT_BYTE_1(src[29])];\
dst[26] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 0])];\
dst[27] ^= groestl_T1[groestl_EXT_BYTE_1(src[31])];\
dst[28] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 2])];\
dst[29] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 1])];\
dst[30] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 4])];\
dst[31] ^= groestl_T1[groestl_EXT_BYTE_1(src[ 3])];\
dst[ 0] ^= groestl_T2[groestl_EXT_BYTE_2(src[10])];\
dst[ 1] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 9])];\
dst[ 2] ^= groestl_T2[groestl_EXT_BYTE_2(src[12])];\
dst[ 3] ^= groestl_T2[groestl_EXT_BYTE_2(src[11])];\
dst[ 4] ^= groestl_T2[groestl_EXT_BYTE_2(src[14])];\
dst[ 5] ^= groestl_T2[groestl_EXT_BYTE_2(src[13])];\
dst[ 6] ^= groestl_T2[groestl_EXT_BYTE_2(src[16])];\
dst[ 7] ^= groestl_T2[groestl_EXT_BYTE_2(src[15])];\
dst[ 8] ^= groestl_T2[groestl_EXT_BYTE_2(src[18])];\
dst[ 9] ^= groestl_T2[groestl_EXT_BYTE_2(src[17])];\
dst[10] ^= groestl_T2[groestl_EXT_BYTE_2(src[20])];\
dst[11] ^= groestl_T2[groestl_EXT_BYTE_2(src[19])];\
dst[12] ^= groestl_T2[groestl_EXT_BYTE_2(src[22])];\
dst[13] ^= groestl_T2[groestl_EXT_BYTE_2(src[21])];\
dst[14] ^= groestl_T2[groestl_EXT_BYTE_2(src[24])];\
dst[15] ^= groestl_T2[groestl_EXT_BYTE_2(src[23])];\
dst[16] ^= groestl_T2[groestl_EXT_BYTE_2(src[26])];\
dst[17] ^= groestl_T2[groestl_EXT_BYTE_2(src[25])];\
dst[18] ^= groestl_T2[groestl_EXT_BYTE_2(src[28])];\
dst[19] ^= groestl_T2[groestl_EXT_BYTE_2(src[27])];\
dst[20] ^= groestl_T2[groestl_EXT_BYTE_2(src[30])];\
dst[21] ^= groestl_T2[groestl_EXT_BYTE_2(src[29])];\
dst[22] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 0])];\
dst[23] ^= groestl_T2[groestl_EXT_BYTE_2(src[31])];\
dst[24] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 2])];\
dst[25] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 1])];\
dst[26] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 4])];\
dst[27] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 3])];\
dst[28] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 6])];\
dst[29] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 5])];\
dst[30] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 8])];\
dst[31] ^= groestl_T2[groestl_EXT_BYTE_2(src[ 7])];\
dst[ 0] ^= groestl_T3[groestl_EXT_BYTE_3(src[22])];\
dst[ 1] ^= groestl_T3[groestl_EXT_BYTE_3(src[13])];\
dst[ 2] ^= groestl_T3[groestl_EXT_BYTE_3(src[24])];\
dst[ 3] ^= groestl_T3[groestl_EXT_BYTE_3(src[15])];\
dst[ 4] ^= groestl_T3[groestl_EXT_BYTE_3(src[26])];\
dst[ 5] ^= groestl_T3[groestl_EXT_BYTE_3(src[17])];\
dst[ 6] ^= groestl_T3[groestl_EXT_BYTE_3(src[28])];\
dst[ 7] ^= groestl_T3[groestl_EXT_BYTE_3(src[19])];\
dst[ 8] ^= groestl_T3[groestl_EXT_BYTE_3(src[30])];\
dst[ 9] ^= groestl_T3[groestl_EXT_BYTE_3(src[21])];\
dst[10] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 0])];\
dst[11] ^= groestl_T3[groestl_EXT_BYTE_3(src[23])];\
dst[12] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 2])];\
dst[13] ^= groestl_T3[groestl_EXT_BYTE_3(src[25])];\
dst[14] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 4])];\
dst[15] ^= groestl_T3[groestl_EXT_BYTE_3(src[27])];\
dst[16] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 6])];\
dst[17] ^= groestl_T3[groestl_EXT_BYTE_3(src[29])];\
dst[18] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 8])];\
dst[19] ^= groestl_T3[groestl_EXT_BYTE_3(src[31])];\
dst[20] ^= groestl_T3[groestl_EXT_BYTE_3(src[10])];\
dst[21] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 1])];\
dst[22] ^= groestl_T3[groestl_EXT_BYTE_3(src[12])];\
dst[23] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 3])];\
dst[24] ^= groestl_T3[groestl_EXT_BYTE_3(src[14])];\
dst[25] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 5])];\
dst[26] ^= groestl_T3[groestl_EXT_BYTE_3(src[16])];\
dst[27] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 7])];\
dst[28] ^= groestl_T3[groestl_EXT_BYTE_3(src[18])];\
dst[29] ^= groestl_T3[groestl_EXT_BYTE_3(src[ 9])];\
dst[30] ^= groestl_T3[groestl_EXT_BYTE_3(src[20])];\
dst[31] ^= groestl_T3[groestl_EXT_BYTE_3(src[11])];\
dst[ 0] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 1])];\
dst[ 1] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 2])];\
dst[ 2] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 3])];\
dst[ 3] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 4])];\
dst[ 4] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 5])];\
dst[ 5] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 6])];\
dst[ 6] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 7])];\
dst[ 7] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 8])];\
dst[ 8] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 9])];\
dst[ 9] ^= groestl_T4[groestl_EXT_BYTE_0(src[10])];\
dst[10] ^= groestl_T4[groestl_EXT_BYTE_0(src[11])];\
dst[11] ^= groestl_T4[groestl_EXT_BYTE_0(src[12])];\
dst[12] ^= groestl_T4[groestl_EXT_BYTE_0(src[13])];\
dst[13] ^= groestl_T4[groestl_EXT_BYTE_0(src[14])];\
dst[14] ^= groestl_T4[groestl_EXT_BYTE_0(src[15])];\
dst[15] ^= groestl_T4[groestl_EXT_BYTE_0(src[16])];\
dst[16] ^= groestl_T4[groestl_EXT_BYTE_0(src[17])];\
dst[17] ^= groestl_T4[groestl_EXT_BYTE_0(src[18])];\
dst[18] ^= groestl_T4[groestl_EXT_BYTE_0(src[19])];\
dst[19] ^= groestl_T4[groestl_EXT_BYTE_0(src[20])];\
dst[20] ^= groestl_T4[groestl_EXT_BYTE_0(src[21])];\
dst[21] ^= groestl_T4[groestl_EXT_BYTE_0(src[22])];\
dst[22] ^= groestl_T4[groestl_EXT_BYTE_0(src[23])];\
dst[23] ^= groestl_T4[groestl_EXT_BYTE_0(src[24])];\
dst[24] ^= groestl_T4[groestl_EXT_BYTE_0(src[25])];\
dst[25] ^= groestl_T4[groestl_EXT_BYTE_0(src[26])];\
dst[26] ^= groestl_T4[groestl_EXT_BYTE_0(src[27])];\
dst[27] ^= groestl_T4[groestl_EXT_BYTE_0(src[28])];\
dst[28] ^= groestl_T4[groestl_EXT_BYTE_0(src[29])];\
dst[29] ^= groestl_T4[groestl_EXT_BYTE_0(src[30])];\
dst[30] ^= groestl_T4[groestl_EXT_BYTE_0(src[31])];\
dst[31] ^= groestl_T4[groestl_EXT_BYTE_0(src[ 0])];\
dst[ 0] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 5])];\
dst[ 1] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 6])];\
dst[ 2] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 7])];\
dst[ 3] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 8])];\
dst[ 4] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 9])];\
dst[ 5] ^= groestl_T5[groestl_EXT_BYTE_1(src[10])];\
dst[ 6] ^= groestl_T5[groestl_EXT_BYTE_1(src[11])];\
dst[ 7] ^= groestl_T5[groestl_EXT_BYTE_1(src[12])];\
dst[ 8] ^= groestl_T5[groestl_EXT_BYTE_1(src[13])];\
dst[ 9] ^= groestl_T5[groestl_EXT_BYTE_1(src[14])];\
dst[10] ^= groestl_T5[groestl_EXT_BYTE_1(src[15])];\
dst[11] ^= groestl_T5[groestl_EXT_BYTE_1(src[16])];\
dst[12] ^= groestl_T5[groestl_EXT_BYTE_1(src[17])];\
dst[13] ^= groestl_T5[groestl_EXT_BYTE_1(src[18])];\
dst[14] ^= groestl_T5[groestl_EXT_BYTE_1(src[19])];\
dst[15] ^= groestl_T5[groestl_EXT_BYTE_1(src[20])];\
dst[16] ^= groestl_T5[groestl_EXT_BYTE_1(src[21])];\
dst[17] ^= groestl_T5[groestl_EXT_BYTE_1(src[22])];\
dst[18] ^= groestl_T5[groestl_EXT_BYTE_1(src[23])];\
dst[19] ^= groestl_T5[groestl_EXT_BYTE_1(src[24])];\
dst[20] ^= groestl_T5[groestl_EXT_BYTE_1(src[25])];\
dst[21] ^= groestl_T5[groestl_EXT_BYTE_1(src[26])];\
dst[22] ^= groestl_T5[groestl_EXT_BYTE_1(src[27])];\
dst[23] ^= groestl_T5[groestl_EXT_BYTE_1(src[28])];\
dst[24] ^= groestl_T5[groestl_EXT_BYTE_1(src[29])];\
dst[25] ^= groestl_T5[groestl_EXT_BYTE_1(src[30])];\
dst[26] ^= groestl_T5[groestl_EXT_BYTE_1(src[31])];\
dst[27] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 0])];\
dst[28] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 1])];\
dst[29] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 2])];\
dst[30] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 3])];\
dst[31] ^= groestl_T5[groestl_EXT_BYTE_1(src[ 4])];\
dst[ 0] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 9])];\
dst[ 1] ^= groestl_T6[groestl_EXT_BYTE_2(src[10])];\
dst[ 2] ^= groestl_T6[groestl_EXT_BYTE_2(src[11])];\
dst[ 3] ^= groestl_T6[groestl_EXT_BYTE_2(src[12])];\
dst[ 4] ^= groestl_T6[groestl_EXT_BYTE_2(src[13])];\
dst[ 5] ^= groestl_T6[groestl_EXT_BYTE_2(src[14])];\
dst[ 6] ^= groestl_T6[groestl_EXT_BYTE_2(src[15])];\
dst[ 7] ^= groestl_T6[groestl_EXT_BYTE_2(src[16])];\
dst[ 8] ^= groestl_T6[groestl_EXT_BYTE_2(src[17])];\
dst[ 9] ^= groestl_T6[groestl_EXT_BYTE_2(src[18])];\
dst[10] ^= groestl_T6[groestl_EXT_BYTE_2(src[19])];\
dst[11] ^= groestl_T6[groestl_EXT_BYTE_2(src[20])];\
dst[12] ^= groestl_T6[groestl_EXT_BYTE_2(src[21])];\
dst[13] ^= groestl_T6[groestl_EXT_BYTE_2(src[22])];\
dst[14] ^= groestl_T6[groestl_EXT_BYTE_2(src[23])];\
dst[15] ^= groestl_T6[groestl_EXT_BYTE_2(src[24])];\
dst[16] ^= groestl_T6[groestl_EXT_BYTE_2(src[25])];\
dst[17] ^= groestl_T6[groestl_EXT_BYTE_2(src[26])];\
dst[18] ^= groestl_T6[groestl_EXT_BYTE_2(src[27])];\
dst[19] ^= groestl_T6[groestl_EXT_BYTE_2(src[28])];\
dst[20] ^= groestl_T6[groestl_EXT_BYTE_2(src[29])];\
dst[21] ^= groestl_T6[groestl_EXT_BYTE_2(src[30])];\
dst[22] ^= groestl_T6[groestl_EXT_BYTE_2(src[31])];\
dst[23] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 0])];\
dst[24] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 1])];\
dst[25] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 2])];\
dst[26] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 3])];\
dst[27] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 4])];\
dst[28] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 5])];\
dst[29] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 6])];\
dst[30] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 7])];\
dst[31] ^= groestl_T6[groestl_EXT_BYTE_2(src[ 8])];\
dst[ 0] ^= groestl_T7[groestl_EXT_BYTE_3(src[13])];\
dst[ 1] ^= groestl_T7[groestl_EXT_BYTE_3(src[22])];\
dst[ 2] ^= groestl_T7[groestl_EXT_BYTE_3(src[15])];\
dst[ 3] ^= groestl_T7[groestl_EXT_BYTE_3(src[24])];\
dst[ 4] ^= groestl_T7[groestl_EXT_BYTE_3(src[17])];\
dst[ 5] ^= groestl_T7[groestl_EXT_BYTE_3(src[26])];\
dst[ 6] ^= groestl_T7[groestl_EXT_BYTE_3(src[19])];\
dst[ 7] ^= groestl_T7[groestl_EXT_BYTE_3(src[28])];\
dst[ 8] ^= groestl_T7[groestl_EXT_BYTE_3(src[21])];\
dst[ 9] ^= groestl_T7[groestl_EXT_BYTE_3(src[30])];\
dst[10] ^= groestl_T7[groestl_EXT_BYTE_3(src[23])];\
dst[11] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 0])];\
dst[12] ^= groestl_T7[groestl_EXT_BYTE_3(src[25])];\
dst[13] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 2])];\
dst[14] ^= groestl_T7[groestl_EXT_BYTE_3(src[27])];\
dst[15] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 4])];\
dst[16] ^= groestl_T7[groestl_EXT_BYTE_3(src[29])];\
dst[17] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 6])];\
dst[18] ^= groestl_T7[groestl_EXT_BYTE_3(src[31])];\
dst[19] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 8])];\
dst[20] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 1])];\
dst[21] ^= groestl_T7[groestl_EXT_BYTE_3(src[10])];\
dst[22] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 3])];\
dst[23] ^= groestl_T7[groestl_EXT_BYTE_3(src[12])];\
dst[24] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 5])];\
dst[25] ^= groestl_T7[groestl_EXT_BYTE_3(src[14])];\
dst[26] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 7])];\
dst[27] ^= groestl_T7[groestl_EXT_BYTE_3(src[16])];\
dst[28] ^= groestl_T7[groestl_EXT_BYTE_3(src[ 9])];\
dst[29] ^= groestl_T7[groestl_EXT_BYTE_3(src[18])];\
dst[30] ^= groestl_T7[groestl_EXT_BYTE_3(src[11])];\
dst[31] ^= groestl_T7[groestl_EXT_BYTE_3(src[20])];

// global
const CONSTANT UINT32 groestl_T_init[256*8] =
{
0xa5f432c6UL,0x84976ff8UL,0x99b05eeeUL,0x8d8c7af6UL,0x0d17e8ffUL,0xbddc0ad6UL,0xb1c816deUL,0x54fc6d91UL,0x50f09060UL,0x03050702UL,0xa9e02eceUL,0x7d87d156UL,0x192bcce7UL,0x62a613b5UL,0xe6317c4dUL,0x9ab559ecUL,0x45cf408fUL,0x9dbca31fUL,0x40c04989UL,0x879268faUL,0x153fd0efUL,0xeb2694b2UL,0xc940ce8eUL,0x0b1de6fbUL,0xec2f6e41UL,0x67a91ab3UL,0xfd1c435fUL,0xea256045UL,0xbfdaf923UL,0xf7025153UL,0x96a145e4UL,0x5bed769bUL,0xc25d2875UL,0x1c24c5e1UL,0xaee9d43dUL,0x6abef24cUL,0x5aee826cUL,0x41c3bd7eUL,0x0206f3f5UL,0x4fd15283UL,0x5ce48c68UL,0xf4075651UL,0x345c8dd1UL,0x0818e1f9UL,0x93ae4ce2UL,0x73953eabUL,0x53f59762UL,0x3f416b2aUL,0x0c141c08UL,0x52f66395UL,0x65afe946UL,0x5ee27f9dUL,0x28784830UL,0xa1f8cf37UL,0x0f111b0aUL,0xb5c4eb2fUL,0x091b150eUL,0x365a7e24UL,0x9bb6ad1bUL,0x3d4798dfUL,0x266aa7cdUL,0x69bbf54eUL,0xcd4c337fUL,0x9fba50eaUL,0x1b2d3f12UL,0x9eb9a41dUL,0x749cc458UL,0x2e724634UL,0x2d774136UL,0xb2cd11dcUL,0xee299db4UL,0xfb164d5bUL,0xf601a5a4UL,0x4dd7a176UL,0x61a314b7UL,0xce49347dUL,0x7b8ddf52UL,0x3e429fddUL,0x7193cd5eUL,0x97a2b113UL,0xf504a2a6UL,0x68b801b9UL,0x00000000UL,0x2c74b5c1UL,0x60a0e040UL,0x1f21c2e3UL,0xc8433a79UL,0xed2c9ab6UL,0xbed90dd4UL,0x46ca478dUL,0xd9701767UL,0x4bddaf72UL,0xde79ed94UL,0xd467ff98UL,0xe82393b0UL,0x4ade5b85UL,0x6bbd06bbUL,0x2a7ebbc5UL,0xe5347b4fUL,0x163ad7edUL,0xc554d286UL,0xd762f89aUL,0x55ff9966UL,0x94a7b611UL,0xcf4ac08aUL,0x1030d9e9UL,0x060a0e04UL,0x819866feUL,0xf00baba0UL,0x44ccb478UL,0xbad5f025UL,0xe33e754bUL,0xf30eaca2UL,0xfe19445dUL,0xc05bdb80UL,0x8a858005UL,0xadecd33fUL,0xbcdffe21UL,0x48d8a870UL,0x040cfdf1UL,0xdf7a1963UL,0xc1582f77UL,0x759f30afUL,0x63a5e742UL,0x30507020UL,0x1a2ecbe5UL,0x0e12effdUL,0x6db708bfUL,0x4cd45581UL,0x143c2418UL,0x355f7926UL,0x2f71b2c3UL,0xe13886beUL,0xa2fdc835UL,0xcc4fc788UL,0x394b652eUL,0x57f96a93UL,0xf20d5855UL,0x829d61fcUL,0x47c9b37aUL,0xacef27c8UL,0xe73288baUL,0x2b7d4f32UL,0x95a442e6UL,0xa0fb3bc0UL,0x98b3aa19UL,0xd168f69eUL,0x7f8122a3UL,0x66aaee44UL,0x7e82d654UL,0xabe6dd3bUL,0x839e950bUL,0xca45c98cUL,0x297bbcc7UL,0xd36e056bUL,0x3c446c28UL,0x798b2ca7UL,0xe23d81bcUL,0x1d273116UL,0x769a37adUL,0x3b4d96dbUL,0x56fa9e64UL,0x4ed2a674UL,0x1e223614UL,0xdb76e492UL,0x0a1e120cUL,0x6cb4fc48UL,0xe4378fb8UL,0x5de7789fUL,0x6eb20fbdUL,0xef2a6943UL,0xa6f135c4UL,0xa8e3da39UL,0xa4f7c631UL,0x37598ad3UL,0x8b8674f2UL,0x325683d5UL,0x43c54e8bUL,0x59eb856eUL,0xb7c218daUL,0x8c8f8e01UL,0x64ac1db1UL,0xd26df19cUL,0xe03b7249UL,0xb4c71fd8UL,0xfa15b9acUL,0x0709faf3UL,0x256fa0cfUL,0xafea20caUL,0x8e897df4UL,0xe9206747UL,0x18283810UL,0xd5640b6fUL,0x888373f0UL,0x6fb1fb4aUL,0x7296ca5cUL,0x246c5438UL,0xf1085f57UL,0xc7522173UL,0x51f36497UL,0x2365aecbUL,0x7c8425a1UL,0x9cbf57e8UL,0x21635d3eUL,0xdd7cea96UL,0xdc7f1e61UL,0x86919c0dUL,0x85949b0fUL,0x90ab4be0UL,0x42c6ba7cUL,0xc4572671UL,0xaae529ccUL,0xd873e390UL,0x050f0906UL,0x0103f4f7UL,0x12362a1cUL,0xa3fe3cc2UL,0x5fe18b6aUL,0xf910beaeUL,0xd06b0269UL,0x91a8bf17UL,0x58e87199UL,0x2769533aUL,0xb9d0f727UL,0x384891d9UL,0x1335deebUL,0xb3cee52bUL,0x33557722UL,0xbbd604d2UL,0x709039a9UL,0x89808707UL,0xa7f2c133UL,0xb6c1ec2dUL,0x22665a3cUL,0x92adb815UL,0x2060a9c9UL,0x49db5c87UL,0xff1ab0aaUL,0x7888d850UL,0x7a8e2ba5UL,0x8f8a8903UL,0xf8134a59UL,0x809b9209UL,0x1739231aUL,0xda751065UL,0x315384d7UL,0xc651d584UL,0xb8d303d0UL,0xc35edc82UL,0xb0cbe229UL,0x7799c35aUL,0x11332d1eUL,0xcb463d7bUL,0xfc1fb7a8UL,0xd6610c6dUL,0x3a4e622cUL,
0xf432c6c6UL,0x976ff8f8UL,0xb05eeeeeUL,0x8c7af6f6UL,0x17e8ffffUL,0xdc0ad6d6UL,0xc816dedeUL,0xfc6d9191UL,0xf0906060UL,0x05070202UL,0xe02ececeUL,0x87d15656UL,0x2bcce7e7UL,0xa613b5b5UL,0x317c4d4dUL,0xb559ececUL,0xcf408f8fUL,0xbca31f1fUL,0xc0498989UL,0x9268fafaUL,0x3fd0efefUL,0x2694b2b2UL,0x40ce8e8eUL,0x1de6fbfbUL,0x2f6e4141UL,0xa91ab3b3UL,0x1c435f5fUL,0x25604545UL,0xdaf92323UL,0x02515353UL,0xa145e4e4UL,0xed769b9bUL,0x5d287575UL,0x24c5e1e1UL,0xe9d43d3dUL,0xbef24c4cUL,0xee826c6cUL,0xc3bd7e7eUL,0x06f3f5f5UL,0xd1528383UL,0xe48c6868UL,0x07565151UL,0x5c8dd1d1UL,0x18e1f9f9UL,0xae4ce2e2UL,0x953eababUL,0xf5976262UL,0x416b2a2aUL,0x141c0808UL,0xf6639595UL,0xafe94646UL,0xe27f9d9dUL,0x78483030UL,0xf8cf3737UL,0x111b0a0aUL,0xc4eb2f2fUL,0x1b150e0eUL,0x5a7e2424UL,0xb6ad1b1bUL,0x4798dfdfUL,0x6aa7cdcdUL,0xbbf54e4eUL,0x4c337f7fUL,0xba50eaeaUL,0x2d3f1212UL,0xb9a41d1dUL,0x9cc45858UL,0x72463434UL,0x77413636UL,0xcd11dcdcUL,0x299db4b4UL,0x164d5b5bUL,0x01a5a4a4UL,0xd7a17676UL,0xa314b7b7UL,0x49347d7dUL,0x8ddf5252UL,0x429fddddUL,0x93cd5e5eUL,0xa2b11313UL,0x04a2a6a6UL,0xb801b9b9UL,0x00000000UL,0x74b5c1c1UL,0xa0e04040UL,0x21c2e3e3UL,0x433a7979UL,0x2c9ab6b6UL,0xd90dd4d4UL,0xca478d8dUL,0x70176767UL,0xddaf7272UL,0x79ed9494UL,0x67ff9898UL,0x2393b0b0UL,0xde5b8585UL,0xbd06bbbbUL,0x7ebbc5c5UL,0x347b4f4fUL,0x3ad7ededUL,0x54d28686UL,0x62f89a9aUL,0xff996666UL,0xa7b61111UL,0x4ac08a8aUL,0x30d9e9e9UL,0x0a0e0404UL,0x9866fefeUL,0x0baba0a0UL,0xccb47878UL,0xd5f02525UL,0x3e754b4bUL,0x0eaca2a2UL,0x19445d5dUL,0x5bdb8080UL,0x85800505UL,0xecd33f3fUL,0xdffe2121UL,0xd8a87070UL,0x0cfdf1f1UL,0x7a196363UL,0x582f7777UL,0x9f30afafUL,0xa5e74242UL,0x50702020UL,0x2ecbe5e5UL,0x12effdfdUL,0xb708bfbfUL,0xd4558181UL,0x3c241818UL,0x5f792626UL,0x71b2c3c3UL,0x3886bebeUL,0xfdc83535UL,0x4fc78888UL,0x4b652e2eUL,0xf96a9393UL,0x0d585555UL,0x9d61fcfcUL,0xc9b37a7aUL,0xef27c8c8UL,0x3288babaUL,0x7d4f3232UL,0xa442e6e6UL,0xfb3bc0c0UL,0xb3aa1919UL,0x68f69e9eUL,0x8122a3a3UL,0xaaee4444UL,0x82d65454UL,0xe6dd3b3bUL,0x9e950b0bUL,0x45c98c8cUL,0x7bbcc7c7UL,0x6e056b6bUL,0x446c2828UL,0x8b2ca7a7UL,0x3d81bcbcUL,0x27311616UL,0x9a37adadUL,0x4d96dbdbUL,0xfa9e6464UL,0xd2a67474UL,0x22361414UL,0x76e49292UL,0x1e120c0cUL,0xb4fc4848UL,0x378fb8b8UL,0xe7789f9fUL,0xb20fbdbdUL,0x2a694343UL,0xf135c4c4UL,0xe3da3939UL,0xf7c63131UL,0x598ad3d3UL,0x8674f2f2UL,0x5683d5d5UL,0xc54e8b8bUL,0xeb856e6eUL,0xc218dadaUL,0x8f8e0101UL,0xac1db1b1UL,0x6df19c9cUL,0x3b724949UL,0xc71fd8d8UL,0x15b9acacUL,0x09faf3f3UL,0x6fa0cfcfUL,0xea20cacaUL,0x897df4f4UL,0x20674747UL,0x28381010UL,0x640b6f6fUL,0x8373f0f0UL,0xb1fb4a4aUL,0x96ca5c5cUL,0x6c543838UL,0x085f5757UL,0x52217373UL,0xf3649797UL,0x65aecbcbUL,0x8425a1a1UL,0xbf57e8e8UL,0x635d3e3eUL,0x7cea9696UL,0x7f1e6161UL,0x919c0d0dUL,0x949b0f0fUL,0xab4be0e0UL,0xc6ba7c7cUL,0x57267171UL,0xe529ccccUL,0x73e39090UL,0x0f090606UL,0x03f4f7f7UL,0x362a1c1cUL,0xfe3cc2c2UL,0xe18b6a6aUL,0x10beaeaeUL,0x6b026969UL,0xa8bf1717UL,0xe8719999UL,0x69533a3aUL,0xd0f72727UL,0x4891d9d9UL,0x35deebebUL,0xcee52b2bUL,0x55772222UL,0xd604d2d2UL,0x9039a9a9UL,0x80870707UL,0xf2c13333UL,0xc1ec2d2dUL,0x665a3c3cUL,0xadb81515UL,0x60a9c9c9UL,0xdb5c8787UL,0x1ab0aaaaUL,0x88d85050UL,0x8e2ba5a5UL,0x8a890303UL,0x134a5959UL,0x9b920909UL,0x39231a1aUL,0x75106565UL,0x5384d7d7UL,0x51d58484UL,0xd303d0d0UL,0x5edc8282UL,0xcbe22929UL,0x99c35a5aUL,0x332d1e1eUL,0x463d7b7bUL,0x1fb7a8a8UL,0x610c6d6dUL,0x4e622c2cUL,
0x32c6c6a5UL,0x6ff8f884UL,0x5eeeee99UL,0x7af6f68dUL,0xe8ffff0dUL,0x0ad6d6bdUL,0x16dedeb1UL,0x6d919154UL,0x90606050UL,0x07020203UL,0x2ececea9UL,0xd156567dUL,0xcce7e719UL,0x13b5b562UL,0x7c4d4de6UL,0x59ecec9aUL,0x408f8f45UL,0xa31f1f9dUL,0x49898940UL,0x68fafa87UL,0xd0efef15UL,0x94b2b2ebUL,0xce8e8ec9UL,0xe6fbfb0bUL,0x6e4141ecUL,0x1ab3b367UL,0x435f5ffdUL,0x604545eaUL,0xf92323bfUL,0x515353f7UL,0x45e4e496UL,0x769b9b5bUL,0x287575c2UL,0xc5e1e11cUL,0xd43d3daeUL,0xf24c4c6aUL,0x826c6c5aUL,0xbd7e7e41UL,0xf3f5f502UL,0x5283834fUL,0x8c68685cUL,0x565151f4UL,0x8dd1d134UL,0xe1f9f908UL,0x4ce2e293UL,0x3eabab73UL,0x97626253UL,0x6b2a2a3fUL,0x1c08080cUL,0x63959552UL,0xe9464665UL,0x7f9d9d5eUL,0x48303028UL,0xcf3737a1UL,0x1b0a0a0fUL,0xeb2f2fb5UL,0x150e0e09UL,0x7e242436UL,0xad1b1b9bUL,0x98dfdf3dUL,0xa7cdcd26UL,0xf54e4e69UL,0x337f7fcdUL,0x50eaea9fUL,0x3f12121bUL,0xa41d1d9eUL,0xc4585874UL,0x4634342eUL,0x4136362dUL,0x11dcdcb2UL,0x9db4b4eeUL,0x4d5b5bfbUL,0xa5a4a4f6UL,0xa176764dUL,0x14b7b761UL,0x347d7dceUL,0xdf52527bUL,0x9fdddd3eUL,0xcd5e5e71UL,0xb1131397UL,0xa2a6a6f5UL,0x01b9b968UL,0x00000000UL,0xb5c1c12cUL,0xe0404060UL,0xc2e3e31fUL,0x3a7979c8UL,0x9ab6b6edUL,0x0dd4d4beUL,0x478d8d46UL,0x176767d9UL,0xaf72724bUL,0xed9494deUL,0xff9898d4UL,0x93b0b0e8UL,0x5b85854aUL,0x06bbbb6bUL,0xbbc5c52aUL,0x7b4f4fe5UL,0xd7eded16UL,0xd28686c5UL,0xf89a9ad7UL,0x99666655UL,0xb6111194UL,0xc08a8acfUL,0xd9e9e910UL,0x0e040406UL,0x66fefe81UL,0xaba0a0f0UL,0xb4787844UL,0xf02525baUL,0x754b4be3UL,0xaca2a2f3UL,0x445d5dfeUL,0xdb8080c0UL,0x8005058aUL,0xd33f3fadUL,0xfe2121bcUL,0xa8707048UL,0xfdf1f104UL,0x196363dfUL,0x2f7777c1UL,0x30afaf75UL,0xe7424263UL,0x70202030UL,0xcbe5e51aUL,0xeffdfd0eUL,0x08bfbf6dUL,0x5581814cUL,0x24181814UL,0x79262635UL,0xb2c3c32fUL,0x86bebee1UL,0xc83535a2UL,0xc78888ccUL,0x652e2e39UL,0x6a939357UL,0x585555f2UL,0x61fcfc82UL,0xb37a7a47UL,0x27c8c8acUL,0x88babae7UL,0x4f32322bUL,0x42e6e695UL,0x3bc0c0a0UL,0xaa191998UL,0xf69e9ed1UL,0x22a3a37fUL,0xee444466UL,0xd654547eUL,0xdd3b3babUL,0x950b0b83UL,0xc98c8ccaUL,0xbcc7c729UL,0x056b6bd3UL,0x6c28283cUL,0x2ca7a779UL,0x81bcbce2UL,0x3116161dUL,0x37adad76UL,0x96dbdb3bUL,0x9e646456UL,0xa674744eUL,0x3614141eUL,0xe49292dbUL,0x120c0c0aUL,0xfc48486cUL,0x8fb8b8e4UL,0x789f9f5dUL,0x0fbdbd6eUL,0x694343efUL,0x35c4c4a6UL,0xda3939a8UL,0xc63131a4UL,0x8ad3d337UL,0x74f2f28bUL,0x83d5d532UL,0x4e8b8b43UL,0x856e6e59UL,0x18dadab7UL,0x8e01018cUL,0x1db1b164UL,0xf19c9cd2UL,0x724949e0UL,0x1fd8d8b4UL,0xb9acacfaUL,0xfaf3f307UL,0xa0cfcf25UL,0x20cacaafUL,0x7df4f48eUL,0x674747e9UL,0x38101018UL,0x0b6f6fd5UL,0x73f0f088UL,0xfb4a4a6fUL,0xca5c5c72UL,0x54383824UL,0x5f5757f1UL,0x217373c7UL,0x64979751UL,0xaecbcb23UL,0x25a1a17cUL,0x57e8e89cUL,0x5d3e3e21UL,0xea9696ddUL,0x1e6161dcUL,0x9c0d0d86UL,0x9b0f0f85UL,0x4be0e090UL,0xba7c7c42UL,0x267171c4UL,0x29ccccaaUL,0xe39090d8UL,0x09060605UL,0xf4f7f701UL,0x2a1c1c12UL,0x3cc2c2a3UL,0x8b6a6a5fUL,0xbeaeaef9UL,0x026969d0UL,0xbf171791UL,0x71999958UL,0x533a3a27UL,0xf72727b9UL,0x91d9d938UL,0xdeebeb13UL,0xe52b2bb3UL,0x77222233UL,0x04d2d2bbUL,0x39a9a970UL,0x87070789UL,0xc13333a7UL,0xec2d2db6UL,0x5a3c3c22UL,0xb8151592UL,0xa9c9c920UL,0x5c878749UL,0xb0aaaaffUL,0xd8505078UL,0x2ba5a57aUL,0x8903038fUL,0x4a5959f8UL,0x92090980UL,0x231a1a17UL,0x106565daUL,0x84d7d731UL,0xd58484c6UL,0x03d0d0b8UL,0xdc8282c3UL,0xe22929b0UL,0xc35a5a77UL,0x2d1e1e11UL,0x3d7b7bcbUL,0xb7a8a8fcUL,0x0c6d6dd6UL,0x622c2c3aUL,
0xc6c6a597UL,0xf8f884ebUL,0xeeee99c7UL,0xf6f68df7UL,0xffff0de5UL,0xd6d6bdb7UL,0xdedeb1a7UL,0x91915439UL,0x606050c0UL,0x02020304UL,0xcecea987UL,0x56567dacUL,0xe7e719d5UL,0xb5b56271UL,0x4d4de69aUL,0xecec9ac3UL,0x8f8f4505UL,0x1f1f9d3eUL,0x89894009UL,0xfafa87efUL,0xefef15c5UL,0xb2b2eb7fUL,0x8e8ec907UL,0xfbfb0bedUL,0x4141ec82UL,0xb3b3677dUL,0x5f5ffdbeUL,0x4545ea8aUL,0x2323bf46UL,0x5353f7a6UL,0xe4e496d3UL,0x9b9b5b2dUL,0x7575c2eaUL,0xe1e11cd9UL,0x3d3dae7aUL,0x4c4c6a98UL,0x6c6c5ad8UL,0x7e7e41fcUL,0xf5f502f1UL,0x83834f1dUL,0x68685cd0UL,0x5151f4a2UL,0xd1d134b9UL,0xf9f908e9UL,0xe2e293dfUL,0xabab734dUL,0x626253c4UL,0x2a2a3f54UL,0x08080c10UL,0x95955231UL,0x4646658cUL,0x9d9d5e21UL,0x30302860UL,0x3737a16eUL,0x0a0a0f14UL,0x2f2fb55eUL,0x0e0e091cUL,0x24243648UL,0x1b1b9b36UL,0xdfdf3da5UL,0xcdcd2681UL,0x4e4e699cUL,0x7f7fcdfeUL,0xeaea9fcfUL,0x12121b24UL,0x1d1d9e3aUL,0x585874b0UL,0x34342e68UL,0x36362d6cUL,0xdcdcb2a3UL,0xb4b4ee73UL,0x5b5bfbb6UL,0xa4a4f653UL,0x76764decUL,0xb7b76175UL,0x7d7dcefaUL,0x52527ba4UL,0xdddd3ea1UL,0x5e5e71bcUL,0x13139726UL,0xa6a6f557UL,0xb9b96869UL,0x00000000UL,0xc1c12c99UL,0x40406080UL,0xe3e31fddUL,0x7979c8f2UL,0xb6b6ed77UL,0xd4d4beb3UL,0x8d8d4601UL,0x6767d9ceUL,0x72724be4UL,0x9494de33UL,0x9898d42bUL,0xb0b0e87bUL,0x85854a11UL,0xbbbb6b6dUL,0xc5c52a91UL,0x4f4fe59eUL,0xeded16c1UL,0x8686c517UL,0x9a9ad72fUL,0x666655ccUL,0x11119422UL,0x8a8acf0fUL,0xe9e910c9UL,0x04040608UL,0xfefe81e7UL,0xa0a0f05bUL,0x787844f0UL,0x2525ba4aUL,0x4b4be396UL,0xa2a2f35fUL,0x5d5dfebaUL,0x8080c01bUL,0x05058a0aUL,0x3f3fad7eUL,0x2121bc42UL,0x707048e0UL,0xf1f104f9UL,0x6363dfc6UL,0x7777c1eeUL,0xafaf7545UL,0x42426384UL,0x20203040UL,0xe5e51ad1UL,0xfdfd0ee1UL,0xbfbf6d65UL,0x81814c19UL,0x18181430UL,0x2626354cUL,0xc3c32f9dUL,0xbebee167UL,0x3535a26aUL,0x8888cc0bUL,0x2e2e395cUL,0x9393573dUL,0x5555f2aaUL,0xfcfc82e3UL,0x7a7a47f4UL,0xc8c8ac8bUL,0xbabae76fUL,0x32322b64UL,0xe6e695d7UL,0xc0c0a09bUL,0x19199832UL,0x9e9ed127UL,0xa3a37f5dUL,0x44446688UL,0x54547ea8UL,0x3b3bab76UL,0x0b0b8316UL,0x8c8cca03UL,0xc7c72995UL,0x6b6bd3d6UL,0x28283c50UL,0xa7a77955UL,0xbcbce263UL,0x16161d2cUL,0xadad7641UL,0xdbdb3badUL,0x646456c8UL,0x74744ee8UL,0x14141e28UL,0x9292db3fUL,0x0c0c0a18UL,0x48486c90UL,0xb8b8e46bUL,0x9f9f5d25UL,0xbdbd6e61UL,0x4343ef86UL,0xc4c4a693UL,0x3939a872UL,0x3131a462UL,0xd3d337bdUL,0xf2f28bffUL,0xd5d532b1UL,0x8b8b430dUL,0x6e6e59dcUL,0xdadab7afUL,0x01018c02UL,0xb1b16479UL,0x9c9cd223UL,0x4949e092UL,0xd8d8b4abUL,0xacacfa43UL,0xf3f307fdUL,0xcfcf2585UL,0xcacaaf8fUL,0xf4f48ef3UL,0x4747e98eUL,0x10101820UL,0x6f6fd5deUL,0xf0f088fbUL,0x4a4a6f94UL,0x5c5c72b8UL,0x38382470UL,0x5757f1aeUL,0x7373c7e6UL,0x97975135UL,0xcbcb238dUL,0xa1a17c59UL,0xe8e89ccbUL,0x3e3e217cUL,0x9696dd37UL,0x6161dcc2UL,0x0d0d861aUL,0x0f0f851eUL,0xe0e090dbUL,0x7c7c42f8UL,0x7171c4e2UL,0xccccaa83UL,0x9090d83bUL,0x0606050cUL,0xf7f701f5UL,0x1c1c1238UL,0xc2c2a39fUL,0x6a6a5fd4UL,0xaeaef947UL,0x6969d0d2UL,0x1717912eUL,0x99995829UL,0x3a3a2774UL,0x2727b94eUL,0xd9d938a9UL,0xebeb13cdUL,0x2b2bb356UL,0x22223344UL,0xd2d2bbbfUL,0xa9a97049UL,0x0707890eUL,0x3333a766UL,0x2d2db65aUL,0x3c3c2278UL,0x1515922aUL,0xc9c92089UL,0x87874915UL,0xaaaaff4fUL,0x505078a0UL,0xa5a57a51UL,0x03038f06UL,0x5959f8b2UL,0x09098012UL,0x1a1a1734UL,0x6565dacaUL,0xd7d731b5UL,0x8484c613UL,0xd0d0b8bbUL,0x8282c31fUL,0x2929b052UL,0x5a5a77b4UL,0x1e1e113cUL,0x7b7bcbf6UL,0xa8a8fc4bUL,0x6d6dd6daUL,0x2c2c3a58UL,
0xc6a597f4UL,0xf884eb97UL,0xee99c7b0UL,0xf68df78cUL,0xff0de517UL,0xd6bdb7dcUL,0xdeb1a7c8UL,0x915439fcUL,0x6050c0f0UL,0x02030405UL,0xcea987e0UL,0x567dac87UL,0xe719d52bUL,0xb56271a6UL,0x4de69a31UL,0xec9ac3b5UL,0x8f4505cfUL,0x1f9d3ebcUL,0x894009c0UL,0xfa87ef92UL,0xef15c53fUL,0xb2eb7f26UL,0x8ec90740UL,0xfb0bed1dUL,0x41ec822fUL,0xb3677da9UL,0x5ffdbe1cUL,0x45ea8a25UL,0x23bf46daUL,0x53f7a602UL,0xe496d3a1UL,0x9b5b2dedUL,0x75c2ea5dUL,0xe11cd924UL,0x3dae7ae9UL,0x4c6a98beUL,0x6c5ad8eeUL,0x7e41fcc3UL,0xf502f106UL,0x834f1dd1UL,0x685cd0e4UL,0x51f4a207UL,0xd134b95cUL,0xf908e918UL,0xe293dfaeUL,0xab734d95UL,0x6253c4f5UL,0x2a3f5441UL,0x080c1014UL,0x955231f6UL,0x46658cafUL,0x9d5e21e2UL,0x30286078UL,0x37a16ef8UL,0x0a0f1411UL,0x2fb55ec4UL,0x0e091c1bUL,0x2436485aUL,0x1b9b36b6UL,0xdf3da547UL,0xcd26816aUL,0x4e699cbbUL,0x7fcdfe4cUL,0xea9fcfbaUL,0x121b242dUL,0x1d9e3ab9UL,0x5874b09cUL,0x342e6872UL,0x362d6c77UL,0xdcb2a3cdUL,0xb4ee7329UL,0x5bfbb616UL,0xa4f65301UL,0x764decd7UL,0xb76175a3UL,0x7dcefa49UL,0x527ba48dUL,0xdd3ea142UL,0x5e71bc93UL,0x139726a2UL,0xa6f55704UL,0xb96869b8UL,0x00000000UL,0xc12c9974UL,0x406080a0UL,0xe31fdd21UL,0x79c8f243UL,0xb6ed772cUL,0xd4beb3d9UL,0x8d4601caUL,0x67d9ce70UL,0x724be4ddUL,0x94de3379UL,0x98d42b67UL,0xb0e87b23UL,0x854a11deUL,0xbb6b6dbdUL,0xc52a917eUL,0x4fe59e34UL,0xed16c13aUL,0x86c51754UL,0x9ad72f62UL,0x6655ccffUL,0x119422a7UL,0x8acf0f4aUL,0xe910c930UL,0x0406080aUL,0xfe81e798UL,0xa0f05b0bUL,0x7844f0ccUL,0x25ba4ad5UL,0x4be3963eUL,0xa2f35f0eUL,0x5dfeba19UL,0x80c01b5bUL,0x058a0a85UL,0x3fad7eecUL,0x21bc42dfUL,0x7048e0d8UL,0xf104f90cUL,0x63dfc67aUL,0x77c1ee58UL,0xaf75459fUL,0x426384a5UL,0x20304050UL,0xe51ad12eUL,0xfd0ee112UL,0xbf6d65b7UL,0x814c19d4UL,0x1814303cUL,0x26354c5fUL,0xc32f9d71UL,0xbee16738UL,0x35a26afdUL,0x88cc0b4fUL,0x2e395c4bUL,0x93573df9UL,0x55f2aa0dUL,0xfc82e39dUL,0x7a47f4c9UL,0xc8ac8befUL,0xbae76f32UL,0x322b647dUL,0xe695d7a4UL,0xc0a09bfbUL,0x199832b3UL,0x9ed12768UL,0xa37f5d81UL,0x446688aaUL,0x547ea882UL,0x3bab76e6UL,0x0b83169eUL,0x8cca0345UL,0xc729957bUL,0x6bd3d66eUL,0x283c5044UL,0xa779558bUL,0xbce2633dUL,0x161d2c27UL,0xad76419aUL,0xdb3bad4dUL,0x6456c8faUL,0x744ee8d2UL,0x141e2822UL,0x92db3f76UL,0x0c0a181eUL,0x486c90b4UL,0xb8e46b37UL,0x9f5d25e7UL,0xbd6e61b2UL,0x43ef862aUL,0xc4a693f1UL,0x39a872e3UL,0x31a462f7UL,0xd337bd59UL,0xf28bff86UL,0xd532b156UL,0x8b430dc5UL,0x6e59dcebUL,0xdab7afc2UL,0x018c028fUL,0xb16479acUL,0x9cd2236dUL,0x49e0923bUL,0xd8b4abc7UL,0xacfa4315UL,0xf307fd09UL,0xcf25856fUL,0xcaaf8feaUL,0xf48ef389UL,0x47e98e20UL,0x10182028UL,0x6fd5de64UL,0xf088fb83UL,0x4a6f94b1UL,0x5c72b896UL,0x3824706cUL,0x57f1ae08UL,0x73c7e652UL,0x975135f3UL,0xcb238d65UL,0xa17c5984UL,0xe89ccbbfUL,0x3e217c63UL,0x96dd377cUL,0x61dcc27fUL,0x0d861a91UL,0x0f851e94UL,0xe090dbabUL,0x7c42f8c6UL,0x71c4e257UL,0xccaa83e5UL,0x90d83b73UL,0x06050c0fUL,0xf701f503UL,0x1c123836UL,0xc2a39ffeUL,0x6a5fd4e1UL,0xaef94710UL,0x69d0d26bUL,0x17912ea8UL,0x995829e8UL,0x3a277469UL,0x27b94ed0UL,0xd938a948UL,0xeb13cd35UL,0x2bb356ceUL,0x22334455UL,0xd2bbbfd6UL,0xa9704990UL,0x07890e80UL,0x33a766f2UL,0x2db65ac1UL,0x3c227866UL,0x15922aadUL,0xc9208960UL,0x874915dbUL,0xaaff4f1aUL,0x5078a088UL,0xa57a518eUL,0x038f068aUL,0x59f8b213UL,0x0980129bUL,0x1a173439UL,0x65daca75UL,0xd731b553UL,0x84c61351UL,0xd0b8bbd3UL,0x82c31f5eUL,0x29b052cbUL,0x5a77b499UL,0x1e113c33UL,0x7bcbf646UL,0xa8fc4b1fUL,0x6dd6da61UL,0x2c3a584eUL,
0xa597f4a5UL,0x84eb9784UL,0x99c7b099UL,0x8df78c8dUL,0x0de5170dUL,0xbdb7dcbdUL,0xb1a7c8b1UL,0x5439fc54UL,0x50c0f050UL,0x03040503UL,0xa987e0a9UL,0x7dac877dUL,0x19d52b19UL,0x6271a662UL,0xe69a31e6UL,0x9ac3b59aUL,0x4505cf45UL,0x9d3ebc9dUL,0x4009c040UL,0x87ef9287UL,0x15c53f15UL,0xeb7f26ebUL,0xc90740c9UL,0x0bed1d0bUL,0xec822fecUL,0x677da967UL,0xfdbe1cfdUL,0xea8a25eaUL,0xbf46dabfUL,0xf7a602f7UL,0x96d3a196UL,0x5b2ded5bUL,0xc2ea5dc2UL,0x1cd9241cUL,0xae7ae9aeUL,0x6a98be6aUL,0x5ad8ee5aUL,0x41fcc341UL,0x02f10602UL,0x4f1dd14fUL,0x5cd0e45cUL,0xf4a207f4UL,0x34b95c34UL,0x08e91808UL,0x93dfae93UL,0x734d9573UL,0x53c4f553UL,0x3f54413fUL,0x0c10140cUL,0x5231f652UL,0x658caf65UL,0x5e21e25eUL,0x28607828UL,0xa16ef8a1UL,0x0f14110fUL,0xb55ec4b5UL,0x091c1b09UL,0x36485a36UL,0x9b36b69bUL,0x3da5473dUL,0x26816a26UL,0x699cbb69UL,0xcdfe4ccdUL,0x9fcfba9fUL,0x1b242d1bUL,0x9e3ab99eUL,0x74b09c74UL,0x2e68722eUL,0x2d6c772dUL,0xb2a3cdb2UL,0xee7329eeUL,0xfbb616fbUL,0xf65301f6UL,0x4decd74dUL,0x6175a361UL,0xcefa49ceUL,0x7ba48d7bUL,0x3ea1423eUL,0x71bc9371UL,0x9726a297UL,0xf55704f5UL,0x6869b868UL,0x00000000UL,0x2c99742cUL,0x6080a060UL,0x1fdd211fUL,0xc8f243c8UL,0xed772cedUL,0xbeb3d9beUL,0x4601ca46UL,0xd9ce70d9UL,0x4be4dd4bUL,0xde3379deUL,0xd42b67d4UL,0xe87b23e8UL,0x4a11de4aUL,0x6b6dbd6bUL,0x2a917e2aUL,0xe59e34e5UL,0x16c13a16UL,0xc51754c5UL,0xd72f62d7UL,0x55ccff55UL,0x9422a794UL,0xcf0f4acfUL,0x10c93010UL,0x06080a06UL,0x81e79881UL,0xf05b0bf0UL,0x44f0cc44UL,0xba4ad5baUL,0xe3963ee3UL,0xf35f0ef3UL,0xfeba19feUL,0xc01b5bc0UL,0x8a0a858aUL,0xad7eecadUL,0xbc42dfbcUL,0x48e0d848UL,0x04f90c04UL,0xdfc67adfUL,0xc1ee58c1UL,0x75459f75UL,0x6384a563UL,0x30405030UL,0x1ad12e1aUL,0x0ee1120eUL,0x6d65b76dUL,0x4c19d44cUL,0x14303c14UL,0x354c5f35UL,0x2f9d712fUL,0xe16738e1UL,0xa26afda2UL,0xcc0b4fccUL,0x395c4b39UL,0x573df957UL,0xf2aa0df2UL,0x82e39d82UL,0x47f4c947UL,0xac8befacUL,0xe76f32e7UL,0x2b647d2bUL,0x95d7a495UL,0xa09bfba0UL,0x9832b398UL,0xd12768d1UL,0x7f5d817fUL,0x6688aa66UL,0x7ea8827eUL,0xab76e6abUL,0x83169e83UL,0xca0345caUL,0x29957b29UL,0xd3d66ed3UL,0x3c50443cUL,0x79558b79UL,0xe2633de2UL,0x1d2c271dUL,0x76419a76UL,0x3bad4d3bUL,0x56c8fa56UL,0x4ee8d24eUL,0x1e28221eUL,0xdb3f76dbUL,0x0a181e0aUL,0x6c90b46cUL,0xe46b37e4UL,0x5d25e75dUL,0x6e61b26eUL,0xef862aefUL,0xa693f1a6UL,0xa872e3a8UL,0xa462f7a4UL,0x37bd5937UL,0x8bff868bUL,0x32b15632UL,0x430dc543UL,0x59dceb59UL,0xb7afc2b7UL,0x8c028f8cUL,0x6479ac64UL,0xd2236dd2UL,0xe0923be0UL,0xb4abc7b4UL,0xfa4315faUL,0x07fd0907UL,0x25856f25UL,0xaf8feaafUL,0x8ef3898eUL,0xe98e20e9UL,0x18202818UL,0xd5de64d5UL,0x88fb8388UL,0x6f94b16fUL,0x72b89672UL,0x24706c24UL,0xf1ae08f1UL,0xc7e652c7UL,0x5135f351UL,0x238d6523UL,0x7c59847cUL,0x9ccbbf9cUL,0x217c6321UL,0xdd377cddUL,0xdcc27fdcUL,0x861a9186UL,0x851e9485UL,0x90dbab90UL,0x42f8c642UL,0xc4e257c4UL,0xaa83e5aaUL,0xd83b73d8UL,0x050c0f05UL,0x01f50301UL,0x12383612UL,0xa39ffea3UL,0x5fd4e15fUL,0xf94710f9UL,0xd0d26bd0UL,0x912ea891UL,0x5829e858UL,0x27746927UL,0xb94ed0b9UL,0x38a94838UL,0x13cd3513UL,0xb356ceb3UL,0x33445533UL,0xbbbfd6bbUL,0x70499070UL,0x890e8089UL,0xa766f2a7UL,0xb65ac1b6UL,0x22786622UL,0x922aad92UL,0x20896020UL,0x4915db49UL,0xff4f1affUL,0x78a08878UL,0x7a518e7aUL,0x8f068a8fUL,0xf8b213f8UL,0x80129b80UL,0x17343917UL,0xdaca75daUL,0x31b55331UL,0xc61351c6UL,0xb8bbd3b8UL,0xc31f5ec3UL,0xb052cbb0UL,0x77b49977UL,0x113c3311UL,0xcbf646cbUL,0xfc4b1ffcUL,0xd6da61d6UL,0x3a584e3aUL,
0x97f4a5f4UL,0xeb978497UL,0xc7b099b0UL,0xf78c8d8cUL,0xe5170d17UL,0xb7dcbddcUL,0xa7c8b1c8UL,0x39fc54fcUL,0xc0f050f0UL,0x04050305UL,0x87e0a9e0UL,0xac877d87UL,0xd52b192bUL,0x71a662a6UL,0x9a31e631UL,0xc3b59ab5UL,0x05cf45cfUL,0x3ebc9dbcUL,0x09c040c0UL,0xef928792UL,0xc53f153fUL,0x7f26eb26UL,0x0740c940UL,0xed1d0b1dUL,0x822fec2fUL,0x7da967a9UL,0xbe1cfd1cUL,0x8a25ea25UL,0x46dabfdaUL,0xa602f702UL,0xd3a196a1UL,0x2ded5bedUL,0xea5dc25dUL,0xd9241c24UL,0x7ae9aee9UL,0x98be6abeUL,0xd8ee5aeeUL,0xfcc341c3UL,0xf1060206UL,0x1dd14fd1UL,0xd0e45ce4UL,0xa207f407UL,0xb95c345cUL,0xe9180818UL,0xdfae93aeUL,0x4d957395UL,0xc4f553f5UL,0x54413f41UL,0x10140c14UL,0x31f652f6UL,0x8caf65afUL,0x21e25ee2UL,0x60782878UL,0x6ef8a1f8UL,0x14110f11UL,0x5ec4b5c4UL,0x1c1b091bUL,0x485a365aUL,0x36b69bb6UL,0xa5473d47UL,0x816a266aUL,0x9cbb69bbUL,0xfe4ccd4cUL,0xcfba9fbaUL,0x242d1b2dUL,0x3ab99eb9UL,0xb09c749cUL,0x68722e72UL,0x6c772d77UL,0xa3cdb2cdUL,0x7329ee29UL,0xb616fb16UL,0x5301f601UL,0xecd74dd7UL,0x75a361a3UL,0xfa49ce49UL,0xa48d7b8dUL,0xa1423e42UL,0xbc937193UL,0x26a297a2UL,0x5704f504UL,0x69b868b8UL,0x00000000UL,0x99742c74UL,0x80a060a0UL,0xdd211f21UL,0xf243c843UL,0x772ced2cUL,0xb3d9bed9UL,0x01ca46caUL,0xce70d970UL,0xe4dd4bddUL,0x3379de79UL,0x2b67d467UL,0x7b23e823UL,0x11de4adeUL,0x6dbd6bbdUL,0x917e2a7eUL,0x9e34e534UL,0xc13a163aUL,0x1754c554UL,0x2f62d762UL,0xccff55ffUL,0x22a794a7UL,0x0f4acf4aUL,0xc9301030UL,0x080a060aUL,0xe7988198UL,0x5b0bf00bUL,0xf0cc44ccUL,0x4ad5bad5UL,0x963ee33eUL,0x5f0ef30eUL,0xba19fe19UL,0x1b5bc05bUL,0x0a858a85UL,0x7eecadecUL,0x42dfbcdfUL,0xe0d848d8UL,0xf90c040cUL,0xc67adf7aUL,0xee58c158UL,0x459f759fUL,0x84a563a5UL,0x40503050UL,0xd12e1a2eUL,0xe1120e12UL,0x65b76db7UL,0x19d44cd4UL,0x303c143cUL,0x4c5f355fUL,0x9d712f71UL,0x6738e138UL,0x6afda2fdUL,0x0b4fcc4fUL,0x5c4b394bUL,0x3df957f9UL,0xaa0df20dUL,0xe39d829dUL,0xf4c947c9UL,0x8befacefUL,0x6f32e732UL,0x647d2b7dUL,0xd7a495a4UL,0x9bfba0fbUL,0x32b398b3UL,0x2768d168UL,0x5d817f81UL,0x88aa66aaUL,0xa8827e82UL,0x76e6abe6UL,0x169e839eUL,0x0345ca45UL,0x957b297bUL,0xd66ed36eUL,0x50443c44UL,0x558b798bUL,0x633de23dUL,0x2c271d27UL,0x419a769aUL,0xad4d3b4dUL,0xc8fa56faUL,0xe8d24ed2UL,0x28221e22UL,0x3f76db76UL,0x181e0a1eUL,0x90b46cb4UL,0x6b37e437UL,0x25e75de7UL,0x61b26eb2UL,0x862aef2aUL,0x93f1a6f1UL,0x72e3a8e3UL,0x62f7a4f7UL,0xbd593759UL,0xff868b86UL,0xb1563256UL,0x0dc543c5UL,0xdceb59ebUL,0xafc2b7c2UL,0x028f8c8fUL,0x79ac64acUL,0x236dd26dUL,0x923be03bUL,0xabc7b4c7UL,0x4315fa15UL,0xfd090709UL,0x856f256fUL,0x8feaafeaUL,0xf3898e89UL,0x8e20e920UL,0x20281828UL,0xde64d564UL,0xfb838883UL,0x94b16fb1UL,0xb8967296UL,0x706c246cUL,0xae08f108UL,0xe652c752UL,0x35f351f3UL,0x8d652365UL,0x59847c84UL,0xcbbf9cbfUL,0x7c632163UL,0x377cdd7cUL,0xc27fdc7fUL,0x1a918691UL,0x1e948594UL,0xdbab90abUL,0xf8c642c6UL,0xe257c457UL,0x83e5aae5UL,0x3b73d873UL,0x0c0f050fUL,0xf5030103UL,0x38361236UL,0x9ffea3feUL,0xd4e15fe1UL,0x4710f910UL,0xd26bd06bUL,0x2ea891a8UL,0x29e858e8UL,0x74692769UL,0x4ed0b9d0UL,0xa9483848UL,0xcd351335UL,0x56ceb3ceUL,0x44553355UL,0xbfd6bbd6UL,0x49907090UL,0x0e808980UL,0x66f2a7f2UL,0x5ac1b6c1UL,0x78662266UL,0x2aad92adUL,0x89602060UL,0x15db49dbUL,0x4f1aff1aUL,0xa0887888UL,0x518e7a8eUL,0x068a8f8aUL,0xb213f813UL,0x129b809bUL,0x34391739UL,0xca75da75UL,0xb5533153UL,0x1351c651UL,0xbbd3b8d3UL,0x1f5ec35eUL,0x52cbb0cbUL,0xb4997799UL,0x3c331133UL,0xf646cb46UL,0x4b1ffc1fUL,0xda61d661UL,0x584e3a4eUL,
0xf4a5f432UL,0x9784976fUL,0xb099b05eUL,0x8c8d8c7aUL,0x170d17e8UL,0xdcbddc0aUL,0xc8b1c816UL,0xfc54fc6dUL,0xf050f090UL,0x05030507UL,0xe0a9e02eUL,0x877d87d1UL,0x2b192bccUL,0xa662a613UL,0x31e6317cUL,0xb59ab559UL,0xcf45cf40UL,0xbc9dbca3UL,0xc040c049UL,0x92879268UL,0x3f153fd0UL,0x26eb2694UL,0x40c940ceUL,0x1d0b1de6UL,0x2fec2f6eUL,0xa967a91aUL,0x1cfd1c43UL,0x25ea2560UL,0xdabfdaf9UL,0x02f70251UL,0xa196a145UL,0xed5bed76UL,0x5dc25d28UL,0x241c24c5UL,0xe9aee9d4UL,0xbe6abef2UL,0xee5aee82UL,0xc341c3bdUL,0x060206f3UL,0xd14fd152UL,0xe45ce48cUL,0x07f40756UL,0x5c345c8dUL,0x180818e1UL,0xae93ae4cUL,0x9573953eUL,0xf553f597UL,0x413f416bUL,0x140c141cUL,0xf652f663UL,0xaf65afe9UL,0xe25ee27fUL,0x78287848UL,0xf8a1f8cfUL,0x110f111bUL,0xc4b5c4ebUL,0x1b091b15UL,0x5a365a7eUL,0xb69bb6adUL,0x473d4798UL,0x6a266aa7UL,0xbb69bbf5UL,0x4ccd4c33UL,0xba9fba50UL,0x2d1b2d3fUL,0xb99eb9a4UL,0x9c749cc4UL,0x722e7246UL,0x772d7741UL,0xcdb2cd11UL,0x29ee299dUL,0x16fb164dUL,0x01f601a5UL,0xd74dd7a1UL,0xa361a314UL,0x49ce4934UL,0x8d7b8ddfUL,0x423e429fUL,0x937193cdUL,0xa297a2b1UL,0x04f504a2UL,0xb868b801UL,0x00000000UL,0x742c74b5UL,0xa060a0e0UL,0x211f21c2UL,0x43c8433aUL,0x2ced2c9aUL,0xd9bed90dUL,0xca46ca47UL,0x70d97017UL,0xdd4bddafUL,0x79de79edUL,0x67d467ffUL,0x23e82393UL,0xde4ade5bUL,0xbd6bbd06UL,0x7e2a7ebbUL,0x34e5347bUL,0x3a163ad7UL,0x54c554d2UL,0x62d762f8UL,0xff55ff99UL,0xa794a7b6UL,0x4acf4ac0UL,0x301030d9UL,0x0a060a0eUL,0x98819866UL,0x0bf00babUL,0xcc44ccb4UL,0xd5bad5f0UL,0x3ee33e75UL,0x0ef30eacUL,0x19fe1944UL,0x5bc05bdbUL,0x858a8580UL,0xecadecd3UL,0xdfbcdffeUL,0xd848d8a8UL,0x0c040cfdUL,0x7adf7a19UL,0x58c1582fUL,0x9f759f30UL,0xa563a5e7UL,0x50305070UL,0x2e1a2ecbUL,0x120e12efUL,0xb76db708UL,0xd44cd455UL,0x3c143c24UL,0x5f355f79UL,0x712f71b2UL,0x38e13886UL,0xfda2fdc8UL,0x4fcc4fc7UL,0x4b394b65UL,0xf957f96aUL,0x0df20d58UL,0x9d829d61UL,0xc947c9b3UL,0xefacef27UL,0x32e73288UL,0x7d2b7d4fUL,0xa495a442UL,0xfba0fb3bUL,0xb398b3aaUL,0x68d168f6UL,0x817f8122UL,0xaa66aaeeUL,0x827e82d6UL,0xe6abe6ddUL,0x9e839e95UL,0x45ca45c9UL,0x7b297bbcUL,0x6ed36e05UL,0x443c446cUL,0x8b798b2cUL,0x3de23d81UL,0x271d2731UL,0x9a769a37UL,0x4d3b4d96UL,0xfa56fa9eUL,0xd24ed2a6UL,0x221e2236UL,0x76db76e4UL,0x1e0a1e12UL,0xb46cb4fcUL,0x37e4378fUL,0xe75de778UL,0xb26eb20fUL,0x2aef2a69UL,0xf1a6f135UL,0xe3a8e3daUL,0xf7a4f7c6UL,0x5937598aUL,0x868b8674UL,0x56325683UL,0xc543c54eUL,0xeb59eb85UL,0xc2b7c218UL,0x8f8c8f8eUL,0xac64ac1dUL,0x6dd26df1UL,0x3be03b72UL,0xc7b4c71fUL,0x15fa15b9UL,0x090709faUL,0x6f256fa0UL,0xeaafea20UL,0x898e897dUL,0x20e92067UL,0x28182838UL,0x64d5640bUL,0x83888373UL,0xb16fb1fbUL,0x967296caUL,0x6c246c54UL,0x08f1085fUL,0x52c75221UL,0xf351f364UL,0x652365aeUL,0x847c8425UL,0xbf9cbf57UL,0x6321635dUL,0x7cdd7ceaUL,0x7fdc7f1eUL,0x9186919cUL,0x9485949bUL,0xab90ab4bUL,0xc642c6baUL,0x57c45726UL,0xe5aae529UL,0x73d873e3UL,0x0f050f09UL,0x030103f4UL,0x3612362aUL,0xfea3fe3cUL,0xe15fe18bUL,0x10f910beUL,0x6bd06b02UL,0xa891a8bfUL,0xe858e871UL,0x69276953UL,0xd0b9d0f7UL,0x48384891UL,0x351335deUL,0xceb3cee5UL,0x55335577UL,0xd6bbd604UL,0x90709039UL,0x80898087UL,0xf2a7f2c1UL,0xc1b6c1ecUL,0x6622665aUL,0xad92adb8UL,0x602060a9UL,0xdb49db5cUL,0x1aff1ab0UL,0x887888d8UL,0x8e7a8e2bUL,0x8a8f8a89UL,0x13f8134aUL,0x9b809b92UL,0x39173923UL,0x75da7510UL,0x53315384UL,0x51c651d5UL,0xd3b8d303UL,0x5ec35edcUL,0xcbb0cbe2UL,0x997799c3UL,0x3311332dUL,0x46cb463dUL,0x1ffc1fb7UL,0x61d6610cUL,0x4e3a4e62UL
};

// local table
LOCAL UINT32 groestl_T_local[256*8];
const UINT32 LOCAL *groestl_T0 = &groestl_T_local[0 * 256];
const UINT32 LOCAL *groestl_T1 = &groestl_T_local[1 * 256];
const UINT32 LOCAL *groestl_T2 = &groestl_T_local[2 * 256];
const UINT32 LOCAL *groestl_T3 = &groestl_T_local[3 * 256];
const UINT32 LOCAL *groestl_T4 = &groestl_T_local[4 * 256];
const UINT32 LOCAL *groestl_T5 = &groestl_T_local[5 * 256];
const UINT32 LOCAL *groestl_T6 = &groestl_T_local[6 * 256];
const UINT32 LOCAL *groestl_T7 = &groestl_T_local[7 * 256];

// init, once per kernel
UINT32 nLocalId = LOCALID;
{
for(i = 0; i < 256 * 8; i += WORKSIZE)
groestl_T_local[i + nLocalId ] = groestl_T_init[i + nLocalId];
}

// declarations
UINT32 hash[32]; // hash[16..31] - scratch buffer

UINT32 groestl_BuffB[32];
UINT32 groestl_BuffC[32];
unsigned groestl_i;
unsigned index;

// inlined function body
groestl_BuffC[16] = hash[16] = 0x80;
groestl_BuffC[17] = hash[17] = 0;
groestl_BuffC[18] = hash[18] = 0;
groestl_BuffC[19] = hash[19] = 0;
groestl_BuffC[20] = hash[20] = 0;
groestl_BuffC[21] = hash[21] = 0;
groestl_BuffC[22] = hash[22] = 0;
groestl_BuffC[23] = hash[23] = 0;
groestl_BuffC[24] = hash[24] = 0;
groestl_BuffC[25] = hash[25] = 0;
groestl_BuffC[26] = hash[26] = 0;
groestl_BuffC[27] = hash[27] = 0;
groestl_BuffC[28] = hash[28] = 0;
groestl_BuffC[29] = hash[29] = 0;
groestl_BuffC[30] = hash[30] = 0;
hash[31] = 0x01000000;
groestl_BuffC[31] = 0x01020000L;

#pragma unroll 16
for (groestl_i = 0; groestl_i < 16; groestl_i++)
{
groestl_BuffC[groestl_i] = hash[groestl_i];
}

for(groestl_i=0; groestl_i < 0x0d000000u; groestl_i+=0x01000000u)
{
groestl_QMIX(hash, groestl_BuffB, groestl_i)
groestl_i+=0x01000000u;
groestl_QMIX(groestl_BuffB, hash, groestl_i)
}

for(groestl_i=0; groestl_i<13; ++groestl_i)
{
groestl_PMIX(groestl_BuffC, groestl_BuffB, groestl_i)
++groestl_i;
groestl_PMIX(groestl_BuffB, groestl_BuffC, groestl_i)
}

#pragma unroll 32
for(groestl_i = 0; groestl_i < 32-1; groestl_i++)
{
hash[groestl_i] ^= groestl_BuffC[groestl_i];
groestl_BuffB[groestl_i] = hash[groestl_i];
}
hash[31] ^= 0x00020000UL ^ groestl_BuffC[31];
groestl_BuffB[31] = hash[31];

for(groestl_i = 0; groestl_i < 14;)
{
groestl_PMIX(groestl_BuffB, groestl_BuffC, groestl_i)
++groestl_i;
groestl_PMIX(groestl_BuffC, groestl_BuffB, groestl_i)
++groestl_i;
}

#pragma unroll 16
for(groestl_i = 0; groestl_i < 16; ++groestl_i)
{
hash[groestl_i] = groestl_BuffB[16+groestl_i] ^ hash[16+groestl_i];
}
Wolf0
Legendary
*
Offline Offline

Activity: 1652


Miner Developer


View Profile
March 05, 2015, 07:16:49 AM
 #278


Pallas,

Are you planning on adding myriad-groestl support in the future? If not, could you explain why not? Is it because your groestl kernel is already faster than the myriad-groestl?

Also, are you planning on putting your work on github? Again, if not, could you explain why not?

It seems to me that both are important ways to further your efforts and establish your reputation.

Best regards as always.

HR

Myr-Groestl must do SHA256 as well, IIRC - of course pure Groestl is faster.

Code:
Donations: BTC: 1WoLFdwcfNEg64fTYsX1P25KUzzSjtEZC -- XMR: 45SLUTzk7UXYHmzJ7bFN6FPfzTusdUVAZjPRgmEDw7G3SeimWM2kCdnDQXwDBYGUWaBtZNgjYtEYA22aMQT4t8KfU3vHLHG
pallas
Legendary
*
Offline Offline

Activity: 1386


Black Belt Developer


View Profile
March 05, 2015, 08:45:32 AM
 #279


Pallas,

Are you planning on adding myriad-groestl support in the future? If not, could you explain why not? Is it because your groestl kernel is already faster than the myriad-groestl?

Also, are you planning on putting your work on github? Again, if not, could you explain why not?

It seems to me that both are important ways to further your efforts and establish your reputation.

Best regards as always.

HR

Myr-Groestl must do SHA256 as well, IIRC - of course pure Groestl is faster.

myr-groestl should be faster because its has a single round of groestl (14 iterations) + sha; groestlcoin is groestl + groestl again, so slower.
it's just that I do not have enough free time to work on all these algos.....
Now wolf0 just did a fantastic job on whirlpoolx and I want to understand the magic ;-)

Wolf0
Legendary
*
Offline Offline

Activity: 1652


Miner Developer


View Profile
March 05, 2015, 08:48:41 AM
 #280


Pallas,

Are you planning on adding myriad-groestl support in the future? If not, could you explain why not? Is it because your groestl kernel is already faster than the myriad-groestl?

Also, are you planning on putting your work on github? Again, if not, could you explain why not?

It seems to me that both are important ways to further your efforts and establish your reputation.

Best regards as always.

HR

Myr-Groestl must do SHA256 as well, IIRC - of course pure Groestl is faster.

myr-groestl should be faster because its has a single round of groestl (14 iterations) + sha; groestlcoin is groestl + groestl again, so slower.
it's just that I do not have enough free time to work on all these algos.....
Now wolf0 just did a fantastic job on whirlpoolx and I want to understand the magic ;-)

Haha, you ain't seen impressive yet! Check the thread, I'm about to post again!

Code:
Donations: BTC: 1WoLFdwcfNEg64fTYsX1P25KUzzSjtEZC -- XMR: 45SLUTzk7UXYHmzJ7bFN6FPfzTusdUVAZjPRgmEDw7G3SeimWM2kCdnDQXwDBYGUWaBtZNgjYtEYA22aMQT4t8KfU3vHLHG
Pages: « 1 2 3 4 5 6 7 8 9 10 11 12 13 [14] 15 16 17 18 19 20 21 »  All
  Print  
 
Jump to:  

Sponsored by , a Bitcoin-accepting VPN.
Powered by MySQL Powered by PHP Powered by SMF 1.1.19 | SMF © 2006-2009, Simple Machines Valid XHTML 1.0! Valid CSS!