[ANN] sph-sgminer: multi-coin multi-algorithm GPU miner | added MaruCoin

[janos666]

How do I set the clock outside of sgminer in Linux?

aticonfig --adapter=all --odsc=1000,1450

You might need to be root. The clocks are [GPU,VRAM]. You can set unique clocks for every cards (starting from adapter=0).

aticonfig --adapter=all --odcc

This applies the earlier "scratchpad" settings (makes them permanent through reboots).

Use can issue this command to read your actual clocks (and limits + current GPU load percentage which can also be very useful!):

aticonfig --adapter=all --odgc

[CryptoGretzky]

How do I set the clock outside of sgminer in Linux?

aticonfig --adapter=all --odsc=1000,1450

You might need to be root. The clocks are [GPU,VRAM]. You can set unique clocks for every cards (starting from adapter=0).

aticonfig --adapter=all --odcc

This applies the earlier "scratchpad" settings (makes them permanent through reboots).

Use can issue this command to read your actual clocks (and limits + current GPU load percentage which can also be very useful!):

aticonfig --adapter=all --odgc

Thanks for the info.   I could run --odgc fine, but when I use --odsc, it ask me to also set powerstate and performance level.   I tried --lsp, --list-powerstate, but somehow it's not recognizing the command even though they are clearly shown in man for aticonfig....     I did get some more info that I am posting in another post now though.

[CryptoGretzky]

I ran aticonfig --odgc --adapter=all for the rig that drops after a min. 

On boot, it's still running at 2.5-2.6 Mh, all 4 GPU's data are the same.  eg. core, mem are 1025,1400, bus speed is 5000:



After a few mins it started dropping down to 2.3 Mh.   ALL GPU's data remained the same.  eg. core, mem are 1025,1400, bus speed is 5000:



So, nothing really changed at all from the GPU's setting perspective....


On my other rig that held steady at 2.77 Mh/s now over 15 hours at this speed, core, mem is the same as my other rig, BUT bus speed is only 2500



Any one know how to set bus speed independently so I can test further to see if that has anything to do with it??

[DJstone]

I ran aticonfig --odgc --adapter=all for the rig that drops after a min. 

On boot, it's still running at 2.5-2.6 Mh, all 4 GPU's data are the same.  eg. core, mem are 1025,1400, bus speed is 5000:

https://i.imgur.com/WBSL7Gr.png

After a few mins it started dropping down to 2.3 Mh.   ALL GPU's data remained the same.  eg. core, mem are 1025,1400, bus speed is 5000:

https://i.imgur.com/7N8EPpu.png

So, nothing really changed at all from the GPU's setting perspective....


On my other rig that held steady at 2.77 Mh/s now over 15 hours at this speed, core, mem is the same as my other rig, BUT bus speed is only 2500

https://i.imgur.com/ypvTYVh.png

Any one know how to set bus speed independently so I can test further to see if that has anything to do with it??

As far as i know, the Bus speed is set over the Bios.

You may have the option to set the Bus speed via any Interface is used by f.e. AMD OverDriver.

[janos666]

Any one know how to set bus speed independently so I can test further to see if that has anything to do with it??

The closest thing which comes to my mind is this kernel config:
"Bus Options" \ "PCI Support" \ "PCI Express AMPS Control" \ "Default AMPS Policy" ---> "BIOS Default" / "Performance" / "Powersave"

I keep this at "Performance".
But I don't have this number in my aticonfig --odgc output.

However, this gentoo setup is highly unstable in many ways and levels. I am just trying to make it work because I had to get rid of Xubuntu asap.
For example, I have to run sgminer with --no-adl or I get segmentation fault. It finds 24 ADL devices for 4 physical cards. And my desktop is all black.

[CryptoGretzky]

I used the following /etc/X11/xorg.conf:

Code:

Section "ServerLayout"
	Identifier     "aticonfig Layout"
	Screen      0  "aticonfig-Screen[0]-0" 0 0
EndSection

Section "Module"
EndSection

Section "Monitor"
	Identifier   "aticonfig-Monitor[0]-0"
	Option	    "VendorName" "ATI Proprietary Driver"
	Option	    "ModelName" "Generic Autodetecting Monitor"
	Option	    "DPMS" "true"
EndSection

Section "Device"
	Identifier  "aticonfig-Device[0]-0"
	Driver      "fglrx"
	BusID       "PCI:1:0:0"
EndSection

Section "Screen"
	Identifier "aticonfig-Screen[0]-0"
	Device     "aticonfig-Device[0]-0"
	Monitor    "aticonfig-Monitor[0]-0"
	DefaultDepth     24
	SubSection "Display"
		Viewport   0 0
		Depth     24
	EndSubSection
EndSection

to get 3 out of 4 of my 290x GPU to be hashing stable at 2.65Mh:



aticonfig --adapter=all --odgc still reported the same 1025/1400, 5000 bus speed for all GPU as before.   Only this time, the hash rate stays the same for the 3 GPU that I took out of xorg.conf.....

This is ultra baffling....

[flipme]

I run the pimp distro since a couple of hours.
Also on AsRock H81 BTC and nothing changed, still dropping.

[CryptoGretzky]

I run the pimp distro since a couple of hours.
Also on AsRock H81 BTC and nothing changed, still dropping.

Did you try the xorg.conf that I posted the message before yours to see if it helps with the rate dropping?

[flipme]

I run the pimp distro since a couple of hours.
Also on AsRock H81 BTC and nothing changed, still dropping.

Did you try the xorg.conf that I posted the message before yours to see if it helps with the rate dropping?

I do since 15 minutes and so far no drop!
Is there a screensaver starting on X?

[CryptoGretzky]

I run the pimp distro since a couple of hours.
Also on AsRock H81 BTC and nothing changed, still dropping.

Did you try the xorg.conf that I posted the message before yours to see if it helps with the rate dropping?

I do since 15 minutes and so far no drop!
Is there a screensaver starting on X?

The screen does black out after a few mins.   Another interesting thing about my other rig that didn't drop at all is that I unplugged the monitor cable from the GPU so I could use it on my Asrock H81 board that do need a monitor plugged in to start.   Don't know if that has anything to with it, but... let me do some test about that monitor being powered off by X scenario....

[boubou]

I just replace my xorg.conf, but did not understand what theses lines means.

I have 4 cards on my rigg, 1 270 with no drop and 3 290 (with drop)
Should I put the 270 on gpu0?

EDIT: I'm under BAMT, the xorg.conf changes back to the original one after a coldreboot, so I can't test if there is some changes.

[CryptoGretzky]

I just replace my xorg.conf, but did not understand what theses lines means.

I have 4 cards on my rigg, 1 270 with no drop and 3 290 (with drop)
Should I put the 270 on gpu0?

EDIT: I'm under BAMT, the xorg.conf changes back to the original one after a coldreboot, so I can't test if there is some changes.

I have SOLVED the problem.  Both in PiMP 1.1beta and BAMT 1.6.   Probably will work on other versions too.

Ok, you can replace your original xorg.conf that detects all cards or run "aticonfig --adapter=all --initial -f".

add these 4 lines into /etc/rc.local before exit 0

Code:

export DISPLAY=:0
xset s off
xset -dpms
xset s noblank

Just in case it's not already executable, run:

Code:

chmod +x /etc/rc.local

Reboot, then the drop in hash will not happen anymore.

Here's my DRK address for the bounty: XbUfQkPJNvA8yntkW6JRXZcsYCcSDNL3Hh  

btc: 12WYGCtEPSsYJsgsseKehS5Gu7QMDgSgGJ
hiro: HSLN3cBQ5P5Jv6FokJVMWaiY9gPjPh58gP

[flipme]

Yeah, its due to DPMS (Display Power Management Signaling)

You can also apply the settings in /etc/X11/xorg.conf

In Section "ServerLayout" add
    

Code:

    
Option "StandbyTime" "0"
Option "SuspendTime" "0"
Option "BlankTime" "0"

In "SectionMonitor" change

Code:

Option  "DPMS" "false"

It ran smooth the whole night, no crashes, very steady.
As far as I can see the PIMP distro from Bee doesn't have the "donation mining" pool switcher included.
Quite nice new interface altogether, boots very fast.

Don't know who'll get the bounty for this, as I had the idea it was the screensaver pulling it down.
In case: DRK XsRa7BwaWHTSacRyE27Npxm56CuLGTZvwb

[smolen]

Long, long ago I've promised to open-source my smelter miner and still found no time to do it properly Here is groestl code from smelter, not sure if it fit into sgminer as is, but you can get the tricks for your project if you wish.

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];
	}

[phm]

Is it possible future development could add support for Simple Inflation Coin which is the father of Quark, secure and all similar clones

Done.

[phm]

The screen does black out after a few mins.   Another interesting thing about my other rig that didn't drop at all is that I unplugged the monitor cable from the GPU so I could use it on my Asrock H81 board that do need a monitor plugged in to start.   Don't know if that has anything to with it, but... let me do some test about that monitor being powered off by X scenario....

No wonder I didn't experience the drop, my rig has no monitor connected.

[phm]

Long, long ago I've promised to open-source my smelter miner and still found no time to do it properly Here is groestl code from smelter, not sure if it fit into sgminer as is, but you can get the tricks for your project if you wish.

Great, thank you!

[subhead]

The screen does black out after a few mins.   Another interesting thing about my other rig that didn't drop at all is that I unplugged the monitor cable from the GPU so I could use it on my Asrock H81 board that do need a monitor plugged in to start.   Don't know if that has anything to with it, but... let me do some test about that monitor being powered off by X scenario....

No wonder I didn't experience the drop, my rig has no monitor connected.

the drop also happens on some of my rigs without a monitor attached. all rigs running 290 and 290x cards.

[boubou]

The screen does black out after a few mins.   Another interesting thing about my other rig that didn't drop at all is that I unplugged the monitor cable from the GPU so I could use it on my Asrock H81 board that do need a monitor plugged in to start.   Don't know if that has anything to with it, but... let me do some test about that monitor being powered off by X scenario....

No wonder I didn't experience the drop, my rig has no monitor connected.

the drop also happens on some of my rigs without a monitor attached. all rigs running 290 and 290x cards.

run "aticonfig --adapter=all --initial -f".

add these 4 lines into /etc/rc.local before exit 0

Code:

export DISPLAY=:0
xset s off
xset -dpms
xset s noblank

Just in case it's not already executable, run:

Code:

chmod +x /etc/rc.local

I have sent a 30 Darkcoin Bounty to cryptogretzky by myself.

As I only have 3 r9 290 I'll make 3 month to recover that bounty with my increased hasrate, but I'll be in peace with this bug.

If any-one want to pay me back a part of that bounty, here is my address to share the cost.
XavaNMidPmofBUt4RWLf6YHhNDUNA7TjY6

[janos666]

I tried to set the parameters mentioned above but I am still experiencing speed drops over time.
My drops always happen randomly, not after fix 10 minutes (could take just a few or even a couple thousand minutes until one or more cards slow down). And I never enabled either a screensaver/lock, nor any energy saving features (dpms was enabled in the xorg.conf by aticonfig as default but I think that depends on either the screensaver daemon or something else I never enabled -or even installed- during the gentoo setup).