diff --git a/Makefile.am b/Makefile.am
index 5a0a806..e886198 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -16,8 +16,8 @@ bin_PROGRAMS	= ccminer
 ccminer_SOURCES		= elist.h miner.h compat.h \
 			  compat/inttypes.h compat/stdbool.h compat/unistd.h \
 			  compat/sys/time.h compat/getopt/getopt.h \
-			  cpu-miner.c util.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c hefty1.c scrypt.c sha2.c \
-			  sph/sph_blake.h sph/sph_groestl.h sph/sph_jh.h sph/sph_keccak.h sph/sph_skein.h sph/sph_types.h \
+			  cpu-miner.c util.c sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c hefty1.c scrypt.c sha2.c \
+			  sph/bmw.h sph/sph_blake.h sph/sph_groestl.h sph/sph_jh.h sph/sph_keccak.h sph/sph_skein.h sph/sph_types.h \
 			  heavy/heavy.cu \
 			  heavy/cuda_blake512.cu heavy/cuda_blake512.h \
 			  heavy/cuda_combine.cu heavy/cuda_combine.h \
@@ -30,6 +30,8 @@ ccminer_SOURCES		= elist.h miner.h compat.h \
 			  JHA/jackpotcoin.cu JHA/cuda_jha_keccak512.cu \
 			  JHA/cuda_jha_compactionTest.cu quark/cuda_quark_checkhash.cu \
 			  quark/cuda_jh512.cu quark/cuda_quark_blake512.cu quark/cuda_quark_groestl512.cu quark/cuda_skein512.cu \
+			  quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu quark/quarkcoin.cu quark/animecoin.cu \
+			  quark/cuda_quark_compactionTest.cu \
 			  myriadgroestl.cpp cuda_myriadgroestl.cu
 
 ccminer_LDFLAGS		= $(PTHREAD_FLAGS) @CUDA_LDFLAGS@
@@ -42,5 +44,8 @@ ccminer_CPPFLAGS	= -msse2 @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(PTHREAD_FLAGS) -f
 
 # ABI requiring code modules
 # this module doesn't compile with Compute 2.0 unfortunately
+quark/cuda_quark_compactionTest.o: quark/cuda_quark_compactionTest.cu
+	$(NVCC) -I . -I cudpp-2.1/include @CFLAGS@ -Xptxas "-abi=yes -v" -gencode=arch=compute_30,code=\"sm_30,compute_30\" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --maxrregcount=80 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<
+
 JHA/cuda_jha_compactionTest.o: JHA/cuda_jha_compactionTest.cu
 	$(NVCC) -I . -I cudpp-2.1/include @CFLAGS@ -Xptxas "-abi=yes -v" -gencode=arch=compute_30,code=\"sm_30,compute_30\" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --maxrregcount=80 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<
diff --git a/Makefile.in b/Makefile.in
index 93fc56e..9da7878 100644
--- a/Makefile.in
+++ b/Makefile.in
@@ -55,18 +55,26 @@ am__installdirs = "$(DESTDIR)$(bindir)"
 PROGRAMS = $(bin_PROGRAMS)
 am__dirstamp = $(am__leading_dot)dirstamp
 am_ccminer_OBJECTS = ccminer-cpu-miner.$(OBJEXT) \
-	ccminer-util.$(OBJEXT) ccminer-blake.$(OBJEXT) \
-	ccminer-groestl.$(OBJEXT) ccminer-keccak.$(OBJEXT) \
-	ccminer-hefty1.$(OBJEXT) ccminer-scrypt.$(OBJEXT) \
-	ccminer-sha2.$(OBJEXT) heavy.$(OBJEXT) cuda_blake512.$(OBJEXT) \
-	cuda_combine.$(OBJEXT) cuda_groestl512.$(OBJEXT) \
-	cuda_hefty1.$(OBJEXT) cuda_keccak512.$(OBJEXT) \
-	cuda_sha256.$(OBJEXT) ccminer-fuguecoin.$(OBJEXT) \
+	ccminer-util.$(OBJEXT) ccminer-bmw.$(OBJEXT) \
+	ccminer-blake.$(OBJEXT) ccminer-groestl.$(OBJEXT) \
+	ccminer-jh.$(OBJEXT) ccminer-keccak.$(OBJEXT) \
+	ccminer-skein.$(OBJEXT) ccminer-hefty1.$(OBJEXT) \
+	ccminer-scrypt.$(OBJEXT) ccminer-sha2.$(OBJEXT) \
+	heavy/heavy.$(OBJEXT) heavy/cuda_blake512.$(OBJEXT) \
+	heavy/cuda_combine.$(OBJEXT) heavy/cuda_groestl512.$(OBJEXT) \
+	heavy/cuda_hefty1.$(OBJEXT) heavy/cuda_keccak512.$(OBJEXT) \
+	heavy/cuda_sha256.$(OBJEXT) ccminer-fuguecoin.$(OBJEXT) \
 	cuda_fugue256.$(OBJEXT) ccminer-fugue.$(OBJEXT) \
 	ccminer-groestlcoin.$(OBJEXT) cuda_groestlcoin.$(OBJEXT) \
 	JHA/jackpotcoin.$(OBJEXT) JHA/cuda_jha_keccak512.$(OBJEXT) \
-	ccminer-jh.$(OBJEXT) ccminer-skein.$(OBJEXT) \
+	JHA/cuda_jha_compactionTest.$(OBJEXT) \
 	quark/cuda_quark_checkhash.$(OBJEXT) \
+	quark/cuda_jh512.$(OBJEXT) quark/cuda_quark_blake512.$(OBJEXT) \
+	quark/cuda_quark_groestl512.$(OBJEXT) \
+	quark/cuda_skein512.$(OBJEXT) quark/cuda_bmw512.$(OBJEXT) \
+	quark/cuda_quark_keccak512.$(OBJEXT) quark/quarkcoin.$(OBJEXT) \
+	quark/animecoin.$(OBJEXT) \
+	quark/cuda_quark_compactionTest.$(OBJEXT) \
 	ccminer-myriadgroestl.$(OBJEXT) cuda_myriadgroestl.$(OBJEXT)
 ccminer_OBJECTS = $(am_ccminer_OBJECTS)
 ccminer_DEPENDENCIES =
@@ -272,19 +280,22 @@ SUBDIRS = compat
 ccminer_SOURCES = elist.h miner.h compat.h \
 			  compat/inttypes.h compat/stdbool.h compat/unistd.h \
 			  compat/sys/time.h compat/getopt/getopt.h \
-			  cpu-miner.c util.c sph/blake.c sph/groestl.c sph/keccak.c hefty1.c scrypt.c sha2.c \
-			  sph/sph_blake.h sph/sph_groestl.h sph/sph_keccak.h sph/sph_types.h \
-			  heavy.cu \
-			  cuda_blake512.cu cuda_blake512.h \
-			  cuda_combine.cu cuda_combine.h \
-			  cuda_groestl512.cu cuda_groestl512.h \
-			  cuda_hefty1.cu cuda_hefty1.h \
-			  cuda_keccak512.cu cuda_keccak512.h \
-			  cuda_sha256.cu cuda_sha256.h \
+			  cpu-miner.c util.c sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c hefty1.c scrypt.c sha2.c \
+			  sph/bmw.h sph/sph_blake.h sph/sph_groestl.h sph/sph_jh.h sph/sph_keccak.h sph/sph_skein.h sph/sph_types.h \
+			  heavy/heavy.cu \
+			  heavy/cuda_blake512.cu heavy/cuda_blake512.h \
+			  heavy/cuda_combine.cu heavy/cuda_combine.h \
+			  heavy/cuda_groestl512.cu heavy/cuda_groestl512.h \
+			  heavy/cuda_hefty1.cu heavy/cuda_hefty1.h \
+			  heavy/cuda_keccak512.cu heavy/cuda_keccak512.h \
+			  heavy/cuda_sha256.cu heavy/cuda_sha256.h \
 			  fuguecoin.cpp cuda_fugue256.cu sph/fugue.c sph/sph_fugue.h uint256.h \
 			  groestlcoin.cpp cuda_groestlcoin.cu cuda_groestlcoin.h \
-			  JHA/jackpotcoin.cu JHA/cuda_jha_keccak512.cu sph/jh.c sph/skein.c \
-			  sph/sph_jh.h sph/sph_skein.h quark/cuda_quark_checkhash.cu \
+			  JHA/jackpotcoin.cu JHA/cuda_jha_keccak512.cu \
+			  JHA/cuda_jha_compactionTest.cu quark/cuda_quark_checkhash.cu \
+			  quark/cuda_jh512.cu quark/cuda_quark_blake512.cu quark/cuda_quark_groestl512.cu quark/cuda_skein512.cu \
+			  quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu quark/quarkcoin.cu quark/animecoin.cu \
+			  quark/cuda_quark_compactionTest.cu \
 			  myriadgroestl.cpp cuda_myriadgroestl.cu
 
 ccminer_LDFLAGS = $(PTHREAD_FLAGS) @CUDA_LDFLAGS@
@@ -381,6 +392,26 @@ uninstall-binPROGRAMS:
 
 clean-binPROGRAMS:
 	-test -z "$(bin_PROGRAMS)" || rm -f $(bin_PROGRAMS)
+heavy/$(am__dirstamp):
+	@$(MKDIR_P) heavy
+	@: > heavy/$(am__dirstamp)
+heavy/$(DEPDIR)/$(am__dirstamp):
+	@$(MKDIR_P) heavy/$(DEPDIR)
+	@: > heavy/$(DEPDIR)/$(am__dirstamp)
+heavy/heavy.$(OBJEXT): heavy/$(am__dirstamp) \
+	heavy/$(DEPDIR)/$(am__dirstamp)
+heavy/cuda_blake512.$(OBJEXT): heavy/$(am__dirstamp) \
+	heavy/$(DEPDIR)/$(am__dirstamp)
+heavy/cuda_combine.$(OBJEXT): heavy/$(am__dirstamp) \
+	heavy/$(DEPDIR)/$(am__dirstamp)
+heavy/cuda_groestl512.$(OBJEXT): heavy/$(am__dirstamp) \
+	heavy/$(DEPDIR)/$(am__dirstamp)
+heavy/cuda_hefty1.$(OBJEXT): heavy/$(am__dirstamp) \
+	heavy/$(DEPDIR)/$(am__dirstamp)
+heavy/cuda_keccak512.$(OBJEXT): heavy/$(am__dirstamp) \
+	heavy/$(DEPDIR)/$(am__dirstamp)
+heavy/cuda_sha256.$(OBJEXT): heavy/$(am__dirstamp) \
+	heavy/$(DEPDIR)/$(am__dirstamp)
 JHA/$(am__dirstamp):
 	@$(MKDIR_P) JHA
 	@: > JHA/$(am__dirstamp)
@@ -391,6 +422,8 @@ JHA/jackpotcoin.$(OBJEXT): JHA/$(am__dirstamp) \
 	JHA/$(DEPDIR)/$(am__dirstamp)
 JHA/cuda_jha_keccak512.$(OBJEXT): JHA/$(am__dirstamp) \
 	JHA/$(DEPDIR)/$(am__dirstamp)
+JHA/cuda_jha_compactionTest.$(OBJEXT): JHA/$(am__dirstamp) \
+	JHA/$(DEPDIR)/$(am__dirstamp)
 quark/$(am__dirstamp):
 	@$(MKDIR_P) quark
 	@: > quark/$(am__dirstamp)
@@ -399,20 +432,56 @@ quark/$(DEPDIR)/$(am__dirstamp):
 	@: > quark/$(DEPDIR)/$(am__dirstamp)
 quark/cuda_quark_checkhash.$(OBJEXT): quark/$(am__dirstamp) \
 	quark/$(DEPDIR)/$(am__dirstamp)
+quark/cuda_jh512.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/cuda_quark_blake512.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/cuda_quark_groestl512.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/cuda_skein512.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/cuda_bmw512.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/cuda_quark_keccak512.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/quarkcoin.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/animecoin.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
+quark/cuda_quark_compactionTest.$(OBJEXT): quark/$(am__dirstamp) \
+	quark/$(DEPDIR)/$(am__dirstamp)
 ccminer$(EXEEXT): $(ccminer_OBJECTS) $(ccminer_DEPENDENCIES) $(EXTRA_ccminer_DEPENDENCIES) 
 	@rm -f ccminer$(EXEEXT)
 	$(ccminer_LINK) $(ccminer_OBJECTS) $(ccminer_LDADD) $(LIBS)
 
 mostlyclean-compile:
 	-rm -f *.$(OBJEXT)
+	-rm -f JHA/cuda_jha_compactionTest.$(OBJEXT)
 	-rm -f JHA/cuda_jha_keccak512.$(OBJEXT)
 	-rm -f JHA/jackpotcoin.$(OBJEXT)
+	-rm -f heavy/cuda_blake512.$(OBJEXT)
+	-rm -f heavy/cuda_combine.$(OBJEXT)
+	-rm -f heavy/cuda_groestl512.$(OBJEXT)
+	-rm -f heavy/cuda_hefty1.$(OBJEXT)
+	-rm -f heavy/cuda_keccak512.$(OBJEXT)
+	-rm -f heavy/cuda_sha256.$(OBJEXT)
+	-rm -f heavy/heavy.$(OBJEXT)
+	-rm -f quark/animecoin.$(OBJEXT)
+	-rm -f quark/cuda_bmw512.$(OBJEXT)
+	-rm -f quark/cuda_jh512.$(OBJEXT)
+	-rm -f quark/cuda_quark_blake512.$(OBJEXT)
 	-rm -f quark/cuda_quark_checkhash.$(OBJEXT)
+	-rm -f quark/cuda_quark_compactionTest.$(OBJEXT)
+	-rm -f quark/cuda_quark_groestl512.$(OBJEXT)
+	-rm -f quark/cuda_quark_keccak512.$(OBJEXT)
+	-rm -f quark/cuda_skein512.$(OBJEXT)
+	-rm -f quark/quarkcoin.$(OBJEXT)
 
 distclean-compile:
 	-rm -f *.tab.c
 
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ccminer-blake.Po@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ccminer-bmw.Po@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ccminer-cpu-miner.Po@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ccminer-fugue.Po@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/ccminer-fuguecoin.Po@am__quote@
@@ -469,6 +538,20 @@ ccminer-util.obj: util.c
 @AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
 @am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-util.obj `if test -f 'util.c'; then $(CYGPATH_W) 'util.c'; else $(CYGPATH_W) '$(srcdir)/util.c'; fi`
 
+ccminer-bmw.o: sph/bmw.c
+@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-bmw.o -MD -MP -MF $(DEPDIR)/ccminer-bmw.Tpo -c -o ccminer-bmw.o `test -f 'sph/bmw.c' || echo '$(srcdir)/'`sph/bmw.c
+@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-bmw.Tpo $(DEPDIR)/ccminer-bmw.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/bmw.c' object='ccminer-bmw.o' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-bmw.o `test -f 'sph/bmw.c' || echo '$(srcdir)/'`sph/bmw.c
+
+ccminer-bmw.obj: sph/bmw.c
+@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-bmw.obj -MD -MP -MF $(DEPDIR)/ccminer-bmw.Tpo -c -o ccminer-bmw.obj `if test -f 'sph/bmw.c'; then $(CYGPATH_W) 'sph/bmw.c'; else $(CYGPATH_W) '$(srcdir)/sph/bmw.c'; fi`
+@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-bmw.Tpo $(DEPDIR)/ccminer-bmw.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/bmw.c' object='ccminer-bmw.obj' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-bmw.obj `if test -f 'sph/bmw.c'; then $(CYGPATH_W) 'sph/bmw.c'; else $(CYGPATH_W) '$(srcdir)/sph/bmw.c'; fi`
+
 ccminer-blake.o: sph/blake.c
 @am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-blake.o -MD -MP -MF $(DEPDIR)/ccminer-blake.Tpo -c -o ccminer-blake.o `test -f 'sph/blake.c' || echo '$(srcdir)/'`sph/blake.c
 @am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-blake.Tpo $(DEPDIR)/ccminer-blake.Po
@@ -497,6 +580,20 @@ ccminer-groestl.obj: sph/groestl.c
 @AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
 @am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-groestl.obj `if test -f 'sph/groestl.c'; then $(CYGPATH_W) 'sph/groestl.c'; else $(CYGPATH_W) '$(srcdir)/sph/groestl.c'; fi`
 
+ccminer-jh.o: sph/jh.c
+@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-jh.o -MD -MP -MF $(DEPDIR)/ccminer-jh.Tpo -c -o ccminer-jh.o `test -f 'sph/jh.c' || echo '$(srcdir)/'`sph/jh.c
+@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-jh.Tpo $(DEPDIR)/ccminer-jh.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/jh.c' object='ccminer-jh.o' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-jh.o `test -f 'sph/jh.c' || echo '$(srcdir)/'`sph/jh.c
+
+ccminer-jh.obj: sph/jh.c
+@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-jh.obj -MD -MP -MF $(DEPDIR)/ccminer-jh.Tpo -c -o ccminer-jh.obj `if test -f 'sph/jh.c'; then $(CYGPATH_W) 'sph/jh.c'; else $(CYGPATH_W) '$(srcdir)/sph/jh.c'; fi`
+@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-jh.Tpo $(DEPDIR)/ccminer-jh.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/jh.c' object='ccminer-jh.obj' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-jh.obj `if test -f 'sph/jh.c'; then $(CYGPATH_W) 'sph/jh.c'; else $(CYGPATH_W) '$(srcdir)/sph/jh.c'; fi`
+
 ccminer-keccak.o: sph/keccak.c
 @am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-keccak.o -MD -MP -MF $(DEPDIR)/ccminer-keccak.Tpo -c -o ccminer-keccak.o `test -f 'sph/keccak.c' || echo '$(srcdir)/'`sph/keccak.c
 @am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-keccak.Tpo $(DEPDIR)/ccminer-keccak.Po
@@ -511,6 +608,20 @@ ccminer-keccak.obj: sph/keccak.c
 @AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
 @am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-keccak.obj `if test -f 'sph/keccak.c'; then $(CYGPATH_W) 'sph/keccak.c'; else $(CYGPATH_W) '$(srcdir)/sph/keccak.c'; fi`
 
+ccminer-skein.o: sph/skein.c
+@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-skein.o -MD -MP -MF $(DEPDIR)/ccminer-skein.Tpo -c -o ccminer-skein.o `test -f 'sph/skein.c' || echo '$(srcdir)/'`sph/skein.c
+@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-skein.Tpo $(DEPDIR)/ccminer-skein.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/skein.c' object='ccminer-skein.o' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-skein.o `test -f 'sph/skein.c' || echo '$(srcdir)/'`sph/skein.c
+
+ccminer-skein.obj: sph/skein.c
+@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-skein.obj -MD -MP -MF $(DEPDIR)/ccminer-skein.Tpo -c -o ccminer-skein.obj `if test -f 'sph/skein.c'; then $(CYGPATH_W) 'sph/skein.c'; else $(CYGPATH_W) '$(srcdir)/sph/skein.c'; fi`
+@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-skein.Tpo $(DEPDIR)/ccminer-skein.Po
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/skein.c' object='ccminer-skein.obj' libtool=no @AMDEPBACKSLASH@
+@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
+@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-skein.obj `if test -f 'sph/skein.c'; then $(CYGPATH_W) 'sph/skein.c'; else $(CYGPATH_W) '$(srcdir)/sph/skein.c'; fi`
+
 ccminer-hefty1.o: hefty1.c
 @am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-hefty1.o -MD -MP -MF $(DEPDIR)/ccminer-hefty1.Tpo -c -o ccminer-hefty1.o `test -f 'hefty1.c' || echo '$(srcdir)/'`hefty1.c
 @am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-hefty1.Tpo $(DEPDIR)/ccminer-hefty1.Po
@@ -567,34 +678,6 @@ ccminer-fugue.obj: sph/fugue.c
 @AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
 @am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-fugue.obj `if test -f 'sph/fugue.c'; then $(CYGPATH_W) 'sph/fugue.c'; else $(CYGPATH_W) '$(srcdir)/sph/fugue.c'; fi`
 
-ccminer-jh.o: sph/jh.c
-@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-jh.o -MD -MP -MF $(DEPDIR)/ccminer-jh.Tpo -c -o ccminer-jh.o `test -f 'sph/jh.c' || echo '$(srcdir)/'`sph/jh.c
-@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-jh.Tpo $(DEPDIR)/ccminer-jh.Po
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/jh.c' object='ccminer-jh.o' libtool=no @AMDEPBACKSLASH@
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
-@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-jh.o `test -f 'sph/jh.c' || echo '$(srcdir)/'`sph/jh.c
-
-ccminer-jh.obj: sph/jh.c
-@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-jh.obj -MD -MP -MF $(DEPDIR)/ccminer-jh.Tpo -c -o ccminer-jh.obj `if test -f 'sph/jh.c'; then $(CYGPATH_W) 'sph/jh.c'; else $(CYGPATH_W) '$(srcdir)/sph/jh.c'; fi`
-@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-jh.Tpo $(DEPDIR)/ccminer-jh.Po
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/jh.c' object='ccminer-jh.obj' libtool=no @AMDEPBACKSLASH@
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
-@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-jh.obj `if test -f 'sph/jh.c'; then $(CYGPATH_W) 'sph/jh.c'; else $(CYGPATH_W) '$(srcdir)/sph/jh.c'; fi`
-
-ccminer-skein.o: sph/skein.c
-@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-skein.o -MD -MP -MF $(DEPDIR)/ccminer-skein.Tpo -c -o ccminer-skein.o `test -f 'sph/skein.c' || echo '$(srcdir)/'`sph/skein.c
-@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-skein.Tpo $(DEPDIR)/ccminer-skein.Po
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/skein.c' object='ccminer-skein.o' libtool=no @AMDEPBACKSLASH@
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
-@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-skein.o `test -f 'sph/skein.c' || echo '$(srcdir)/'`sph/skein.c
-
-ccminer-skein.obj: sph/skein.c
-@am__fastdepCC_TRUE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -MT ccminer-skein.obj -MD -MP -MF $(DEPDIR)/ccminer-skein.Tpo -c -o ccminer-skein.obj `if test -f 'sph/skein.c'; then $(CYGPATH_W) 'sph/skein.c'; else $(CYGPATH_W) '$(srcdir)/sph/skein.c'; fi`
-@am__fastdepCC_TRUE@	$(am__mv) $(DEPDIR)/ccminer-skein.Tpo $(DEPDIR)/ccminer-skein.Po
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	source='sph/skein.c' object='ccminer-skein.obj' libtool=no @AMDEPBACKSLASH@
-@AMDEP_TRUE@@am__fastdepCC_FALSE@	DEPDIR=$(DEPDIR) $(CCDEPMODE) $(depcomp) @AMDEPBACKSLASH@
-@am__fastdepCC_FALSE@	$(CC) $(DEFS) $(DEFAULT_INCLUDES) $(INCLUDES) $(ccminer_CPPFLAGS) $(CPPFLAGS) $(AM_CFLAGS) $(CFLAGS) -c -o ccminer-skein.obj `if test -f 'sph/skein.c'; then $(CYGPATH_W) 'sph/skein.c'; else $(CYGPATH_W) '$(srcdir)/sph/skein.c'; fi`
-
 .cpp.o:
 @am__fastdepCXX_TRUE@	$(CXXCOMPILE) -MT $@ -MD -MP -MF $(DEPDIR)/$*.Tpo -c -o $@ $<
 @am__fastdepCXX_TRUE@	$(am__mv) $(DEPDIR)/$*.Tpo $(DEPDIR)/$*.Po
@@ -1012,6 +1095,8 @@ distclean-generic:
 	-test . = "$(srcdir)" || test -z "$(CONFIG_CLEAN_VPATH_FILES)" || rm -f $(CONFIG_CLEAN_VPATH_FILES)
 	-rm -f JHA/$(DEPDIR)/$(am__dirstamp)
 	-rm -f JHA/$(am__dirstamp)
+	-rm -f heavy/$(DEPDIR)/$(am__dirstamp)
+	-rm -f heavy/$(am__dirstamp)
 	-rm -f quark/$(DEPDIR)/$(am__dirstamp)
 	-rm -f quark/$(am__dirstamp)
 
@@ -1115,6 +1200,14 @@ uninstall-am: uninstall-binPROGRAMS
 .cu.o:
 	$(NVCC) @CFLAGS@ -I . -Xptxas "-abi=no -v" -gencode=arch=compute_20,code=\"sm_20,compute_20\" -gencode=arch=compute_30,code=\"sm_30,compute_30\" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --maxrregcount=80 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<
 
+# ABI requiring code modules
+# this module doesn't compile with Compute 2.0 unfortunately
+quark/cuda_quark_compactionTest.o: quark/cuda_quark_compactionTest.cu
+	$(NVCC) -I . -I cudpp-2.1/include @CFLAGS@ -Xptxas "-abi=yes -v" -gencode=arch=compute_30,code=\"sm_30,compute_30\" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --maxrregcount=80 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<
+
+JHA/cuda_jha_compactionTest.o: JHA/cuda_jha_compactionTest.cu
+	$(NVCC) -I . -I cudpp-2.1/include @CFLAGS@ -Xptxas "-abi=yes -v" -gencode=arch=compute_30,code=\"sm_30,compute_30\" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --maxrregcount=80 --ptxas-options=-v $(JANSSON_INCLUDES) -o $@ -c $<
+
 # Tell versions [3.59,3.63) of GNU make to not export all variables.
 # Otherwise a system limit (for SysV at least) may be exceeded.
 .NOEXPORT:
diff --git a/README.txt b/README.txt
index 666101b..dd9e936 100644
--- a/README.txt
+++ b/README.txt
@@ -1,5 +1,5 @@
 
-ccMiner release 0.7 (May 01th 2014) - "Jackpot II"
+ccMiner release 0.9 (May 06th 2014) - "Say Hi to Quark, Anime"
 -------------------------------------------------------------
 
 ***************************************************************
@@ -14,6 +14,8 @@ If you find this tool useful and like to support its continued
   DOGE donation address: DT9ghsGmez6ojVdEZgvaZbT2Z3TruXG6yP
    HVC donation address: HNN3PyyTMkDo4RkEjkWSGMwqia1yD8mwJN
    GRS donation address: FmJKJAhvyHWPeEVeLQHefr2naqgWc9ABTM
+   JPC donation address: JYFBypVDkk583yKWY4M46TG5vXG8hfgD2U
+   MNC donation address: MShgNUSYwybEbXLvJUtdNg1a7rUeiNgooK
 ***************************************************************
 
 >>> Introduction <<<
@@ -39,6 +41,8 @@ its command line interface and options.
                           groestl     use to mine Groestlcoin
                           myr-gr      use to mine Myriad-Groestl
                           jackpot     use to mine Jackpotcoin
+                          quark       use to mine Quarkcoin
+                          anime       use to mine Animecoin
 
   -d, --devices         gives a comma separated list of CUDA device IDs
                         to operate on. Device IDs start counting from 0!
@@ -117,6 +121,8 @@ from your old clunkers.
 
 >>> RELEASE HISTORY <<<
 
+  May 6th 2014    this adds the quark and animecoin algorithms.
+
   May 3rd 2014    add the MjollnirCoin hash algorithm for the upcomin
                   MjollnirCoin relaunch.
 
diff --git a/configure b/configure
index 741cf90..f7b8197 100755
--- a/configure
+++ b/configure
@@ -1,6 +1,6 @@
 #! /bin/sh
 # Guess values for system-dependent variables and create Makefiles.
-# Generated by GNU Autoconf 2.68 for ccminer 2014.04.27.
+# Generated by GNU Autoconf 2.68 for ccminer 2014.05.03.
 #
 #
 # Copyright (C) 1992, 1993, 1994, 1995, 1996, 1998, 1999, 2000, 2001,
@@ -557,8 +557,8 @@ MAKEFLAGS=
 # Identity of this package.
 PACKAGE_NAME='ccminer'
 PACKAGE_TARNAME='ccminer'
-PACKAGE_VERSION='2014.04.27'
-PACKAGE_STRING='ccminer 2014.04.27'
+PACKAGE_VERSION='2014.05.03'
+PACKAGE_STRING='ccminer 2014.05.03'
 PACKAGE_BUGREPORT=''
 PACKAGE_URL=''
 
@@ -1297,7 +1297,7 @@ if test "$ac_init_help" = "long"; then
   # Omit some internal or obsolete options to make the list less imposing.
   # This message is too long to be a string in the A/UX 3.1 sh.
   cat <<_ACEOF
-\`configure' configures ccminer 2014.04.27 to adapt to many kinds of systems.
+\`configure' configures ccminer 2014.05.03 to adapt to many kinds of systems.
 
 Usage: $0 [OPTION]... [VAR=VALUE]...
 
@@ -1368,7 +1368,7 @@ fi
 
 if test -n "$ac_init_help"; then
   case $ac_init_help in
-     short | recursive ) echo "Configuration of ccminer 2014.04.27:";;
+     short | recursive ) echo "Configuration of ccminer 2014.05.03:";;
    esac
   cat <<\_ACEOF
 
@@ -1469,7 +1469,7 @@ fi
 test -n "$ac_init_help" && exit $ac_status
 if $ac_init_version; then
   cat <<\_ACEOF
-ccminer configure 2014.04.27
+ccminer configure 2014.05.03
 generated by GNU Autoconf 2.68
 
 Copyright (C) 2010 Free Software Foundation, Inc.
@@ -1972,7 +1972,7 @@ cat >config.log <<_ACEOF
 This file contains any messages produced by compilers while
 running configure, to aid debugging if configure makes a mistake.
 
-It was created by ccminer $as_me 2014.04.27, which was
+It was created by ccminer $as_me 2014.05.03, which was
 generated by GNU Autoconf 2.68.  Invocation command line was
 
   $ $0 $@
@@ -2901,7 +2901,7 @@ fi
 
 # Define the identity of the package.
  PACKAGE='ccminer'
- VERSION='2014.04.27'
+ VERSION='2014.05.03'
 
 
 cat >>confdefs.h <<_ACEOF
@@ -7118,7 +7118,7 @@ cat >>$CONFIG_STATUS <<\_ACEOF || ac_write_fail=1
 # report actual input values of CONFIG_FILES etc. instead of their
 # values after options handling.
 ac_log="
-This file was extended by ccminer $as_me 2014.04.27, which was
+This file was extended by ccminer $as_me 2014.05.03, which was
 generated by GNU Autoconf 2.68.  Invocation command line was
 
   CONFIG_FILES    = $CONFIG_FILES
@@ -7184,7 +7184,7 @@ _ACEOF
 cat >>$CONFIG_STATUS <<_ACEOF || ac_write_fail=1
 ac_cs_config="`$as_echo "$ac_configure_args" | sed 's/^ //; s/[\\""\`\$]/\\\\&/g'`"
 ac_cs_version="\\
-ccminer config.status 2014.04.27
+ccminer config.status 2014.05.03
 configured by $0, generated by GNU Autoconf 2.68,
   with options \\"\$ac_cs_config\\"
 
diff --git a/configure.ac b/configure.ac
index 9fa301d..b6c9ae4 100644
--- a/configure.ac
+++ b/configure.ac
@@ -1,4 +1,4 @@
-AC_INIT([ccminer], [2014.05.03])
+AC_INIT([ccminer], [2014.05.06])
 
 AC_PREREQ([2.59c])
 AC_CANONICAL_SYSTEM
diff --git a/cpu-miner.c b/cpu-miner.c
index 3e625da..b0ca293 100644
--- a/cpu-miner.c
+++ b/cpu-miner.c
@@ -126,7 +126,9 @@ typedef enum {
 	ALGO_FUGUE256,		/* Fugue256 */
 	ALGO_GROESTL,
 	ALGO_MYR_GR,
-	ALGO_JACKPOT
+	ALGO_JACKPOT,
+	ALGO_QUARK,
+	ALGO_ANIME
 } sha256_algos;
 
 static const char *algo_names[] = {
@@ -135,7 +137,9 @@ static const char *algo_names[] = {
 	"fugue256",
 	"groestl",
 	"myr-gr",
-	"jackpot"
+	"jackpot",
+	"quark",
+	"anime"
 };
 
 bool opt_debug = false;
@@ -203,6 +207,8 @@ Options:\n\
                         groestl   Groestlcoin hash\n\
                         myr-gr    Myriad-Groestl hash\n\
                         jackpot   Jackpot hash\n\
+                        quark     Quark hash\n\
+                        anime     Animecoin hash\n\
   -d, --devices         takes a comma separated list of CUDA devices to use.\n\
                         Device IDs start counting from 0! Alternatively takes\n\
                         string names of your cards like gtx780ti or gt640#2\n\
@@ -877,6 +883,16 @@ static void *miner_thread(void *userdata)
 			                      max_nonce, &hashes_done);
 			break;
 
+                case ALGO_QUARK:
+                        rc = scanhash_quark(thr_id, work.data, work.target,
+                                              max_nonce, &hashes_done);
+                        break;
+
+                case ALGO_ANIME:
+                        rc = scanhash_anime(thr_id, work.data, work.target,
+                                              max_nonce, &hashes_done);
+                        break;
+
 		default:
 			/* should never happen */
 			goto out;
@@ -1429,7 +1445,7 @@ static void signal_handler(int sig)
 }
 #endif
 
-#define PROGRAM_VERSION "0.8"
+#define PROGRAM_VERSION "0.9"
 int main(int argc, char *argv[])
 {
 	struct thr_info *thr;
diff --git a/cpuminer-config.h b/cpuminer-config.h
index 31f94fc..c5c3ce5 100644
--- a/cpuminer-config.h
+++ b/cpuminer-config.h
@@ -152,7 +152,7 @@
 #define PACKAGE_NAME "ccminer"
 
 /* Define to the full name and version of this package. */
-#define PACKAGE_STRING "ccminer 2014.05.03"
+#define PACKAGE_STRING "ccminer 2014.05.06"
 
 /* Define to the one symbol short name of this package. */
 #undef PACKAGE_TARNAME
@@ -161,7 +161,7 @@
 #undef PACKAGE_URL
 
 /* Define to the version of this package. */
-#define PACKAGE_VERSION "2014.05.03"
+#define PACKAGE_VERSION "2014.05.06"
 
 /* If using the C implementation of alloca, define if you know the
    direction of stack growth for your system; otherwise it will be
diff --git a/miner.h b/miner.h
index 7a68b37..e3f4b8f 100644
--- a/miner.h
+++ b/miner.h
@@ -223,6 +223,14 @@ extern int scanhash_jackpot(int thr_id, uint32_t *pdata,
 	const uint32_t *ptarget, uint32_t max_nonce,
 	unsigned long *hashes_done);
 
+extern int scanhash_quark(int thr_id, uint32_t *pdata,
+	const uint32_t *ptarget, uint32_t max_nonce,
+	unsigned long *hashes_done);
+
+extern int scanhash_anime(int thr_id, uint32_t *pdata,
+	const uint32_t *ptarget, uint32_t max_nonce,
+	unsigned long *hashes_done);
+
 extern void fugue256_hash(unsigned char* output, const unsigned char* input, int len);
 extern void heavycoin_hash(unsigned char* output, const unsigned char* input, int len);
 extern void groestlcoin_hash(unsigned char* output, const unsigned char* input, int len);
diff --git a/quark/animecoin.cu b/quark/animecoin.cu
new file mode 100644
index 0000000..fa771aa
--- /dev/null
+++ b/quark/animecoin.cu
@@ -0,0 +1,292 @@
+
+extern "C"
+{
+#include "sph/sph_blake.h"
+#include "sph/sph_bmw.h"
+#include "sph/sph_groestl.h"
+#include "sph/sph_skein.h"
+#include "sph/sph_jh.h"
+#include "sph/sph_keccak.h"
+#include "miner.h"
+}
+
+#include <stdint.h>
+
+// aus cpu-miner.c
+extern int device_map[8];
+
+// Speicher für Input/Output der verketteten Hashfunktionen
+static uint32_t *d_hash[8];
+
+// Speicher zur Generierung der Noncevektoren für die bedingten Hashes
+static uint32_t *d_animeNonces[8];
+static uint32_t *d_branch1Nonces[8];
+static uint32_t *d_branch2Nonces[8];
+static uint32_t *d_branch3Nonces[8];
+
+extern void quark_blake512_cpu_init(int thr_id, int threads);
+extern void quark_blake512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_bmw512_cpu_init(int thr_id, int threads);
+extern void quark_bmw512_cpu_setBlock_80(void *pdata);
+extern void quark_bmw512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_outputHash, int order);
+extern void quark_bmw512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_outputHash, int order);
+
+extern void quark_groestl512_cpu_init(int thr_id, int threads);
+extern void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+extern void quark_doublegroestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_skein512_cpu_init(int thr_id, int threads);
+extern void quark_skein512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_keccak512_cpu_init(int thr_id, int threads);
+extern void quark_keccak512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_jh512_cpu_init(int thr_id, int threads);
+extern void quark_jh512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_check_cpu_init(int thr_id, int threads);
+extern void quark_check_cpu_setTarget(const void *ptarget);
+extern uint32_t quark_check_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order);
+
+extern void quark_compactTest_cpu_init(int thr_id, int threads);
+extern void quark_compactTest_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
+											uint32_t *d_nonces1, size_t *nrm1,
+											uint32_t *d_nonces2, size_t *nrm2,
+											int order);
+extern void quark_compactTest_single_false_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
+											uint32_t *d_nonces1, size_t *nrm1,
+											int order);
+
+// Original Quarkhash Funktion aus einem miner Quelltext
+inline void animehash(void *state, const void *input)
+{
+    sph_blake512_context ctx_blake;
+    sph_bmw512_context ctx_bmw;
+    sph_groestl512_context ctx_groestl;
+    sph_jh512_context ctx_jh;
+    sph_keccak512_context ctx_keccak;
+    sph_skein512_context ctx_skein;
+    
+    unsigned char hash[64];
+
+    sph_bmw512_init(&ctx_bmw);
+    // ZBMW;
+    sph_bmw512 (&ctx_bmw, (const void*) input, 80);
+    sph_bmw512_close(&ctx_bmw, (void*) hash);
+
+    sph_blake512_init(&ctx_blake);
+    // ZBLAKE;
+    sph_blake512 (&ctx_blake, hash, 64);
+    sph_blake512_close(&ctx_blake, (void*) hash);
+    
+    if (hash[0] & 0x8)
+    {
+        sph_groestl512_init(&ctx_groestl);
+        // ZGROESTL;
+        sph_groestl512 (&ctx_groestl, (const void*) hash, 64);
+        sph_groestl512_close(&ctx_groestl, (void*) hash);
+    }
+    else
+    {
+        sph_skein512_init(&ctx_skein);
+        // ZSKEIN;
+        sph_skein512 (&ctx_skein, (const void*) hash, 64);
+        sph_skein512_close(&ctx_skein, (void*) hash);
+    }
+    
+    sph_groestl512_init(&ctx_groestl);
+    // ZGROESTL;
+    sph_groestl512 (&ctx_groestl, (const void*) hash, 64);
+    sph_groestl512_close(&ctx_groestl, (void*) hash);
+
+    sph_jh512_init(&ctx_jh);
+    // ZJH;
+    sph_jh512 (&ctx_jh, (const void*) hash, 64);
+    sph_jh512_close(&ctx_jh, (void*) hash);
+
+    if (hash[0] & 0x8)
+    {
+        sph_blake512_init(&ctx_blake);
+        // ZBLAKE;
+        sph_blake512 (&ctx_blake, (const void*) hash, 64);
+        sph_blake512_close(&ctx_blake, (void*) hash);
+    }
+    else
+    {
+        sph_bmw512_init(&ctx_bmw);
+        // ZBMW;
+        sph_bmw512 (&ctx_bmw, (const void*) hash, 64);
+        sph_bmw512_close(&ctx_bmw, (void*) hash);
+    }
+
+    sph_keccak512_init(&ctx_keccak);
+    // ZKECCAK;
+    sph_keccak512 (&ctx_keccak, (const void*) hash, 64);
+    sph_keccak512_close(&ctx_keccak, (void*) hash);
+
+    sph_skein512_init(&ctx_skein);
+    // SKEIN;
+    sph_skein512 (&ctx_skein, (const void*) hash, 64);
+    sph_skein512_close(&ctx_skein, (void*) hash);
+
+    if (hash[0] & 0x8)
+    {
+        sph_keccak512_init(&ctx_keccak);
+        // ZKECCAK;
+        sph_keccak512 (&ctx_keccak, (const void*) hash, 64);
+        sph_keccak512_close(&ctx_keccak, (void*) hash);
+    }
+    else
+    {
+        sph_jh512_init(&ctx_jh);
+        // ZJH;
+        sph_jh512 (&ctx_jh, (const void*) hash, 64);
+        sph_jh512_close(&ctx_jh, (void*) hash);
+    }
+
+    memcpy(state, hash, 32);
+}
+
+
+struct HashPredicate
+{    
+    HashPredicate(uint32_t *hashes, uint32_t startNonce) :
+        m_hashes(hashes),
+        m_startNonce(startNonce)
+        { }
+
+    __device__
+    bool operator()(const uint32_t x)
+    {
+        uint32_t *hash = &m_hashes[(x - m_startNonce)*16];
+        return hash[0] & 0x8;
+    }
+
+    uint32_t *m_hashes;
+    uint32_t  m_startNonce;
+};
+
+extern bool opt_benchmark;
+
+extern "C" int scanhash_anime(int thr_id, uint32_t *pdata,
+    const uint32_t *ptarget, uint32_t max_nonce,
+    unsigned long *hashes_done)
+{
+	const uint32_t first_nonce = pdata[19];
+
+	// TODO: entfernen für eine Release! Ist nur zum Testen!
+	if (opt_benchmark)
+		((uint32_t*)ptarget)[7] = 0x00000f;
+
+	const uint32_t Htarg = ptarget[7];
+
+	const int throughput = 256*2048; // 100;
+
+	static bool init[8] = {0,0,0,0,0,0,0,0};
+	if (!init[thr_id])
+	{
+		cudaSetDevice(device_map[thr_id]);
+
+		// Konstanten kopieren, Speicher belegen
+		cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput);
+		quark_blake512_cpu_init(thr_id, throughput);
+		quark_groestl512_cpu_init(thr_id, throughput);
+		quark_skein512_cpu_init(thr_id, throughput);
+		quark_bmw512_cpu_init(thr_id, throughput);
+		quark_keccak512_cpu_init(thr_id, throughput);
+		quark_jh512_cpu_init(thr_id, throughput);
+		quark_check_cpu_init(thr_id, throughput);
+		quark_compactTest_cpu_init(thr_id, throughput);
+		cudaMalloc(&d_animeNonces[thr_id], sizeof(uint32_t)*throughput);
+		cudaMalloc(&d_branch1Nonces[thr_id], sizeof(uint32_t)*throughput);
+		cudaMalloc(&d_branch2Nonces[thr_id], sizeof(uint32_t)*throughput);
+		cudaMalloc(&d_branch3Nonces[thr_id], sizeof(uint32_t)*throughput);
+		init[thr_id] = true;
+	}
+
+	uint32_t endiandata[20];
+	for (int k=0; k < 20; k++)
+		be32enc(&endiandata[k], ((uint32_t*)pdata)[k]);
+
+	quark_bmw512_cpu_setBlock_80((void*)endiandata);
+	quark_check_cpu_setTarget(ptarget);
+
+	do {
+		int order = 0;
+		size_t nrm1=0, nrm2=0, nrm3=0;
+
+		// erstes BMW512 Hash mit CUDA
+		quark_bmw512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für Blake512
+		quark_blake512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+
+		quark_compactTest_single_false_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], NULL,
+				d_branch3Nonces[thr_id], &nrm3,
+				order++);
+		
+		// nur den Skein Branch weiterverfolgen
+		quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für Groestl512
+		quark_groestl512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für JH512
+		quark_jh512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// quarkNonces in branch1 und branch2 aufsplitten gemäss if (hash[0] & 0x8)
+		quark_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id],
+			d_branch1Nonces[thr_id], &nrm1,
+			d_branch2Nonces[thr_id], &nrm2,
+			order++);
+
+		// das ist der bedingte Branch für Blake512
+		quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der bedingte Branch für Bmw512
+		quark_bmw512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für Keccak512
+		quark_keccak512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für Skein512
+		quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// quarkNonces in branch1 und branch2 aufsplitten gemäss if (hash[0] & 0x8)
+		quark_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id],
+			d_branch1Nonces[thr_id], &nrm1,
+			d_branch2Nonces[thr_id], &nrm2,
+			order++);
+
+		// das ist der bedingte Branch für Keccak512
+		quark_keccak512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der bedingte Branch für JH512
+		quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++);
+
+		// Scan nach Gewinner Hashes auf der GPU
+		uint32_t foundNonce = quark_check_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+		if  (foundNonce != 0xffffffff)
+		{
+			uint32_t vhash64[8];
+			be32enc(&endiandata[19], foundNonce);
+			animehash(vhash64, endiandata);
+
+			if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) {
+
+				pdata[19] = foundNonce;
+				*hashes_done = (foundNonce - first_nonce + 1)/2;
+				return 1;
+			} else {
+				applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU!", thr_id, foundNonce);
+			}
+		}
+
+		pdata[19] += throughput;
+
+	} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
+
+	*hashes_done = (pdata[19] - first_nonce + 1)/2;
+	return 0;
+}
diff --git a/quark/cuda_bmw512.cu b/quark/cuda_bmw512.cu
new file mode 100644
index 0000000..7c706f2
--- /dev/null
+++ b/quark/cuda_bmw512.cu
@@ -0,0 +1,473 @@
+#if 1
+
+#include <cuda.h>
+#include "cuda_runtime.h"
+#include "device_launch_parameters.h"
+
+#include <stdio.h>
+#include <memory.h>
+
+// Folgende Definitionen später durch header ersetzen
+typedef unsigned char uint8_t;
+typedef unsigned int uint32_t;
+
+// Endian Drehung für 32 Bit Typen
+/*
+static __device__ uint32_t cuda_swab32(uint32_t x)
+{
+    return (((x << 24) & 0xff000000u) | ((x << 8) & 0x00ff0000u)
+          | ((x >> 8) & 0x0000ff00u) | ((x >> 24) & 0x000000ffu));
+}
+*/
+static __device__ uint32_t cuda_swab32(uint32_t x)
+{
+	return __byte_perm(x, 0, 0x0123);
+}
+// Endian Drehung für 64 Bit Typen
+static __device__ unsigned long long cuda_swab64(unsigned long long x) {
+    uint32_t h = (x >> 32);
+    uint32_t l = (x & 0xFFFFFFFFULL);
+    return (((unsigned long long)cuda_swab32(l)) << 32) | ((unsigned long long)cuda_swab32(h));
+}
+
+// das Hi Word aus einem 64 Bit Typen extrahieren
+static __device__ uint32_t HIWORD(const unsigned long long &x) {
+#if __CUDA_ARCH__ >= 130
+	return (uint32_t)__double2hiint(__longlong_as_double(x));
+#else
+	return (uint32_t)(x >> 32);
+#endif
+}
+
+// das Hi Word in einem 64 Bit Typen ersetzen
+static __device__ unsigned long long REPLACE_HIWORD(const unsigned long long &x, const uint32_t &y) {
+	return (x & 0xFFFFFFFFULL) | (((unsigned long long)y) << 32ULL);
+}
+
+// das Lo Word aus einem 64 Bit Typen extrahieren
+static __device__ uint32_t LOWORD(const unsigned long long &x) {
+#if __CUDA_ARCH__ >= 130
+	return (uint32_t)__double2loint(__longlong_as_double(x));
+#else
+	return (uint32_t)(x & 0xFFFFFFFFULL);
+#endif
+}
+
+static __device__ unsigned long long MAKE_ULONGLONG(uint32_t LO, uint32_t HI)
+{
+#if __CUDA_ARCH__ >= 130
+    return __double_as_longlong(__hiloint2double(HI, LO));
+#else
+	return (unsigned long long)LO | (((unsigned long long)HI) << 32ULL);
+#endif
+}
+
+// das Lo Word in einem 64 Bit Typen ersetzen
+static __device__ unsigned long long REPLACE_LOWORD(const unsigned long long &x, const uint32_t &y) {
+	return (x & 0xFFFFFFFF00000000ULL) | ((unsigned long long)y);
+}
+
+// der Versuch, einen Wrapper für einen aus 32 Bit Registern zusammengesetzten uin64_t Typen zu entferfen...
+#if 1
+typedef unsigned long long uint64_t;
+#else
+typedef class uint64
+{
+public:
+	__device__ uint64()
+	{
+	}
+	__device__ uint64(unsigned long long init)
+	{
+		val = make_uint2( LOWORD(init), HIWORD(init) );
+	}
+	__device__ uint64(uint32_t lo, uint32_t hi)
+	{
+		val = make_uint2( lo, hi );
+	}
+	__device__ const uint64 operator^(uint64 const& rhs) const
+	{
+		return uint64(val.x ^ rhs.val.x, val.y ^ rhs.val.y);
+	}
+	__device__ const uint64 operator|(uint64 const& rhs) const
+	{
+		return uint64(val.x | rhs.val.x, val.y | rhs.val.y);
+	}
+	__device__ const uint64 operator+(unsigned long long const& rhs) const
+	{
+		return *this+uint64(rhs);
+	}
+	__device__ const uint64 operator+(uint64 const& rhs) const
+	{
+		uint64 res;
+		asm ("add.cc.u32      %0, %2, %4;\n\t"
+			 "addc.cc.u32     %1, %3, %5;\n\t"
+			 : "=r"(res.val.x), "=r"(res.val.y)
+			 : "r"(    val.x), "r"(    val.y),
+			   "r"(rhs.val.x), "r"(rhs.val.y));
+		return res;
+	}
+	__device__ const uint64 operator-(uint64 const& rhs) const
+	{
+		uint64 res;
+		asm ("sub.cc.u32      %0, %2, %4;\n\t"
+			 "subc.cc.u32     %1, %3, %5;\n\t"
+			 : "=r"(res.val.x), "=r"(res.val.y)
+			 : "r"(    val.x), "r"(    val.y),
+			   "r"(rhs.val.x), "r"(rhs.val.y));
+		return res;
+	}
+	__device__ const uint64 operator<<(int n) const
+	{
+		return uint64(unsigned long long(*this)<<n);
+	}
+	__device__ const uint64 operator>>(int n) const
+	{
+		return uint64(unsigned long long(*this)>>n);
+	}
+	__device__ operator unsigned long long() const
+	{
+		return MAKE_ULONGLONG(val.x, val.y);
+	}
+	uint2 val;
+} uint64_t;
+#endif
+
+// aus heavy.cu
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+// die Message it Padding zur Berechnung auf der GPU
+__constant__ uint64_t c_PaddedMessage80[16]; // padded message (80 bytes + padding)
+
+#define SPH_C64(x)    ((uint64_t)(x ## ULL))
+
+// aus heavy.cu
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+// diese 64 Bit Rotates werden unter Compute 3.5 (und besser) mit dem Funnel Shifter beschleunigt
+#if __CUDA_ARCH__ >= 350
+__forceinline__ __device__ uint64_t ROTL64(const uint64_t value, const int offset) {
+    uint2 result;
+    if(offset >= 32) {
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
+    } else {
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
+    }
+    return  __double_as_longlong(__hiloint2double(result.y, result.x));
+}
+#else
+#define ROTL64(x, n)        (((x) << (n)) | ((x) >> (64 - (n))))
+#endif
+#define SHL(x, n)            ((x) << (n))
+#define SHR(x, n)            ((x) >> (n))
+
+#define CONST_EXP2    q[i+0] + ROTL64(q[i+1], 5)  + q[i+2] + ROTL64(q[i+3], 11) + \
+                    q[i+4] + ROTL64(q[i+5], 27) + q[i+6] + ROTL64(q[i+7], 32) + \
+                    q[i+8] + ROTL64(q[i+9], 37) + q[i+10] + ROTL64(q[i+11], 43) + \
+                    q[i+12] + ROTL64(q[i+13], 53) + (SHR(q[i+14],1) ^ q[i+14]) + (SHR(q[i+15],2) ^ q[i+15])
+
+__device__ void Compression512(uint64_t *msg, uint64_t *hash)
+{
+    // Compression ref. implementation
+    uint64_t tmp;
+    uint64_t q[32];
+
+    tmp = (msg[ 5] ^ hash[ 5]) - (msg[ 7] ^ hash[ 7]) + (msg[10] ^ hash[10]) + (msg[13] ^ hash[13]) + (msg[14] ^ hash[14]);
+    q[0] = (SHR(tmp, 1) ^ SHL(tmp, 3) ^ ROTL64(tmp,  4) ^ ROTL64(tmp, 37)) + hash[1];
+    tmp = (msg[ 6] ^ hash[ 6]) - (msg[ 8] ^ hash[ 8]) + (msg[11] ^ hash[11]) + (msg[14] ^ hash[14]) - (msg[15] ^ hash[15]);
+    q[1] = (SHR(tmp, 1) ^ SHL(tmp, 2) ^ ROTL64(tmp, 13) ^ ROTL64(tmp, 43)) + hash[2];
+    tmp = (msg[ 0] ^ hash[ 0]) + (msg[ 7] ^ hash[ 7]) + (msg[ 9] ^ hash[ 9]) - (msg[12] ^ hash[12]) + (msg[15] ^ hash[15]);
+    q[2] = (SHR(tmp, 2) ^ SHL(tmp, 1) ^ ROTL64(tmp, 19) ^ ROTL64(tmp, 53)) + hash[3];
+    tmp = (msg[ 0] ^ hash[ 0]) - (msg[ 1] ^ hash[ 1]) + (msg[ 8] ^ hash[ 8]) - (msg[10] ^ hash[10]) + (msg[13] ^ hash[13]);
+    q[3] = (SHR(tmp, 2) ^ SHL(tmp, 2) ^ ROTL64(tmp, 28) ^ ROTL64(tmp, 59)) + hash[4];
+    tmp = (msg[ 1] ^ hash[ 1]) + (msg[ 2] ^ hash[ 2]) + (msg[ 9] ^ hash[ 9]) - (msg[11] ^ hash[11]) - (msg[14] ^ hash[14]);
+    q[4] = (SHR(tmp, 1) ^ tmp) + hash[5];
+    tmp = (msg[ 3] ^ hash[ 3]) - (msg[ 2] ^ hash[ 2]) + (msg[10] ^ hash[10]) - (msg[12] ^ hash[12]) + (msg[15] ^ hash[15]);
+    q[5] = (SHR(tmp, 1) ^ SHL(tmp, 3) ^ ROTL64(tmp,  4) ^ ROTL64(tmp, 37)) + hash[6];
+    tmp = (msg[ 4] ^ hash[ 4]) - (msg[ 0] ^ hash[ 0]) - (msg[ 3] ^ hash[ 3]) - (msg[11] ^ hash[11]) + (msg[13] ^ hash[13]);
+    q[6] = (SHR(tmp, 1) ^ SHL(tmp, 2) ^ ROTL64(tmp, 13) ^ ROTL64(tmp, 43)) + hash[7];
+    tmp = (msg[ 1] ^ hash[ 1]) - (msg[ 4] ^ hash[ 4]) - (msg[ 5] ^ hash[ 5]) - (msg[12] ^ hash[12]) - (msg[14] ^ hash[14]);
+    q[7] = (SHR(tmp, 2) ^ SHL(tmp, 1) ^ ROTL64(tmp, 19) ^ ROTL64(tmp, 53)) + hash[8];
+    tmp = (msg[ 2] ^ hash[ 2]) - (msg[ 5] ^ hash[ 5]) - (msg[ 6] ^ hash[ 6]) + (msg[13] ^ hash[13]) - (msg[15] ^ hash[15]);
+    q[8] = (SHR(tmp, 2) ^ SHL(tmp, 2) ^ ROTL64(tmp, 28) ^ ROTL64(tmp, 59)) + hash[9];
+    tmp = (msg[ 0] ^ hash[ 0]) - (msg[ 3] ^ hash[ 3]) + (msg[ 6] ^ hash[ 6]) - (msg[ 7] ^ hash[ 7]) + (msg[14] ^ hash[14]);
+    q[9] = (SHR(tmp, 1) ^ tmp) + hash[10];
+    tmp = (msg[ 8] ^ hash[ 8]) - (msg[ 1] ^ hash[ 1]) - (msg[ 4] ^ hash[ 4]) - (msg[ 7] ^ hash[ 7]) + (msg[15] ^ hash[15]);
+    q[10] = (SHR(tmp, 1) ^ SHL(tmp, 3) ^ ROTL64(tmp,  4) ^ ROTL64(tmp, 37)) + hash[11];
+    tmp = (msg[ 8] ^ hash[ 8]) - (msg[ 0] ^ hash[ 0]) - (msg[ 2] ^ hash[ 2]) - (msg[ 5] ^ hash[ 5]) + (msg[ 9] ^ hash[ 9]);
+    q[11] = (SHR(tmp, 1) ^ SHL(tmp, 2) ^ ROTL64(tmp, 13) ^ ROTL64(tmp, 43)) + hash[12];
+    tmp = (msg[ 1] ^ hash[ 1]) + (msg[ 3] ^ hash[ 3]) - (msg[ 6] ^ hash[ 6]) - (msg[ 9] ^ hash[ 9]) + (msg[10] ^ hash[10]);
+    q[12] = (SHR(tmp, 2) ^ SHL(tmp, 1) ^ ROTL64(tmp, 19) ^ ROTL64(tmp, 53)) + hash[13];
+    tmp = (msg[ 2] ^ hash[ 2]) + (msg[ 4] ^ hash[ 4]) + (msg[ 7] ^ hash[ 7]) + (msg[10] ^ hash[10]) + (msg[11] ^ hash[11]);
+    q[13] = (SHR(tmp, 2) ^ SHL(tmp, 2) ^ ROTL64(tmp, 28) ^ ROTL64(tmp, 59)) + hash[14];
+    tmp = (msg[ 3] ^ hash[ 3]) - (msg[ 5] ^ hash[ 5]) + (msg[ 8] ^ hash[ 8]) - (msg[11] ^ hash[11]) - (msg[12] ^ hash[12]);
+    q[14] = (SHR(tmp, 1) ^ tmp) + hash[15];
+    tmp = (msg[12] ^ hash[12]) - (msg[ 4] ^ hash[ 4]) - (msg[ 6] ^ hash[ 6]) - (msg[ 9] ^ hash[ 9]) + (msg[13] ^ hash[13]);
+    q[15] = (SHR(tmp, 1) ^ SHL(tmp, 3) ^ ROTL64(tmp, 4) ^ ROTL64(tmp, 37)) + hash[0];
+
+    // Expand 1
+#pragma unroll 2
+    for(int i=0;i<2;i++)
+    {
+        q[i+16] =
+        (SHR(q[i], 1) ^ SHL(q[i], 2) ^ ROTL64(q[i], 13) ^ ROTL64(q[i], 43)) +
+        (SHR(q[i+1], 2) ^ SHL(q[i+1], 1) ^ ROTL64(q[i+1], 19) ^ ROTL64(q[i+1], 53)) +
+        (SHR(q[i+2], 2) ^ SHL(q[i+2], 2) ^ ROTL64(q[i+2], 28) ^ ROTL64(q[i+2], 59)) +
+        (SHR(q[i+3], 1) ^ SHL(q[i+3], 3) ^ ROTL64(q[i+3],  4) ^ ROTL64(q[i+3], 37)) +
+        (SHR(q[i+4], 1) ^ SHL(q[i+4], 2) ^ ROTL64(q[i+4], 13) ^ ROTL64(q[i+4], 43)) +
+        (SHR(q[i+5], 2) ^ SHL(q[i+5], 1) ^ ROTL64(q[i+5], 19) ^ ROTL64(q[i+5], 53)) +
+        (SHR(q[i+6], 2) ^ SHL(q[i+6], 2) ^ ROTL64(q[i+6], 28) ^ ROTL64(q[i+6], 59)) +
+        (SHR(q[i+7], 1) ^ SHL(q[i+7], 3) ^ ROTL64(q[i+7],  4) ^ ROTL64(q[i+7], 37)) +
+        (SHR(q[i+8], 1) ^ SHL(q[i+8], 2) ^ ROTL64(q[i+8], 13) ^ ROTL64(q[i+8], 43)) +
+        (SHR(q[i+9], 2) ^ SHL(q[i+9], 1) ^ ROTL64(q[i+9], 19) ^ ROTL64(q[i+9], 53)) +
+        (SHR(q[i+10], 2) ^ SHL(q[i+10], 2) ^ ROTL64(q[i+10], 28) ^ ROTL64(q[i+10], 59)) +
+        (SHR(q[i+11], 1) ^ SHL(q[i+11], 3) ^ ROTL64(q[i+11],  4) ^ ROTL64(q[i+11], 37)) +
+        (SHR(q[i+12], 1) ^ SHL(q[i+12], 2) ^ ROTL64(q[i+12], 13) ^ ROTL64(q[i+12], 43)) +
+        (SHR(q[i+13], 2) ^ SHL(q[i+13], 1) ^ ROTL64(q[i+13], 19) ^ ROTL64(q[i+13], 53)) +
+        (SHR(q[i+14], 2) ^ SHL(q[i+14], 2) ^ ROTL64(q[i+14], 28) ^ ROTL64(q[i+14], 59)) +
+        (SHR(q[i+15], 1) ^ SHL(q[i+15], 3) ^ ROTL64(q[i+15],  4) ^ ROTL64(q[i+15], 37)) +
+        ((    ((i+16)*(0x0555555555555555ull)) + ROTL64(msg[i], i+1) +
+            ROTL64(msg[i+3], i+4) - ROTL64(msg[i+10], i+11) ) ^ hash[i+7]);
+    }
+
+#pragma unroll 4
+    for(int i=2;i<6;i++) {
+        q[i+16] = CONST_EXP2 + 
+        ((    ((i+16)*(0x0555555555555555ull)) + ROTL64(msg[i], i+1) +
+            ROTL64(msg[i+3], i+4) - ROTL64(msg[i+10], i+11) ) ^ hash[i+7]);
+    }
+#pragma unroll 3
+    for(int i=6;i<9;i++) {
+        q[i+16] = CONST_EXP2 + 
+        ((    ((i+16)*(0x0555555555555555ull)) + ROTL64(msg[i], i+1) +
+            ROTL64(msg[i+3], i+4) - ROTL64(msg[i-6], (i-6)+1) ) ^ hash[i+7]);
+    }
+#pragma unroll 4
+    for(int i=9;i<13;i++) {
+        q[i+16] = CONST_EXP2 + 
+        ((    ((i+16)*(0x0555555555555555ull)) + ROTL64(msg[i], i+1) +
+            ROTL64(msg[i+3], i+4) - ROTL64(msg[i-6], (i-6)+1) ) ^ hash[i-9]);
+    }
+#pragma unroll 3
+    for(int i=13;i<16;i++) {
+        q[i+16] = CONST_EXP2 + 
+        ((    ((i+16)*(0x0555555555555555ull)) + ROTL64(msg[i], i+1) +
+            ROTL64(msg[i-13], (i-13)+1) - ROTL64(msg[i-6], (i-6)+1) ) ^ hash[i-9]);
+    }
+
+    uint64_t XL64 = q[16]^q[17]^q[18]^q[19]^q[20]^q[21]^q[22]^q[23];
+    uint64_t XH64 = XL64^q[24]^q[25]^q[26]^q[27]^q[28]^q[29]^q[30]^q[31];
+
+    hash[0] =                       (SHL(XH64, 5) ^ SHR(q[16],5) ^ msg[ 0]) + (    XL64    ^ q[24] ^ q[ 0]);
+    hash[1] =                       (SHR(XH64, 7) ^ SHL(q[17],8) ^ msg[ 1]) + (    XL64    ^ q[25] ^ q[ 1]);
+    hash[2] =                       (SHR(XH64, 5) ^ SHL(q[18],5) ^ msg[ 2]) + (    XL64    ^ q[26] ^ q[ 2]);
+    hash[3] =                       (SHR(XH64, 1) ^ SHL(q[19],5) ^ msg[ 3]) + (    XL64    ^ q[27] ^ q[ 3]);
+    hash[4] =                       (SHR(XH64, 3) ^     q[20]    ^ msg[ 4]) + (    XL64    ^ q[28] ^ q[ 4]);
+    hash[5] =                       (SHL(XH64, 6) ^ SHR(q[21],6) ^ msg[ 5]) + (    XL64    ^ q[29] ^ q[ 5]);
+    hash[6] =                       (SHR(XH64, 4) ^ SHL(q[22],6) ^ msg[ 6]) + (    XL64    ^ q[30] ^ q[ 6]);
+    hash[7] =                       (SHR(XH64,11) ^ SHL(q[23],2) ^ msg[ 7]) + (    XL64    ^ q[31] ^ q[ 7]);
+
+    hash[ 8] = ROTL64(hash[4], 9) + (    XH64     ^     q[24]    ^ msg[ 8]) + (SHL(XL64,8) ^ q[23] ^ q[ 8]);
+    hash[ 9] = ROTL64(hash[5],10) + (    XH64     ^     q[25]    ^ msg[ 9]) + (SHR(XL64,6) ^ q[16] ^ q[ 9]);
+    hash[10] = ROTL64(hash[6],11) + (    XH64     ^     q[26]    ^ msg[10]) + (SHL(XL64,6) ^ q[17] ^ q[10]);
+    hash[11] = ROTL64(hash[7],12) + (    XH64     ^     q[27]    ^ msg[11]) + (SHL(XL64,4) ^ q[18] ^ q[11]);
+    hash[12] = ROTL64(hash[0],13) + (    XH64     ^     q[28]    ^ msg[12]) + (SHR(XL64,3) ^ q[19] ^ q[12]);
+    hash[13] = ROTL64(hash[1],14) + (    XH64     ^     q[29]    ^ msg[13]) + (SHR(XL64,4) ^ q[20] ^ q[13]);
+    hash[14] = ROTL64(hash[2],15) + (    XH64     ^     q[30]    ^ msg[14]) + (SHR(XL64,7) ^ q[21] ^ q[14]);
+    hash[15] = ROTL64(hash[3],16) + (    XH64     ^     q[31]    ^ msg[15]) + (SHR(XL64,2) ^ q[22] ^ q[15]);
+}
+static __constant__ uint64_t d_constMem[16];
+static uint64_t h_constMem[16] = {
+	SPH_C64(0x8081828384858687),
+    SPH_C64(0x88898A8B8C8D8E8F),
+    SPH_C64(0x9091929394959697),
+    SPH_C64(0x98999A9B9C9D9E9F),
+    SPH_C64(0xA0A1A2A3A4A5A6A7),
+    SPH_C64(0xA8A9AAABACADAEAF),
+    SPH_C64(0xB0B1B2B3B4B5B6B7),
+    SPH_C64(0xB8B9BABBBCBDBEBF),
+    SPH_C64(0xC0C1C2C3C4C5C6C7),
+    SPH_C64(0xC8C9CACBCCCDCECF),
+    SPH_C64(0xD0D1D2D3D4D5D6D7),
+    SPH_C64(0xD8D9DADBDCDDDEDF),
+    SPH_C64(0xE0E1E2E3E4E5E6E7),
+    SPH_C64(0xE8E9EAEBECEDEEEF),
+    SPH_C64(0xF0F1F2F3F4F5F6F7),
+    SPH_C64(0xF8F9FAFBFCFDFEFF)
+};
+
+__global__ void quark_bmw512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
+{
+    int thread = (blockDim.x * blockIdx.x + threadIdx.x);
+    if (thread < threads)
+    {
+        uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
+
+        int hashPosition = nounce - startNounce;
+        uint64_t *inpHash = &g_hash[8 * hashPosition];
+
+        // Init
+        uint64_t h[16];
+		/*
+        h[ 0] = SPH_C64(0x8081828384858687);
+        h[ 1] = SPH_C64(0x88898A8B8C8D8E8F);
+        h[ 2] = SPH_C64(0x9091929394959697);
+        h[ 3] = SPH_C64(0x98999A9B9C9D9E9F);
+        h[ 4] = SPH_C64(0xA0A1A2A3A4A5A6A7);
+        h[ 5] = SPH_C64(0xA8A9AAABACADAEAF);
+        h[ 6] = SPH_C64(0xB0B1B2B3B4B5B6B7);
+        h[ 7] = SPH_C64(0xB8B9BABBBCBDBEBF);
+        h[ 8] = SPH_C64(0xC0C1C2C3C4C5C6C7);
+        h[ 9] = SPH_C64(0xC8C9CACBCCCDCECF);
+        h[10] = SPH_C64(0xD0D1D2D3D4D5D6D7);
+        h[11] = SPH_C64(0xD8D9DADBDCDDDEDF);
+        h[12] = SPH_C64(0xE0E1E2E3E4E5E6E7);
+        h[13] = SPH_C64(0xE8E9EAEBECEDEEEF);
+        h[14] = SPH_C64(0xF0F1F2F3F4F5F6F7);
+        h[15] = SPH_C64(0xF8F9FAFBFCFDFEFF);
+		*/
+#pragma unroll 16
+		for(int i=0;i<16;i++)
+			h[i] = d_constMem[i];
+        // Nachricht kopieren (Achtung, die Nachricht hat 64 Byte,
+        // BMW arbeitet mit 128 Byte!!!
+        uint64_t message[16];
+#pragma unroll 8
+        for(int i=0;i<8;i++)
+            message[i] = inpHash[i];
+#pragma unroll 6
+        for(int i=9;i<15;i++)
+            message[i] = 0;
+
+        // Padding einfügen (Byteorder?!?)
+        message[8] = SPH_C64(0x80);
+        // Länge (in Bits, d.h. 64 Byte * 8 = 512 Bits
+        message[15] = SPH_C64(512);
+
+        // Compression 1
+        Compression512(message, h);
+
+        // Final
+#pragma unroll 16
+        for(int i=0;i<16;i++)
+            message[i] = 0xaaaaaaaaaaaaaaa0ull + (uint64_t)i;
+
+        Compression512(h, message);
+
+        // fertig
+        uint64_t *outpHash = &g_hash[8 * hashPosition];
+
+#pragma unroll 8
+        for(int i=0;i<8;i++)
+            outpHash[i] = message[i+8];
+    }
+}
+
+__global__ void quark_bmw512_gpu_hash_80(int threads, uint32_t startNounce, uint64_t *g_hash)
+{
+    int thread = (blockDim.x * blockIdx.x + threadIdx.x);
+    if (thread < threads)
+    {
+        uint32_t nounce = startNounce + thread;
+
+        // Init
+        uint64_t h[16];
+#pragma unroll 16
+		for(int i=0;i<16;i++)
+			h[i] = d_constMem[i];
+
+        // Nachricht kopieren (Achtung, die Nachricht hat 64 Byte,
+        // BMW arbeitet mit 128 Byte!!!
+        uint64_t message[16];
+#pragma unroll 16
+        for(int i=0;i<16;i++)
+            message[i] = c_PaddedMessage80[i];
+
+        // die Nounce durch die thread-spezifische ersetzen
+        message[9] = REPLACE_HIWORD(message[9], cuda_swab32(nounce));
+
+        // Compression 1
+        Compression512(message, h);
+
+        // Final
+#pragma unroll 16
+        for(int i=0;i<16;i++)
+            message[i] = 0xaaaaaaaaaaaaaaa0ull + (uint64_t)i;
+
+        Compression512(h, message);
+
+        // fertig
+        uint64_t *outpHash = &g_hash[8 * thread];
+
+#pragma unroll 8
+        for(int i=0;i<8;i++)
+            outpHash[i] = message[i+8];
+    }
+}
+
+// Setup-Funktionen
+__host__ void quark_bmw512_cpu_init(int thr_id, int threads)
+{
+    // nix zu tun ;-)
+	// jetzt schon :D
+	cudaMemcpyToSymbol( d_constMem,
+                        h_constMem,
+                        sizeof(h_constMem),
+                        0, cudaMemcpyHostToDevice);
+}
+
+// Bmw512 für 80 Byte grosse Eingangsdaten
+__host__ void quark_bmw512_cpu_setBlock_80(void *pdata)
+{
+	// Message mit Padding bereitstellen
+	// lediglich die korrekte Nonce ist noch ab Byte 76 einzusetzen.
+	unsigned char PaddedMessage[128];
+	memcpy(PaddedMessage, pdata, 80);
+	memset(PaddedMessage+80, 0, 48);
+	uint64_t *message = (uint64_t*)PaddedMessage;
+	// Padding einfügen (Byteorder?!?)
+	message[10] = SPH_C64(0x80);
+	// Länge (in Bits, d.h. 80 Byte * 8 = 640 Bits
+	message[15] = SPH_C64(640);
+
+	// die Message zur Berechnung auf der GPU
+	cudaMemcpyToSymbol( c_PaddedMessage80, PaddedMessage, 16*sizeof(uint64_t), 0, cudaMemcpyHostToDevice);
+}
+
+__host__ void quark_bmw512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
+{
+    const int threadsperblock = 256;
+
+    // berechne wie viele Thread Blocks wir brauchen
+    dim3 grid((threads + threadsperblock-1)/threadsperblock);
+    dim3 block(threadsperblock);
+
+    // Größe des dynamischen Shared Memory Bereichs
+    size_t shared_size = 0;
+
+//    fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
+
+    quark_bmw512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
+    MyStreamSynchronize(NULL, order, thr_id);
+}
+
+__host__ void quark_bmw512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_hash, int order)
+{
+    const int threadsperblock = 256;
+
+    // berechne wie viele Thread Blocks wir brauchen
+    dim3 grid((threads + threadsperblock-1)/threadsperblock);
+    dim3 block(threadsperblock);
+
+    // Größe des dynamischen Shared Memory Bereichs
+    size_t shared_size = 0;
+
+//    fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
+
+    quark_bmw512_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash);
+    MyStreamSynchronize(NULL, order, thr_id);
+}
+
+#endif
diff --git a/quark/cuda_jh512.cu b/quark/cuda_jh512.cu
index 75508cf..779af44 100644
--- a/quark/cuda_jh512.cu
+++ b/quark/cuda_jh512.cu
@@ -1,356 +1,358 @@
-#include <stdint.h>
-
-// aus heavy.cu
-extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
-
-typedef struct {
-    uint32_t x[8][4];                     /*the 1024-bit state, ( x[i][0] || x[i][1] || x[i][2] || x[i][3] ) is the ith row of the state in the pseudocode*/
-    uint32_t buffer[16];                  /*the 512-bit message block to be hashed;*/
-} hashState;
-
-/*42 round constants, each round constant is 32-byte (256-bit)*/
-__constant__ uint32_t c_INIT_bitslice[8][4];
-__constant__ unsigned char c_E8_bitslice_roundconstant[42][32];
-
-const uint32_t h_INIT_bitslice[8][4] = {
-	{ 0x964bd16f, 0x17aa003e, 0x052e6a63, 0x43d5157a},
-	{ 0x8d5e228a, 0x0bef970c, 0x591234e9, 0x61c3b3f2},
-	{ 0xc1a01d89, 0x1e806f53, 0x6b05a92a, 0x806d2bea},
-	{ 0xdbcc8e58, 0xa6ba7520, 0x763a0fa9, 0xf73bf8ba},
-	{ 0x05e66901, 0x694ae341, 0x8e8ab546, 0x5ae66f2e},
-	{ 0xd0a74710, 0x243c84c1, 0xb1716e3b, 0x99c15a2d},
-	{ 0xecf657cf, 0x56f8b19d, 0x7c8806a7, 0x56b11657},
-	{ 0xdffcc2e3, 0xfb1785e6, 0x78465a54, 0x4bdd8ccc} };
-
-const unsigned char h_E8_bitslice_roundconstant[42][32]={
-{0x72,0xd5,0xde,0xa2,0xdf,0x15,0xf8,0x67,0x7b,0x84,0x15,0xa,0xb7,0x23,0x15,0x57,0x81,0xab,0xd6,0x90,0x4d,0x5a,0x87,0xf6,0x4e,0x9f,0x4f,0xc5,0xc3,0xd1,0x2b,0x40},
-{0xea,0x98,0x3a,0xe0,0x5c,0x45,0xfa,0x9c,0x3,0xc5,0xd2,0x99,0x66,0xb2,0x99,0x9a,0x66,0x2,0x96,0xb4,0xf2,0xbb,0x53,0x8a,0xb5,0x56,0x14,0x1a,0x88,0xdb,0xa2,0x31},
-{0x3,0xa3,0x5a,0x5c,0x9a,0x19,0xe,0xdb,0x40,0x3f,0xb2,0xa,0x87,0xc1,0x44,0x10,0x1c,0x5,0x19,0x80,0x84,0x9e,0x95,0x1d,0x6f,0x33,0xeb,0xad,0x5e,0xe7,0xcd,0xdc},
-{0x10,0xba,0x13,0x92,0x2,0xbf,0x6b,0x41,0xdc,0x78,0x65,0x15,0xf7,0xbb,0x27,0xd0,0xa,0x2c,0x81,0x39,0x37,0xaa,0x78,0x50,0x3f,0x1a,0xbf,0xd2,0x41,0x0,0x91,0xd3},
-{0x42,0x2d,0x5a,0xd,0xf6,0xcc,0x7e,0x90,0xdd,0x62,0x9f,0x9c,0x92,0xc0,0x97,0xce,0x18,0x5c,0xa7,0xb,0xc7,0x2b,0x44,0xac,0xd1,0xdf,0x65,0xd6,0x63,0xc6,0xfc,0x23},
-{0x97,0x6e,0x6c,0x3,0x9e,0xe0,0xb8,0x1a,0x21,0x5,0x45,0x7e,0x44,0x6c,0xec,0xa8,0xee,0xf1,0x3,0xbb,0x5d,0x8e,0x61,0xfa,0xfd,0x96,0x97,0xb2,0x94,0x83,0x81,0x97},
-{0x4a,0x8e,0x85,0x37,0xdb,0x3,0x30,0x2f,0x2a,0x67,0x8d,0x2d,0xfb,0x9f,0x6a,0x95,0x8a,0xfe,0x73,0x81,0xf8,0xb8,0x69,0x6c,0x8a,0xc7,0x72,0x46,0xc0,0x7f,0x42,0x14},
-{0xc5,0xf4,0x15,0x8f,0xbd,0xc7,0x5e,0xc4,0x75,0x44,0x6f,0xa7,0x8f,0x11,0xbb,0x80,0x52,0xde,0x75,0xb7,0xae,0xe4,0x88,0xbc,0x82,0xb8,0x0,0x1e,0x98,0xa6,0xa3,0xf4},
-{0x8e,0xf4,0x8f,0x33,0xa9,0xa3,0x63,0x15,0xaa,0x5f,0x56,0x24,0xd5,0xb7,0xf9,0x89,0xb6,0xf1,0xed,0x20,0x7c,0x5a,0xe0,0xfd,0x36,0xca,0xe9,0x5a,0x6,0x42,0x2c,0x36},
-{0xce,0x29,0x35,0x43,0x4e,0xfe,0x98,0x3d,0x53,0x3a,0xf9,0x74,0x73,0x9a,0x4b,0xa7,0xd0,0xf5,0x1f,0x59,0x6f,0x4e,0x81,0x86,0xe,0x9d,0xad,0x81,0xaf,0xd8,0x5a,0x9f},
-{0xa7,0x5,0x6,0x67,0xee,0x34,0x62,0x6a,0x8b,0xb,0x28,0xbe,0x6e,0xb9,0x17,0x27,0x47,0x74,0x7,0x26,0xc6,0x80,0x10,0x3f,0xe0,0xa0,0x7e,0x6f,0xc6,0x7e,0x48,0x7b},
-{0xd,0x55,0xa,0xa5,0x4a,0xf8,0xa4,0xc0,0x91,0xe3,0xe7,0x9f,0x97,0x8e,0xf1,0x9e,0x86,0x76,0x72,0x81,0x50,0x60,0x8d,0xd4,0x7e,0x9e,0x5a,0x41,0xf3,0xe5,0xb0,0x62},
-{0xfc,0x9f,0x1f,0xec,0x40,0x54,0x20,0x7a,0xe3,0xe4,0x1a,0x0,0xce,0xf4,0xc9,0x84,0x4f,0xd7,0x94,0xf5,0x9d,0xfa,0x95,0xd8,0x55,0x2e,0x7e,0x11,0x24,0xc3,0x54,0xa5},
-{0x5b,0xdf,0x72,0x28,0xbd,0xfe,0x6e,0x28,0x78,0xf5,0x7f,0xe2,0xf,0xa5,0xc4,0xb2,0x5,0x89,0x7c,0xef,0xee,0x49,0xd3,0x2e,0x44,0x7e,0x93,0x85,0xeb,0x28,0x59,0x7f},
-{0x70,0x5f,0x69,0x37,0xb3,0x24,0x31,0x4a,0x5e,0x86,0x28,0xf1,0x1d,0xd6,0xe4,0x65,0xc7,0x1b,0x77,0x4,0x51,0xb9,0x20,0xe7,0x74,0xfe,0x43,0xe8,0x23,0xd4,0x87,0x8a},
-{0x7d,0x29,0xe8,0xa3,0x92,0x76,0x94,0xf2,0xdd,0xcb,0x7a,0x9,0x9b,0x30,0xd9,0xc1,0x1d,0x1b,0x30,0xfb,0x5b,0xdc,0x1b,0xe0,0xda,0x24,0x49,0x4f,0xf2,0x9c,0x82,0xbf},
-{0xa4,0xe7,0xba,0x31,0xb4,0x70,0xbf,0xff,0xd,0x32,0x44,0x5,0xde,0xf8,0xbc,0x48,0x3b,0xae,0xfc,0x32,0x53,0xbb,0xd3,0x39,0x45,0x9f,0xc3,0xc1,0xe0,0x29,0x8b,0xa0},
-{0xe5,0xc9,0x5,0xfd,0xf7,0xae,0x9,0xf,0x94,0x70,0x34,0x12,0x42,0x90,0xf1,0x34,0xa2,0x71,0xb7,0x1,0xe3,0x44,0xed,0x95,0xe9,0x3b,0x8e,0x36,0x4f,0x2f,0x98,0x4a},
-{0x88,0x40,0x1d,0x63,0xa0,0x6c,0xf6,0x15,0x47,0xc1,0x44,0x4b,0x87,0x52,0xaf,0xff,0x7e,0xbb,0x4a,0xf1,0xe2,0xa,0xc6,0x30,0x46,0x70,0xb6,0xc5,0xcc,0x6e,0x8c,0xe6},
-{0xa4,0xd5,0xa4,0x56,0xbd,0x4f,0xca,0x0,0xda,0x9d,0x84,0x4b,0xc8,0x3e,0x18,0xae,0x73,0x57,0xce,0x45,0x30,0x64,0xd1,0xad,0xe8,0xa6,0xce,0x68,0x14,0x5c,0x25,0x67},
-{0xa3,0xda,0x8c,0xf2,0xcb,0xe,0xe1,0x16,0x33,0xe9,0x6,0x58,0x9a,0x94,0x99,0x9a,0x1f,0x60,0xb2,0x20,0xc2,0x6f,0x84,0x7b,0xd1,0xce,0xac,0x7f,0xa0,0xd1,0x85,0x18},
-{0x32,0x59,0x5b,0xa1,0x8d,0xdd,0x19,0xd3,0x50,0x9a,0x1c,0xc0,0xaa,0xa5,0xb4,0x46,0x9f,0x3d,0x63,0x67,0xe4,0x4,0x6b,0xba,0xf6,0xca,0x19,0xab,0xb,0x56,0xee,0x7e},
-{0x1f,0xb1,0x79,0xea,0xa9,0x28,0x21,0x74,0xe9,0xbd,0xf7,0x35,0x3b,0x36,0x51,0xee,0x1d,0x57,0xac,0x5a,0x75,0x50,0xd3,0x76,0x3a,0x46,0xc2,0xfe,0xa3,0x7d,0x70,0x1},
-{0xf7,0x35,0xc1,0xaf,0x98,0xa4,0xd8,0x42,0x78,0xed,0xec,0x20,0x9e,0x6b,0x67,0x79,0x41,0x83,0x63,0x15,0xea,0x3a,0xdb,0xa8,0xfa,0xc3,0x3b,0x4d,0x32,0x83,0x2c,0x83},
-{0xa7,0x40,0x3b,0x1f,0x1c,0x27,0x47,0xf3,0x59,0x40,0xf0,0x34,0xb7,0x2d,0x76,0x9a,0xe7,0x3e,0x4e,0x6c,0xd2,0x21,0x4f,0xfd,0xb8,0xfd,0x8d,0x39,0xdc,0x57,0x59,0xef},
-{0x8d,0x9b,0xc,0x49,0x2b,0x49,0xeb,0xda,0x5b,0xa2,0xd7,0x49,0x68,0xf3,0x70,0xd,0x7d,0x3b,0xae,0xd0,0x7a,0x8d,0x55,0x84,0xf5,0xa5,0xe9,0xf0,0xe4,0xf8,0x8e,0x65},
-{0xa0,0xb8,0xa2,0xf4,0x36,0x10,0x3b,0x53,0xc,0xa8,0x7,0x9e,0x75,0x3e,0xec,0x5a,0x91,0x68,0x94,0x92,0x56,0xe8,0x88,0x4f,0x5b,0xb0,0x5c,0x55,0xf8,0xba,0xbc,0x4c},
-{0xe3,0xbb,0x3b,0x99,0xf3,0x87,0x94,0x7b,0x75,0xda,0xf4,0xd6,0x72,0x6b,0x1c,0x5d,0x64,0xae,0xac,0x28,0xdc,0x34,0xb3,0x6d,0x6c,0x34,0xa5,0x50,0xb8,0x28,0xdb,0x71},
-{0xf8,0x61,0xe2,0xf2,0x10,0x8d,0x51,0x2a,0xe3,0xdb,0x64,0x33,0x59,0xdd,0x75,0xfc,0x1c,0xac,0xbc,0xf1,0x43,0xce,0x3f,0xa2,0x67,0xbb,0xd1,0x3c,0x2,0xe8,0x43,0xb0},
-{0x33,0xa,0x5b,0xca,0x88,0x29,0xa1,0x75,0x7f,0x34,0x19,0x4d,0xb4,0x16,0x53,0x5c,0x92,0x3b,0x94,0xc3,0xe,0x79,0x4d,0x1e,0x79,0x74,0x75,0xd7,0xb6,0xee,0xaf,0x3f},
-{0xea,0xa8,0xd4,0xf7,0xbe,0x1a,0x39,0x21,0x5c,0xf4,0x7e,0x9,0x4c,0x23,0x27,0x51,0x26,0xa3,0x24,0x53,0xba,0x32,0x3c,0xd2,0x44,0xa3,0x17,0x4a,0x6d,0xa6,0xd5,0xad},
-{0xb5,0x1d,0x3e,0xa6,0xaf,0xf2,0xc9,0x8,0x83,0x59,0x3d,0x98,0x91,0x6b,0x3c,0x56,0x4c,0xf8,0x7c,0xa1,0x72,0x86,0x60,0x4d,0x46,0xe2,0x3e,0xcc,0x8,0x6e,0xc7,0xf6},
-{0x2f,0x98,0x33,0xb3,0xb1,0xbc,0x76,0x5e,0x2b,0xd6,0x66,0xa5,0xef,0xc4,0xe6,0x2a,0x6,0xf4,0xb6,0xe8,0xbe,0xc1,0xd4,0x36,0x74,0xee,0x82,0x15,0xbc,0xef,0x21,0x63},
-{0xfd,0xc1,0x4e,0xd,0xf4,0x53,0xc9,0x69,0xa7,0x7d,0x5a,0xc4,0x6,0x58,0x58,0x26,0x7e,0xc1,0x14,0x16,0x6,0xe0,0xfa,0x16,0x7e,0x90,0xaf,0x3d,0x28,0x63,0x9d,0x3f},
-{0xd2,0xc9,0xf2,0xe3,0x0,0x9b,0xd2,0xc,0x5f,0xaa,0xce,0x30,0xb7,0xd4,0xc,0x30,0x74,0x2a,0x51,0x16,0xf2,0xe0,0x32,0x98,0xd,0xeb,0x30,0xd8,0xe3,0xce,0xf8,0x9a},
-{0x4b,0xc5,0x9e,0x7b,0xb5,0xf1,0x79,0x92,0xff,0x51,0xe6,0x6e,0x4,0x86,0x68,0xd3,0x9b,0x23,0x4d,0x57,0xe6,0x96,0x67,0x31,0xcc,0xe6,0xa6,0xf3,0x17,0xa,0x75,0x5},
-{0xb1,0x76,0x81,0xd9,0x13,0x32,0x6c,0xce,0x3c,0x17,0x52,0x84,0xf8,0x5,0xa2,0x62,0xf4,0x2b,0xcb,0xb3,0x78,0x47,0x15,0x47,0xff,0x46,0x54,0x82,0x23,0x93,0x6a,0x48},
-{0x38,0xdf,0x58,0x7,0x4e,0x5e,0x65,0x65,0xf2,0xfc,0x7c,0x89,0xfc,0x86,0x50,0x8e,0x31,0x70,0x2e,0x44,0xd0,0xb,0xca,0x86,0xf0,0x40,0x9,0xa2,0x30,0x78,0x47,0x4e},
-{0x65,0xa0,0xee,0x39,0xd1,0xf7,0x38,0x83,0xf7,0x5e,0xe9,0x37,0xe4,0x2c,0x3a,0xbd,0x21,0x97,0xb2,0x26,0x1,0x13,0xf8,0x6f,0xa3,0x44,0xed,0xd1,0xef,0x9f,0xde,0xe7},
-{0x8b,0xa0,0xdf,0x15,0x76,0x25,0x92,0xd9,0x3c,0x85,0xf7,0xf6,0x12,0xdc,0x42,0xbe,0xd8,0xa7,0xec,0x7c,0xab,0x27,0xb0,0x7e,0x53,0x8d,0x7d,0xda,0xaa,0x3e,0xa8,0xde},
-{0xaa,0x25,0xce,0x93,0xbd,0x2,0x69,0xd8,0x5a,0xf6,0x43,0xfd,0x1a,0x73,0x8,0xf9,0xc0,0x5f,0xef,0xda,0x17,0x4a,0x19,0xa5,0x97,0x4d,0x66,0x33,0x4c,0xfd,0x21,0x6a},
-{0x35,0xb4,0x98,0x31,0xdb,0x41,0x15,0x70,0xea,0x1e,0xf,0xbb,0xed,0xcd,0x54,0x9b,0x9a,0xd0,0x63,0xa1,0x51,0x97,0x40,0x72,0xf6,0x75,0x9d,0xbf,0x91,0x47,0x6f,0xe2}};
-
-/*swapping bit 2i with bit 2i+1 of 32-bit x*/
-#define SWAP1(x)   (x) = ((((x) & 0x55555555UL) << 1) | (((x) & 0xaaaaaaaaUL) >> 1));
-/*swapping bits 4i||4i+1 with bits 4i+2||4i+3 of 32-bit x*/
-#define SWAP2(x)   (x) = ((((x) & 0x33333333UL) << 2) | (((x) & 0xccccccccUL) >> 2));
-/*swapping bits 8i||8i+1||8i+2||8i+3 with bits 8i+4||8i+5||8i+6||8i+7 of 32-bit x*/
-#define SWAP4(x)   (x) = ((((x) & 0x0f0f0f0fUL) << 4) | (((x) & 0xf0f0f0f0UL) >> 4));
-/*swapping bits 16i||16i+1||......||16i+7  with bits 16i+8||16i+9||......||16i+15 of 32-bit x*/
-#define SWAP8(x)   (x) = ((((x) & 0x00ff00ffUL) << 8) | (((x) & 0xff00ff00UL) >> 8));
-/*swapping bits 32i||32i+1||......||32i+15 with bits 32i+16||32i+17||......||32i+31 of 32-bit x*/
-#define SWAP16(x)  (x) = ((((x) & 0x0000ffffUL) << 16) | (((x) & 0xffff0000UL) >> 16));
-
-/*The MDS transform*/
-#define L(m0,m1,m2,m3,m4,m5,m6,m7) \
-      (m4) ^= (m1);                \
-      (m5) ^= (m2);                \
-      (m6) ^= (m0) ^ (m3);         \
-      (m7) ^= (m0);                \
-      (m0) ^= (m5);                \
-      (m1) ^= (m6);                \
-      (m2) ^= (m4) ^ (m7);         \
-      (m3) ^= (m4);
-
-/*The Sbox*/
-#define Sbox(m0,m1,m2,m3,cc)       \
-      m3  = ~(m3);                 \
-      m0 ^= ((~(m2)) & (cc));      \
-      temp0 = (cc) ^ ((m0) & (m1));\
-      m0 ^= ((m2) & (m3));         \
-      m3 ^= ((~(m1)) & (m2));      \
-      m1 ^= ((m0) & (m2));         \
-      m2 ^= ((m0) & (~(m3)));      \
-      m0 ^= ((m1) | (m3));         \
-      m3 ^= ((m1) & (m2));         \
-      m1 ^= (temp0 & (m0));        \
-      m2 ^= temp0;
-
-__device__ __forceinline__ void Sbox_and_MDS_layer(hashState* state, uint32_t roundnumber)
-{
-    uint32_t temp0;
-	uint32_t cc0, cc1;
-    //Sbox and MDS layer
-#pragma unroll 4
-    for (int i = 0; i < 4; i++) {
-		cc0 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i];
-		cc1 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i+4];
-        Sbox(state->x[0][i],state->x[2][i], state->x[4][i], state->x[6][i], cc0);
-        Sbox(state->x[1][i],state->x[3][i], state->x[5][i], state->x[7][i], cc1);
-        L(state->x[0][i],state->x[2][i],state->x[4][i],state->x[6][i],state->x[1][i],state->x[3][i],state->x[5][i],state->x[7][i]);
-    }
-}
-
-__device__ __forceinline__ void RoundFunction0(hashState* state, uint32_t roundnumber)
-{
-	Sbox_and_MDS_layer(state, roundnumber);
-
-#pragma unroll 4
-	for (int j = 1; j < 8; j = j+2)
-	{
-#pragma unroll 4
-		for (int i = 0; i < 4; i++) SWAP1(state->x[j][i]);
-	}
-}
-
-__device__ __forceinline__ void RoundFunction1(hashState* state, uint32_t roundnumber)
-{
-	Sbox_and_MDS_layer(state, roundnumber);
-
-#pragma unroll 4
-	for (int j = 1; j < 8; j = j+2)
-	{
-#pragma unroll 4
-		for (int i = 0; i < 4; i++) SWAP2(state->x[j][i]);
-	}
-}
-
-__device__ __forceinline__ void RoundFunction2(hashState* state, uint32_t roundnumber)
-{
-	Sbox_and_MDS_layer(state, roundnumber);
-
-#pragma unroll 4
-	for (int j = 1; j < 8; j = j+2)
-	{
-#pragma unroll 4
-		for (int i = 0; i < 4; i++) SWAP4(state->x[j][i]);
-	}
-}
-
-__device__ __forceinline__ void RoundFunction3(hashState* state, uint32_t roundnumber)
-{
-	Sbox_and_MDS_layer(state, roundnumber);
-
-#pragma unroll 4
-	for (int j = 1; j < 8; j = j+2)
-	{
-#pragma unroll 4
-		for (int i = 0; i < 4; i++) SWAP8(state->x[j][i]);
-	}
-}
-
-__device__ __forceinline__ void RoundFunction4(hashState* state, uint32_t roundnumber)
-{
-	Sbox_and_MDS_layer(state, roundnumber);
-
-#pragma unroll 4
-	for (int j = 1; j < 8; j = j+2)
-	{
-#pragma unroll 4
-		for (int i = 0; i < 4; i++) SWAP16(state->x[j][i]);
-	}
-}
-
-__device__ __forceinline__ void RoundFunction5(hashState* state, uint32_t roundnumber)
-{
-	uint32_t temp0;
-
-	Sbox_and_MDS_layer(state, roundnumber);
-
-#pragma unroll 4
-	for (int j = 1; j < 8; j = j+2)
-	{
-#pragma unroll 2
-		for (int i = 0; i < 4; i = i+2) {
-			temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+1]; state->x[j][i+1] = temp0;
-		}
-	}
-}
-
-__device__ __forceinline__ void RoundFunction6(hashState* state, uint32_t roundnumber)
-{
-	uint32_t temp0;
-
-	Sbox_and_MDS_layer(state, roundnumber);
-
-#pragma unroll 4
-	for (int j = 1; j < 8; j = j+2)
-	{
-#pragma unroll 2
-		for (int i = 0; i < 2; i++) {
-			temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+2]; state->x[j][i+2] = temp0;
-		}
-	}
-}
-
-/*The bijective function E8, in bitslice form */
-__device__ __forceinline__ void E8(hashState *state)
-{
-    /*perform 6 rounds*/
-//#pragma unroll 6
-    for (int i = 0; i < 42; i+=7)
-	{
-		RoundFunction0(state, i);
-		RoundFunction1(state, i+1);
-		RoundFunction2(state, i+2);
-		RoundFunction3(state, i+3);
-		RoundFunction4(state, i+4);
-		RoundFunction5(state, i+5);
-		RoundFunction6(state, i+6);
-	}
-}
-
-/*The compression function F8 */
-__device__ __forceinline__ void F8(hashState *state)
-{
-    /*xor the 512-bit message with the fist half of the 1024-bit hash state*/
-#pragma unroll 16
-    for (int i = 0; i < 16; i++)  state->x[i >> 2][i & 3] ^= ((uint32_t*)state->buffer)[i];
-
-    /*the bijective function E8 */
-    E8(state);
-
-    /*xor the 512-bit message with the second half of the 1024-bit hash state*/
-#pragma unroll 16
-    for (int i = 0; i < 16; i++)  state->x[(16+i) >> 2][(16+i) & 3] ^= ((uint32_t*)state->buffer)[i];
-}
-
-
-__device__ __forceinline__ void JHHash(const uint32_t *data, uint32_t *hashval)
-{
-    hashState state;
-
-    /*load the intital hash value H0 into state*/
-	/*
-    #define INIT(a,b,c,d) ((a) | ((b)<<8) | ((c)<<16) | ((d)<<24))
-    state.x[0][0] = INIT(0x6f,0xd1,0x4b,0x96);
-    state.x[0][1] = INIT(0x3e,0x00,0xaa,0x17);
-    state.x[0][2] = INIT(0x63,0x6a,0x2e,0x05);
-    state.x[0][3] = INIT(0x7a,0x15,0xd5,0x43);
-    state.x[1][0] = INIT(0x8a,0x22,0x5e,0x8d);
-    state.x[1][1] = INIT(0x0c,0x97,0xef,0x0b);
-    state.x[1][2] = INIT(0xe9,0x34,0x12,0x59);
-    state.x[1][3] = INIT(0xf2,0xb3,0xc3,0x61);
-    state.x[2][0] = INIT(0x89,0x1d,0xa0,0xc1);
-    state.x[2][1] = INIT(0x53,0x6f,0x80,0x1e);
-    state.x[2][2] = INIT(0x2a,0xa9,0x05,0x6b);
-    state.x[2][3] = INIT(0xea,0x2b,0x6d,0x80);
-    state.x[3][0] = INIT(0x58,0x8e,0xcc,0xdb);
-    state.x[3][1] = INIT(0x20,0x75,0xba,0xa6);
-    state.x[3][2] = INIT(0xa9,0x0f,0x3a,0x76);
-    state.x[3][3] = INIT(0xba,0xf8,0x3b,0xf7);
-    state.x[4][0] = INIT(0x01,0x69,0xe6,0x05);
-    state.x[4][1] = INIT(0x41,0xe3,0x4a,0x69);
-    state.x[4][2] = INIT(0x46,0xb5,0x8a,0x8e);
-    state.x[4][3] = INIT(0x2e,0x6f,0xe6,0x5a);
-    state.x[5][0] = INIT(0x10,0x47,0xa7,0xd0);
-    state.x[5][1] = INIT(0xc1,0x84,0x3c,0x24);
-    state.x[5][2] = INIT(0x3b,0x6e,0x71,0xb1);
-    state.x[5][3] = INIT(0x2d,0x5a,0xc1,0x99);
-    state.x[6][0] = INIT(0xcf,0x57,0xf6,0xec);
-    state.x[6][1] = INIT(0x9d,0xb1,0xf8,0x56);
-    state.x[6][2] = INIT(0xa7,0x06,0x88,0x7c);
-    state.x[6][3] = INIT(0x57,0x16,0xb1,0x56);
-    state.x[7][0] = INIT(0xe3,0xc2,0xfc,0xdf);
-    state.x[7][1] = INIT(0xe6,0x85,0x17,0xfb);
-    state.x[7][2] = INIT(0x54,0x5a,0x46,0x78);
-    state.x[7][3] = INIT(0xcc,0x8c,0xdd,0x4b);
-	*/
-#pragma unroll 8
-	for(int j=0;j<8;j++)
-	{
-#pragma unroll 4
-		for(int i=0;i<4;i++)
-			state.x[j][i] = c_INIT_bitslice[j][i];
-	}
-
-#pragma unroll 16
-    for (int i=0; i < 16; ++i) state.buffer[i] = data[i];
-    F8(&state);
-
-    /*pad the message when databitlen is multiple of 512 bits, then process the padded block*/
-    state.buffer[0] = 0x80;
-#pragma unroll 14
-    for (int i=1; i < 15; i++) state.buffer[i] = 0;
-    state.buffer[15] = 0x00020000;
-    F8(&state);
-
-    /*truncating the final hash value to generate the message digest*/
-#pragma unroll 16
-    for (int i=0; i < 16; ++i) hashval[i] = state.x[4][i];
-}
-
-// Die Hash-Funktion
-__global__ void quark_jh512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
-{
-    int thread = (blockDim.x * blockIdx.x + threadIdx.x);
-    if (thread < threads)
-    {
-        uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
-
-        int hashPosition = nounce - startNounce;
-        uint32_t *Hash = (uint32_t*)&g_hash[8 * hashPosition];
-
-        JHHash(Hash, Hash);
-    }
-}
-
-
-// Setup-Funktionen
-__host__ void quark_jh512_cpu_init(int thr_id, int threads)
-{
-	
-    cudaMemcpyToSymbol( c_E8_bitslice_roundconstant,
-                        h_E8_bitslice_roundconstant,
-                        sizeof(h_E8_bitslice_roundconstant),
-                        0, cudaMemcpyHostToDevice);
-
-	cudaMemcpyToSymbol( c_INIT_bitslice,
-                        h_INIT_bitslice,
-                        sizeof(h_INIT_bitslice),
-                        0, cudaMemcpyHostToDevice);
-}
-
-__host__ void quark_jh512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
-{
-    const int threadsperblock = 256;
-
-    // berechne wie viele Thread Blocks wir brauchen
-    dim3 grid((threads + threadsperblock-1)/threadsperblock);
-    dim3 block(threadsperblock);
-
-    // Größe des dynamischen Shared Memory Bereichs
-    size_t shared_size = 0;
-
-//    fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
-
-    quark_jh512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
-    MyStreamSynchronize(NULL, order, thr_id);
-}
-
+#include <stdint.h>
+
+// aus heavy.cu
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+typedef struct {
+    uint32_t x[8][4];                     /*the 1024-bit state, ( x[i][0] || x[i][1] || x[i][2] || x[i][3] ) is the ith row of the state in the pseudocode*/
+    uint32_t buffer[16];                  /*the 512-bit message block to be hashed;*/
+} hashState;
+
+/*42 round constants, each round constant is 32-byte (256-bit)*/
+__constant__ uint32_t c_INIT_bitslice[8][4];
+__constant__ unsigned char c_E8_bitslice_roundconstant[42][32];
+
+const uint32_t h_INIT_bitslice[8][4] = {
+	{ 0x964bd16f, 0x17aa003e, 0x052e6a63, 0x43d5157a},
+	{ 0x8d5e228a, 0x0bef970c, 0x591234e9, 0x61c3b3f2},
+	{ 0xc1a01d89, 0x1e806f53, 0x6b05a92a, 0x806d2bea},
+	{ 0xdbcc8e58, 0xa6ba7520, 0x763a0fa9, 0xf73bf8ba},
+	{ 0x05e66901, 0x694ae341, 0x8e8ab546, 0x5ae66f2e},
+	{ 0xd0a74710, 0x243c84c1, 0xb1716e3b, 0x99c15a2d},
+	{ 0xecf657cf, 0x56f8b19d, 0x7c8806a7, 0x56b11657},
+	{ 0xdffcc2e3, 0xfb1785e6, 0x78465a54, 0x4bdd8ccc} };
+
+const unsigned char h_E8_bitslice_roundconstant[42][32]={
+{0x72,0xd5,0xde,0xa2,0xdf,0x15,0xf8,0x67,0x7b,0x84,0x15,0xa,0xb7,0x23,0x15,0x57,0x81,0xab,0xd6,0x90,0x4d,0x5a,0x87,0xf6,0x4e,0x9f,0x4f,0xc5,0xc3,0xd1,0x2b,0x40},
+{0xea,0x98,0x3a,0xe0,0x5c,0x45,0xfa,0x9c,0x3,0xc5,0xd2,0x99,0x66,0xb2,0x99,0x9a,0x66,0x2,0x96,0xb4,0xf2,0xbb,0x53,0x8a,0xb5,0x56,0x14,0x1a,0x88,0xdb,0xa2,0x31},
+{0x3,0xa3,0x5a,0x5c,0x9a,0x19,0xe,0xdb,0x40,0x3f,0xb2,0xa,0x87,0xc1,0x44,0x10,0x1c,0x5,0x19,0x80,0x84,0x9e,0x95,0x1d,0x6f,0x33,0xeb,0xad,0x5e,0xe7,0xcd,0xdc},
+{0x10,0xba,0x13,0x92,0x2,0xbf,0x6b,0x41,0xdc,0x78,0x65,0x15,0xf7,0xbb,0x27,0xd0,0xa,0x2c,0x81,0x39,0x37,0xaa,0x78,0x50,0x3f,0x1a,0xbf,0xd2,0x41,0x0,0x91,0xd3},
+{0x42,0x2d,0x5a,0xd,0xf6,0xcc,0x7e,0x90,0xdd,0x62,0x9f,0x9c,0x92,0xc0,0x97,0xce,0x18,0x5c,0xa7,0xb,0xc7,0x2b,0x44,0xac,0xd1,0xdf,0x65,0xd6,0x63,0xc6,0xfc,0x23},
+{0x97,0x6e,0x6c,0x3,0x9e,0xe0,0xb8,0x1a,0x21,0x5,0x45,0x7e,0x44,0x6c,0xec,0xa8,0xee,0xf1,0x3,0xbb,0x5d,0x8e,0x61,0xfa,0xfd,0x96,0x97,0xb2,0x94,0x83,0x81,0x97},
+{0x4a,0x8e,0x85,0x37,0xdb,0x3,0x30,0x2f,0x2a,0x67,0x8d,0x2d,0xfb,0x9f,0x6a,0x95,0x8a,0xfe,0x73,0x81,0xf8,0xb8,0x69,0x6c,0x8a,0xc7,0x72,0x46,0xc0,0x7f,0x42,0x14},
+{0xc5,0xf4,0x15,0x8f,0xbd,0xc7,0x5e,0xc4,0x75,0x44,0x6f,0xa7,0x8f,0x11,0xbb,0x80,0x52,0xde,0x75,0xb7,0xae,0xe4,0x88,0xbc,0x82,0xb8,0x0,0x1e,0x98,0xa6,0xa3,0xf4},
+{0x8e,0xf4,0x8f,0x33,0xa9,0xa3,0x63,0x15,0xaa,0x5f,0x56,0x24,0xd5,0xb7,0xf9,0x89,0xb6,0xf1,0xed,0x20,0x7c,0x5a,0xe0,0xfd,0x36,0xca,0xe9,0x5a,0x6,0x42,0x2c,0x36},
+{0xce,0x29,0x35,0x43,0x4e,0xfe,0x98,0x3d,0x53,0x3a,0xf9,0x74,0x73,0x9a,0x4b,0xa7,0xd0,0xf5,0x1f,0x59,0x6f,0x4e,0x81,0x86,0xe,0x9d,0xad,0x81,0xaf,0xd8,0x5a,0x9f},
+{0xa7,0x5,0x6,0x67,0xee,0x34,0x62,0x6a,0x8b,0xb,0x28,0xbe,0x6e,0xb9,0x17,0x27,0x47,0x74,0x7,0x26,0xc6,0x80,0x10,0x3f,0xe0,0xa0,0x7e,0x6f,0xc6,0x7e,0x48,0x7b},
+{0xd,0x55,0xa,0xa5,0x4a,0xf8,0xa4,0xc0,0x91,0xe3,0xe7,0x9f,0x97,0x8e,0xf1,0x9e,0x86,0x76,0x72,0x81,0x50,0x60,0x8d,0xd4,0x7e,0x9e,0x5a,0x41,0xf3,0xe5,0xb0,0x62},
+{0xfc,0x9f,0x1f,0xec,0x40,0x54,0x20,0x7a,0xe3,0xe4,0x1a,0x0,0xce,0xf4,0xc9,0x84,0x4f,0xd7,0x94,0xf5,0x9d,0xfa,0x95,0xd8,0x55,0x2e,0x7e,0x11,0x24,0xc3,0x54,0xa5},
+{0x5b,0xdf,0x72,0x28,0xbd,0xfe,0x6e,0x28,0x78,0xf5,0x7f,0xe2,0xf,0xa5,0xc4,0xb2,0x5,0x89,0x7c,0xef,0xee,0x49,0xd3,0x2e,0x44,0x7e,0x93,0x85,0xeb,0x28,0x59,0x7f},
+{0x70,0x5f,0x69,0x37,0xb3,0x24,0x31,0x4a,0x5e,0x86,0x28,0xf1,0x1d,0xd6,0xe4,0x65,0xc7,0x1b,0x77,0x4,0x51,0xb9,0x20,0xe7,0x74,0xfe,0x43,0xe8,0x23,0xd4,0x87,0x8a},
+{0x7d,0x29,0xe8,0xa3,0x92,0x76,0x94,0xf2,0xdd,0xcb,0x7a,0x9,0x9b,0x30,0xd9,0xc1,0x1d,0x1b,0x30,0xfb,0x5b,0xdc,0x1b,0xe0,0xda,0x24,0x49,0x4f,0xf2,0x9c,0x82,0xbf},
+{0xa4,0xe7,0xba,0x31,0xb4,0x70,0xbf,0xff,0xd,0x32,0x44,0x5,0xde,0xf8,0xbc,0x48,0x3b,0xae,0xfc,0x32,0x53,0xbb,0xd3,0x39,0x45,0x9f,0xc3,0xc1,0xe0,0x29,0x8b,0xa0},
+{0xe5,0xc9,0x5,0xfd,0xf7,0xae,0x9,0xf,0x94,0x70,0x34,0x12,0x42,0x90,0xf1,0x34,0xa2,0x71,0xb7,0x1,0xe3,0x44,0xed,0x95,0xe9,0x3b,0x8e,0x36,0x4f,0x2f,0x98,0x4a},
+{0x88,0x40,0x1d,0x63,0xa0,0x6c,0xf6,0x15,0x47,0xc1,0x44,0x4b,0x87,0x52,0xaf,0xff,0x7e,0xbb,0x4a,0xf1,0xe2,0xa,0xc6,0x30,0x46,0x70,0xb6,0xc5,0xcc,0x6e,0x8c,0xe6},
+{0xa4,0xd5,0xa4,0x56,0xbd,0x4f,0xca,0x0,0xda,0x9d,0x84,0x4b,0xc8,0x3e,0x18,0xae,0x73,0x57,0xce,0x45,0x30,0x64,0xd1,0xad,0xe8,0xa6,0xce,0x68,0x14,0x5c,0x25,0x67},
+{0xa3,0xda,0x8c,0xf2,0xcb,0xe,0xe1,0x16,0x33,0xe9,0x6,0x58,0x9a,0x94,0x99,0x9a,0x1f,0x60,0xb2,0x20,0xc2,0x6f,0x84,0x7b,0xd1,0xce,0xac,0x7f,0xa0,0xd1,0x85,0x18},
+{0x32,0x59,0x5b,0xa1,0x8d,0xdd,0x19,0xd3,0x50,0x9a,0x1c,0xc0,0xaa,0xa5,0xb4,0x46,0x9f,0x3d,0x63,0x67,0xe4,0x4,0x6b,0xba,0xf6,0xca,0x19,0xab,0xb,0x56,0xee,0x7e},
+{0x1f,0xb1,0x79,0xea,0xa9,0x28,0x21,0x74,0xe9,0xbd,0xf7,0x35,0x3b,0x36,0x51,0xee,0x1d,0x57,0xac,0x5a,0x75,0x50,0xd3,0x76,0x3a,0x46,0xc2,0xfe,0xa3,0x7d,0x70,0x1},
+{0xf7,0x35,0xc1,0xaf,0x98,0xa4,0xd8,0x42,0x78,0xed,0xec,0x20,0x9e,0x6b,0x67,0x79,0x41,0x83,0x63,0x15,0xea,0x3a,0xdb,0xa8,0xfa,0xc3,0x3b,0x4d,0x32,0x83,0x2c,0x83},
+{0xa7,0x40,0x3b,0x1f,0x1c,0x27,0x47,0xf3,0x59,0x40,0xf0,0x34,0xb7,0x2d,0x76,0x9a,0xe7,0x3e,0x4e,0x6c,0xd2,0x21,0x4f,0xfd,0xb8,0xfd,0x8d,0x39,0xdc,0x57,0x59,0xef},
+{0x8d,0x9b,0xc,0x49,0x2b,0x49,0xeb,0xda,0x5b,0xa2,0xd7,0x49,0x68,0xf3,0x70,0xd,0x7d,0x3b,0xae,0xd0,0x7a,0x8d,0x55,0x84,0xf5,0xa5,0xe9,0xf0,0xe4,0xf8,0x8e,0x65},
+{0xa0,0xb8,0xa2,0xf4,0x36,0x10,0x3b,0x53,0xc,0xa8,0x7,0x9e,0x75,0x3e,0xec,0x5a,0x91,0x68,0x94,0x92,0x56,0xe8,0x88,0x4f,0x5b,0xb0,0x5c,0x55,0xf8,0xba,0xbc,0x4c},
+{0xe3,0xbb,0x3b,0x99,0xf3,0x87,0x94,0x7b,0x75,0xda,0xf4,0xd6,0x72,0x6b,0x1c,0x5d,0x64,0xae,0xac,0x28,0xdc,0x34,0xb3,0x6d,0x6c,0x34,0xa5,0x50,0xb8,0x28,0xdb,0x71},
+{0xf8,0x61,0xe2,0xf2,0x10,0x8d,0x51,0x2a,0xe3,0xdb,0x64,0x33,0x59,0xdd,0x75,0xfc,0x1c,0xac,0xbc,0xf1,0x43,0xce,0x3f,0xa2,0x67,0xbb,0xd1,0x3c,0x2,0xe8,0x43,0xb0},
+{0x33,0xa,0x5b,0xca,0x88,0x29,0xa1,0x75,0x7f,0x34,0x19,0x4d,0xb4,0x16,0x53,0x5c,0x92,0x3b,0x94,0xc3,0xe,0x79,0x4d,0x1e,0x79,0x74,0x75,0xd7,0xb6,0xee,0xaf,0x3f},
+{0xea,0xa8,0xd4,0xf7,0xbe,0x1a,0x39,0x21,0x5c,0xf4,0x7e,0x9,0x4c,0x23,0x27,0x51,0x26,0xa3,0x24,0x53,0xba,0x32,0x3c,0xd2,0x44,0xa3,0x17,0x4a,0x6d,0xa6,0xd5,0xad},
+{0xb5,0x1d,0x3e,0xa6,0xaf,0xf2,0xc9,0x8,0x83,0x59,0x3d,0x98,0x91,0x6b,0x3c,0x56,0x4c,0xf8,0x7c,0xa1,0x72,0x86,0x60,0x4d,0x46,0xe2,0x3e,0xcc,0x8,0x6e,0xc7,0xf6},
+{0x2f,0x98,0x33,0xb3,0xb1,0xbc,0x76,0x5e,0x2b,0xd6,0x66,0xa5,0xef,0xc4,0xe6,0x2a,0x6,0xf4,0xb6,0xe8,0xbe,0xc1,0xd4,0x36,0x74,0xee,0x82,0x15,0xbc,0xef,0x21,0x63},
+{0xfd,0xc1,0x4e,0xd,0xf4,0x53,0xc9,0x69,0xa7,0x7d,0x5a,0xc4,0x6,0x58,0x58,0x26,0x7e,0xc1,0x14,0x16,0x6,0xe0,0xfa,0x16,0x7e,0x90,0xaf,0x3d,0x28,0x63,0x9d,0x3f},
+{0xd2,0xc9,0xf2,0xe3,0x0,0x9b,0xd2,0xc,0x5f,0xaa,0xce,0x30,0xb7,0xd4,0xc,0x30,0x74,0x2a,0x51,0x16,0xf2,0xe0,0x32,0x98,0xd,0xeb,0x30,0xd8,0xe3,0xce,0xf8,0x9a},
+{0x4b,0xc5,0x9e,0x7b,0xb5,0xf1,0x79,0x92,0xff,0x51,0xe6,0x6e,0x4,0x86,0x68,0xd3,0x9b,0x23,0x4d,0x57,0xe6,0x96,0x67,0x31,0xcc,0xe6,0xa6,0xf3,0x17,0xa,0x75,0x5},
+{0xb1,0x76,0x81,0xd9,0x13,0x32,0x6c,0xce,0x3c,0x17,0x52,0x84,0xf8,0x5,0xa2,0x62,0xf4,0x2b,0xcb,0xb3,0x78,0x47,0x15,0x47,0xff,0x46,0x54,0x82,0x23,0x93,0x6a,0x48},
+{0x38,0xdf,0x58,0x7,0x4e,0x5e,0x65,0x65,0xf2,0xfc,0x7c,0x89,0xfc,0x86,0x50,0x8e,0x31,0x70,0x2e,0x44,0xd0,0xb,0xca,0x86,0xf0,0x40,0x9,0xa2,0x30,0x78,0x47,0x4e},
+{0x65,0xa0,0xee,0x39,0xd1,0xf7,0x38,0x83,0xf7,0x5e,0xe9,0x37,0xe4,0x2c,0x3a,0xbd,0x21,0x97,0xb2,0x26,0x1,0x13,0xf8,0x6f,0xa3,0x44,0xed,0xd1,0xef,0x9f,0xde,0xe7},
+{0x8b,0xa0,0xdf,0x15,0x76,0x25,0x92,0xd9,0x3c,0x85,0xf7,0xf6,0x12,0xdc,0x42,0xbe,0xd8,0xa7,0xec,0x7c,0xab,0x27,0xb0,0x7e,0x53,0x8d,0x7d,0xda,0xaa,0x3e,0xa8,0xde},
+{0xaa,0x25,0xce,0x93,0xbd,0x2,0x69,0xd8,0x5a,0xf6,0x43,0xfd,0x1a,0x73,0x8,0xf9,0xc0,0x5f,0xef,0xda,0x17,0x4a,0x19,0xa5,0x97,0x4d,0x66,0x33,0x4c,0xfd,0x21,0x6a},
+{0x35,0xb4,0x98,0x31,0xdb,0x41,0x15,0x70,0xea,0x1e,0xf,0xbb,0xed,0xcd,0x54,0x9b,0x9a,0xd0,0x63,0xa1,0x51,0x97,0x40,0x72,0xf6,0x75,0x9d,0xbf,0x91,0x47,0x6f,0xe2}};
+
+/*swapping bit 2i with bit 2i+1 of 32-bit x*/
+#define SWAP1(x)   (x) = ((((x) & 0x55555555UL) << 1) | (((x) & 0xaaaaaaaaUL) >> 1));
+/*swapping bits 4i||4i+1 with bits 4i+2||4i+3 of 32-bit x*/
+#define SWAP2(x)   (x) = ((((x) & 0x33333333UL) << 2) | (((x) & 0xccccccccUL) >> 2));
+/*swapping bits 8i||8i+1||8i+2||8i+3 with bits 8i+4||8i+5||8i+6||8i+7 of 32-bit x*/
+#define SWAP4(x)   (x) = ((((x) & 0x0f0f0f0fUL) << 4) | (((x) & 0xf0f0f0f0UL) >> 4));
+/*swapping bits 16i||16i+1||......||16i+7  with bits 16i+8||16i+9||......||16i+15 of 32-bit x*/
+//#define SWAP8(x)   (x) = ((((x) & 0x00ff00ffUL) << 8) | (((x) & 0xff00ff00UL) >> 8));
+#define SWAP8(x) (x) = __byte_perm(x, x, 0x2301);
+/*swapping bits 32i||32i+1||......||32i+15 with bits 32i+16||32i+17||......||32i+31 of 32-bit x*/
+//#define SWAP16(x)  (x) = ((((x) & 0x0000ffffUL) << 16) | (((x) & 0xffff0000UL) >> 16));
+#define SWAP16(x) (x) = __byte_perm(x, x, 0x1032);
+
+/*The MDS transform*/
+#define L(m0,m1,m2,m3,m4,m5,m6,m7) \
+      (m4) ^= (m1);                \
+      (m5) ^= (m2);                \
+      (m6) ^= (m0) ^ (m3);         \
+      (m7) ^= (m0);                \
+      (m0) ^= (m5);                \
+      (m1) ^= (m6);                \
+      (m2) ^= (m4) ^ (m7);         \
+      (m3) ^= (m4);
+
+/*The Sbox*/
+#define Sbox(m0,m1,m2,m3,cc)       \
+      m3  = ~(m3);                 \
+      m0 ^= ((~(m2)) & (cc));      \
+      temp0 = (cc) ^ ((m0) & (m1));\
+      m0 ^= ((m2) & (m3));         \
+      m3 ^= ((~(m1)) & (m2));      \
+      m1 ^= ((m0) & (m2));         \
+      m2 ^= ((m0) & (~(m3)));      \
+      m0 ^= ((m1) | (m3));         \
+      m3 ^= ((m1) & (m2));         \
+      m1 ^= (temp0 & (m0));        \
+      m2 ^= temp0;
+
+__device__ __forceinline__ void Sbox_and_MDS_layer(hashState* state, uint32_t roundnumber)
+{
+    uint32_t temp0;
+	uint32_t cc0, cc1;
+    //Sbox and MDS layer
+#pragma unroll 4
+    for (int i = 0; i < 4; i++) {
+		cc0 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i];
+		cc1 = ((uint32_t*)c_E8_bitslice_roundconstant[roundnumber])[i+4];
+        Sbox(state->x[0][i],state->x[2][i], state->x[4][i], state->x[6][i], cc0);
+        Sbox(state->x[1][i],state->x[3][i], state->x[5][i], state->x[7][i], cc1);
+        L(state->x[0][i],state->x[2][i],state->x[4][i],state->x[6][i],state->x[1][i],state->x[3][i],state->x[5][i],state->x[7][i]);
+    }
+}
+
+__device__ __forceinline__ void RoundFunction0(hashState* state, uint32_t roundnumber)
+{
+	Sbox_and_MDS_layer(state, roundnumber);
+
+#pragma unroll 4
+	for (int j = 1; j < 8; j = j+2)
+	{
+#pragma unroll 4
+		for (int i = 0; i < 4; i++) SWAP1(state->x[j][i]);
+	}
+}
+
+__device__ __forceinline__ void RoundFunction1(hashState* state, uint32_t roundnumber)
+{
+	Sbox_and_MDS_layer(state, roundnumber);
+
+#pragma unroll 4
+	for (int j = 1; j < 8; j = j+2)
+	{
+#pragma unroll 4
+		for (int i = 0; i < 4; i++) SWAP2(state->x[j][i]);
+	}
+}
+
+__device__ __forceinline__ void RoundFunction2(hashState* state, uint32_t roundnumber)
+{
+	Sbox_and_MDS_layer(state, roundnumber);
+
+#pragma unroll 4
+	for (int j = 1; j < 8; j = j+2)
+	{
+#pragma unroll 4
+		for (int i = 0; i < 4; i++) SWAP4(state->x[j][i]);
+	}
+}
+
+__device__ __forceinline__ void RoundFunction3(hashState* state, uint32_t roundnumber)
+{
+	Sbox_and_MDS_layer(state, roundnumber);
+
+#pragma unroll 4
+	for (int j = 1; j < 8; j = j+2)
+	{
+#pragma unroll 4
+		for (int i = 0; i < 4; i++) SWAP8(state->x[j][i]);
+	}
+}
+
+__device__ __forceinline__ void RoundFunction4(hashState* state, uint32_t roundnumber)
+{
+	Sbox_and_MDS_layer(state, roundnumber);
+
+#pragma unroll 4
+	for (int j = 1; j < 8; j = j+2)
+	{
+#pragma unroll 4
+		for (int i = 0; i < 4; i++) SWAP16(state->x[j][i]);
+	}
+}
+
+__device__ __forceinline__ void RoundFunction5(hashState* state, uint32_t roundnumber)
+{
+	uint32_t temp0;
+
+	Sbox_and_MDS_layer(state, roundnumber);
+
+#pragma unroll 4
+	for (int j = 1; j < 8; j = j+2)
+	{
+#pragma unroll 2
+		for (int i = 0; i < 4; i = i+2) {
+			temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+1]; state->x[j][i+1] = temp0;
+		}
+	}
+}
+
+__device__ __forceinline__ void RoundFunction6(hashState* state, uint32_t roundnumber)
+{
+	uint32_t temp0;
+
+	Sbox_and_MDS_layer(state, roundnumber);
+
+#pragma unroll 4
+	for (int j = 1; j < 8; j = j+2)
+	{
+#pragma unroll 2
+		for (int i = 0; i < 2; i++) {
+			temp0 = state->x[j][i]; state->x[j][i] = state->x[j][i+2]; state->x[j][i+2] = temp0;
+		}
+	}
+}
+
+/*The bijective function E8, in bitslice form */
+__device__ __forceinline__ void E8(hashState *state)
+{
+    /*perform 6 rounds*/
+//#pragma unroll 6
+    for (int i = 0; i < 42; i+=7)
+	{
+		RoundFunction0(state, i);
+		RoundFunction1(state, i+1);
+		RoundFunction2(state, i+2);
+		RoundFunction3(state, i+3);
+		RoundFunction4(state, i+4);
+		RoundFunction5(state, i+5);
+		RoundFunction6(state, i+6);
+	}
+}
+
+/*The compression function F8 */
+__device__ __forceinline__ void F8(hashState *state)
+{
+    /*xor the 512-bit message with the fist half of the 1024-bit hash state*/
+#pragma unroll 16
+    for (int i = 0; i < 16; i++)  state->x[i >> 2][i & 3] ^= ((uint32_t*)state->buffer)[i];
+
+    /*the bijective function E8 */
+    E8(state);
+
+    /*xor the 512-bit message with the second half of the 1024-bit hash state*/
+#pragma unroll 16
+    for (int i = 0; i < 16; i++)  state->x[(16+i) >> 2][(16+i) & 3] ^= ((uint32_t*)state->buffer)[i];
+}
+
+
+__device__ __forceinline__ void JHHash(const uint32_t *data, uint32_t *hashval)
+{
+    hashState state;
+
+    /*load the intital hash value H0 into state*/
+	/*
+    #define INIT(a,b,c,d) ((a) | ((b)<<8) | ((c)<<16) | ((d)<<24))
+    state.x[0][0] = INIT(0x6f,0xd1,0x4b,0x96);
+    state.x[0][1] = INIT(0x3e,0x00,0xaa,0x17);
+    state.x[0][2] = INIT(0x63,0x6a,0x2e,0x05);
+    state.x[0][3] = INIT(0x7a,0x15,0xd5,0x43);
+    state.x[1][0] = INIT(0x8a,0x22,0x5e,0x8d);
+    state.x[1][1] = INIT(0x0c,0x97,0xef,0x0b);
+    state.x[1][2] = INIT(0xe9,0x34,0x12,0x59);
+    state.x[1][3] = INIT(0xf2,0xb3,0xc3,0x61);
+    state.x[2][0] = INIT(0x89,0x1d,0xa0,0xc1);
+    state.x[2][1] = INIT(0x53,0x6f,0x80,0x1e);
+    state.x[2][2] = INIT(0x2a,0xa9,0x05,0x6b);
+    state.x[2][3] = INIT(0xea,0x2b,0x6d,0x80);
+    state.x[3][0] = INIT(0x58,0x8e,0xcc,0xdb);
+    state.x[3][1] = INIT(0x20,0x75,0xba,0xa6);
+    state.x[3][2] = INIT(0xa9,0x0f,0x3a,0x76);
+    state.x[3][3] = INIT(0xba,0xf8,0x3b,0xf7);
+    state.x[4][0] = INIT(0x01,0x69,0xe6,0x05);
+    state.x[4][1] = INIT(0x41,0xe3,0x4a,0x69);
+    state.x[4][2] = INIT(0x46,0xb5,0x8a,0x8e);
+    state.x[4][3] = INIT(0x2e,0x6f,0xe6,0x5a);
+    state.x[5][0] = INIT(0x10,0x47,0xa7,0xd0);
+    state.x[5][1] = INIT(0xc1,0x84,0x3c,0x24);
+    state.x[5][2] = INIT(0x3b,0x6e,0x71,0xb1);
+    state.x[5][3] = INIT(0x2d,0x5a,0xc1,0x99);
+    state.x[6][0] = INIT(0xcf,0x57,0xf6,0xec);
+    state.x[6][1] = INIT(0x9d,0xb1,0xf8,0x56);
+    state.x[6][2] = INIT(0xa7,0x06,0x88,0x7c);
+    state.x[6][3] = INIT(0x57,0x16,0xb1,0x56);
+    state.x[7][0] = INIT(0xe3,0xc2,0xfc,0xdf);
+    state.x[7][1] = INIT(0xe6,0x85,0x17,0xfb);
+    state.x[7][2] = INIT(0x54,0x5a,0x46,0x78);
+    state.x[7][3] = INIT(0xcc,0x8c,0xdd,0x4b);
+	*/
+#pragma unroll 8
+	for(int j=0;j<8;j++)
+	{
+#pragma unroll 4
+		for(int i=0;i<4;i++)
+			state.x[j][i] = c_INIT_bitslice[j][i];
+	}
+
+#pragma unroll 16
+    for (int i=0; i < 16; ++i) state.buffer[i] = data[i];
+    F8(&state);
+
+    /*pad the message when databitlen is multiple of 512 bits, then process the padded block*/
+    state.buffer[0] = 0x80;
+#pragma unroll 14
+    for (int i=1; i < 15; i++) state.buffer[i] = 0;
+    state.buffer[15] = 0x00020000;
+    F8(&state);
+
+    /*truncating the final hash value to generate the message digest*/
+#pragma unroll 16
+    for (int i=0; i < 16; ++i) hashval[i] = state.x[4][i];
+}
+
+// Die Hash-Funktion
+__global__ void quark_jh512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
+{
+    int thread = (blockDim.x * blockIdx.x + threadIdx.x);
+    if (thread < threads)
+    {
+        uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
+
+        int hashPosition = nounce - startNounce;
+        uint32_t *Hash = (uint32_t*)&g_hash[8 * hashPosition];
+
+        JHHash(Hash, Hash);
+    }
+}
+
+
+// Setup-Funktionen
+__host__ void quark_jh512_cpu_init(int thr_id, int threads)
+{
+	
+    cudaMemcpyToSymbol( c_E8_bitslice_roundconstant,
+                        h_E8_bitslice_roundconstant,
+                        sizeof(h_E8_bitslice_roundconstant),
+                        0, cudaMemcpyHostToDevice);
+
+	cudaMemcpyToSymbol( c_INIT_bitslice,
+                        h_INIT_bitslice,
+                        sizeof(h_INIT_bitslice),
+                        0, cudaMemcpyHostToDevice);
+}
+
+__host__ void quark_jh512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
+{
+    const int threadsperblock = 256;
+
+    // berechne wie viele Thread Blocks wir brauchen
+    dim3 grid((threads + threadsperblock-1)/threadsperblock);
+    dim3 block(threadsperblock);
+
+    // Größe des dynamischen Shared Memory Bereichs
+    size_t shared_size = 0;
+
+//    fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
+
+    quark_jh512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
+    MyStreamSynchronize(NULL, order, thr_id);
+}
+
diff --git a/quark/cuda_quark_blake512.cu b/quark/cuda_quark_blake512.cu
index 0e8cb3b..80e500e 100644
--- a/quark/cuda_quark_blake512.cu
+++ b/quark/cuda_quark_blake512.cu
@@ -1,482 +1,434 @@
-#include <cuda.h>
-#include "cuda_runtime.h"
-#include "device_launch_parameters.h"
-
-#include <stdio.h>
-#include <memory.h>
-
-#define USE_SHUFFLE 0
-
-// Folgende Definitionen später durch header ersetzen
-typedef unsigned char uint8_t;
-typedef unsigned int uint32_t;
-typedef unsigned long long uint64_t;
-
-// aus heavy.cu
-extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
-
-// die Message it Padding zur Berechnung auf der GPU
-__constant__ uint64_t c_PaddedMessage80[16]; // padded message (80 bytes + padding)
-
-// ---------------------------- BEGIN CUDA quark_blake512 functions ------------------------------------
-
-__constant__ uint8_t c_sigma[16][16];
-
-const uint8_t host_sigma[16][16] =
-{
-  { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
-  {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
-  {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
-  { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
-  { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
-  { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
-  {12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
-  {13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
-  { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
-  {10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13 , 0 },
-  { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
-  {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
-  {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
-  { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
-  { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
-  { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 }
-};
-
-// das Hi Word aus einem 64 Bit Typen extrahieren
-static __device__ uint32_t HIWORD(const uint64_t &x) {
-#if __CUDA_ARCH__ >= 130
-	return (uint32_t)__double2hiint(__longlong_as_double(x));
-#else
-	return (uint32_t)(x >> 32);
-#endif
-}
-
-// das Hi Word in einem 64 Bit Typen ersetzen
-static __device__ uint64_t REPLACE_HIWORD(const uint64_t &x, const uint32_t &y) {
-	return (x & 0xFFFFFFFFULL) | (((uint64_t)y) << 32ULL);
-}
-
-// das Lo Word aus einem 64 Bit Typen extrahieren
-static __device__ uint32_t LOWORD(const uint64_t &x) {
-#if __CUDA_ARCH__ >= 130
-	return (uint32_t)__double2loint(__longlong_as_double(x));
-#else
-	return (uint32_t)(x & 0xFFFFFFFFULL);
-#endif
-}
-
-// das Lo Word in einem 64 Bit Typen ersetzen
-static __device__ uint64_t REPLACE_LOWORD(const uint64_t &x, const uint32_t &y) {
-	return (x & 0xFFFFFFFF00000000ULL) | ((uint64_t)y);
-}
-
-/*
-#define SWAP32(x) \
-    ((((x) << 24) & 0xff000000u) | (((x) << 8) & 0x00ff0000u)   | \
-      (((x) >> 8) & 0x0000ff00u) | (((x) >> 24) & 0x000000ffu))
-
-#define SWAP64(x) \
-    ((uint64_t)((((uint64_t)(x) & 0xff00000000000000ULL) >> 56) | \
-                (((uint64_t)(x) & 0x00ff000000000000ULL) >> 40) | \
-                (((uint64_t)(x) & 0x0000ff0000000000ULL) >> 24) | \
-                (((uint64_t)(x) & 0x000000ff00000000ULL) >>  8) | \
-                (((uint64_t)(x) & 0x00000000ff000000ULL) <<  8) | \
-                (((uint64_t)(x) & 0x0000000000ff0000ULL) << 24) | \
-                (((uint64_t)(x) & 0x000000000000ff00ULL) << 40) | \
-                (((uint64_t)(x) & 0x00000000000000ffULL) << 56)))
-*/
-
-/*
-__device__ __forceinline__ void SWAP32(uint32_t *x)
-{
-	// Input:	33221100
-	// Output:	00112233
-	x[0] = __byte_perm(x[0], 0, 0x0123);
-}
-*/
-__device__ __forceinline__ uint64_t SWAP64(uint64_t x)
-{
-	// Input:	77665544 33221100
-	// Output:	00112233 44556677
-	uint64_t temp[2];
-	temp[0] = __byte_perm(HIWORD(x), 0, 0x0123);
-	temp[1] = __byte_perm(LOWORD(x), 0, 0x0123);
-
-	return temp[0] | (temp[1]<<32);
-}
-
-__constant__ uint64_t c_u512[16];
-
-const uint64_t host_u512[16] =
-{
-  0x243f6a8885a308d3ULL, 0x13198a2e03707344ULL, 
-  0xa4093822299f31d0ULL, 0x082efa98ec4e6c89ULL,
-  0x452821e638d01377ULL, 0xbe5466cf34e90c6cULL, 
-  0xc0ac29b7c97c50ddULL, 0x3f84d5b5b5470917ULL,
-  0x9216d5d98979fb1bULL, 0xd1310ba698dfb5acULL, 
-  0x2ffd72dbd01adfb7ULL, 0xb8e1afed6a267e96ULL,
-  0xba7c9045f12c7f99ULL, 0x24a19947b3916cf7ULL, 
-  0x0801f2e2858efc16ULL, 0x636920d871574e69ULL
-};
-
-
-// diese 64 Bit Rotates werden unter Compute 3.5 (und besser) mit dem Funnel Shifter beschleunigt
-#if __CUDA_ARCH__ >= 350
-__forceinline__ __device__ uint64_t ROTR(const uint64_t value, const int offset) {
-    uint2 result;
-    if(offset < 32) {
-        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
-        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
-    } else {
-        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
-        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
-    }
-    return  __double_as_longlong(__hiloint2double(result.y, result.x));
-}
-#else
-#define ROTR(x, n)        (((x) >> (n)) | ((x) << (64 - (n))))
-#endif
-
-#define G(a,b,c,d,e)          \
-    v[a] += (m[sigma[i][e]] ^ u512[sigma[i][e+1]]) + v[b];\
-    v[d] = ROTR( v[d] ^ v[a],32);        \
-    v[c] += v[d];           \
-    v[b] = ROTR( v[b] ^ v[c],25);        \
-    v[a] += (m[sigma[i][e+1]] ^ u512[sigma[i][e]])+v[b];  \
-    v[d] = ROTR( v[d] ^ v[a],16);        \
-    v[c] += v[d];           \
-    v[b] = ROTR( v[b] ^ v[c],11);
-
-
-__device__ void quark_blake512_compress( uint64_t *h, const uint64_t *block, const uint8_t ((*sigma)[16]), const uint64_t *u512, const int bits )
-{
-    uint64_t v[16], m[16], i;
-
-#pragma unroll 16
-    for( i = 0; i < 16; ++i )
-	{
-		m[i] = SWAP64(block[i]);		
-	}
-
-#pragma unroll 8
-    for( i = 0; i < 8; ++i )  v[i] = h[i];
-
-    v[ 8] = u512[0];
-    v[ 9] = u512[1];
-    v[10] = u512[2];
-    v[11] = u512[3];
-    v[12] = u512[4];
-    v[13] = u512[5];
-    v[14] = u512[6];
-    v[15] = u512[7];
-
-    v[12] ^= bits;
-    v[13] ^= bits;
-
-//#pragma unroll 16
-    for( i = 0; i < 16; ++i )
-    {
-        /* column step */
-        G( 0, 4, 8, 12, 0 );
-        G( 1, 5, 9, 13, 2 );
-        G( 2, 6, 10, 14, 4 );
-        G( 3, 7, 11, 15, 6 );
-        /* diagonal step */
-        G( 0, 5, 10, 15, 8 );
-        G( 1, 6, 11, 12, 10 );
-        G( 2, 7, 8, 13, 12 );
-        G( 3, 4, 9, 14, 14 );
-    }
-
-#pragma unroll 16
-    for( i = 0; i < 16; ++i )  h[i % 8] ^= v[i];
-}
-
-// Endian Drehung für 32 Bit Typen
-
-static __device__ uint32_t cuda_swab32(uint32_t x)
-{
-	return __byte_perm(x, 0, 0x0123);
-	/*
-    return (((x << 24) & 0xff000000u) | ((x << 8) & 0x00ff0000u)
-          | ((x >> 8) & 0x0000ff00u) | ((x >> 24) & 0x000000ffu));
-		  */
-}
-/*
-// Endian Drehung für 64 Bit Typen
-static __device__ uint64_t cuda_swab64(uint64_t x) {
-    uint32_t h = (x >> 32);
-    uint32_t l = (x & 0xFFFFFFFFULL);
-    return (((uint64_t)cuda_swab32(l)) << 32) | ((uint64_t)cuda_swab32(h));
-}
-*/
-
-static __constant__ uint64_t d_constMem[8];
-static const uint64_t h_constMem[8] = {
-	0x6a09e667f3bcc908ULL,
-	0xbb67ae8584caa73bULL,
-	0x3c6ef372fe94f82bULL,
-	0xa54ff53a5f1d36f1ULL,
-	0x510e527fade682d1ULL,
-	0x9b05688c2b3e6c1fULL,
-	0x1f83d9abfb41bd6bULL,
-	0x5be0cd19137e2179ULL };
-
-// Hash-Padding
-static __constant__ uint64_t d_constHashPadding[8];
-static const uint64_t h_constHashPadding[8] = {
-	0x0000000000000080ull,
-	0,
-	0,
-	0,
-	0,
-	0x0100000000000000ull,
-	0,
-	0x0002000000000000ull };
-
-__global__ void quark_blake512_gpu_hash_64(int threads, uint32_t startNounce, uint32_t *g_nonceVector, uint64_t *g_hash)
-{
-	int thread = (blockDim.x * blockIdx.x + threadIdx.x);
-
-#if USE_SHUFFLE
-	const int warpID = threadIdx.x & 0x0F; // 16 warps
-	const int warpBlockID = (thread + 15)>>4; // aufrunden auf volle Warp-Blöcke
-	const int maxHashPosition = thread<<3;
-#endif
-
-#if USE_SHUFFLE
-	if (warpBlockID < ( (threads+15)>>4 ))
-#else
-	if (thread < threads)
-#endif
-	{
-		// bestimme den aktuellen Zähler
-		uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
-
-		int hashPosition = nounce - startNounce;
-		//uint64_t *inpHash = &g_hash[8 * hashPosition];
-		uint64_t *inpHash = &g_hash[hashPosition<<3];
-
-		// State vorbereiten
-		uint64_t h[8];
-		/*
-		h[0] = 0x6a09e667f3bcc908ULL;
-		h[1] = 0xbb67ae8584caa73bULL;
-		h[2] = 0x3c6ef372fe94f82bULL;
-		h[3] = 0xa54ff53a5f1d36f1ULL;
-		h[4] = 0x510e527fade682d1ULL;
-		h[5] = 0x9b05688c2b3e6c1fULL;
-		h[6] = 0x1f83d9abfb41bd6bULL;
-		h[7] = 0x5be0cd19137e2179ULL;
-		*/
-#pragma unroll 8
-		for(int i=0;i<8;i++)
-			h[i] = d_constMem[i];
-
-		// 128 Byte für die Message
-		uint64_t buf[16];
-
-		// Message für die erste Runde in Register holen
-#pragma unroll 8
-		for (int i=0; i < 8; ++i) buf[i] = inpHash[i];
-
-		/*
-		buf[ 8] = 0x0000000000000080ull;
-		buf[ 9] = 0;
-		buf[10] = 0;
-		buf[11] = 0;
-		buf[12] = 0;
-		buf[13] = 0x0100000000000000ull;
-		buf[14] = 0;
-		buf[15] = 0x0002000000000000ull;
-		*/
-#pragma unroll 8
-		for(int i=0;i<8;i++)
-			buf[i+8] = d_constHashPadding[i];
-
-		// die einzige Hashing-Runde
-		quark_blake512_compress( h, buf, c_sigma, c_u512, 512 );
-
-		// Hash rauslassen
-#if __CUDA_ARCH__ >= 130
-		// ausschliesslich 32 bit Operationen sofern die SM1.3 double intrinsics verfügbar sind
-		uint32_t *outHash = (uint32_t*)&g_hash[8 * hashPosition];
-#pragma unroll 8
-		for (int i=0; i < 8; ++i) {
-			outHash[2*i+0] = cuda_swab32( HIWORD(h[i]) );
-			outHash[2*i+1] = cuda_swab32( LOWORD(h[i]) );
-		}
-#else
-		// in dieser Version passieren auch ein paar 64 Bit Shifts
-		uint64_t *outHash = &g_hash[8 * hashPosition];
-#pragma unroll 8
-		for (int i=0; i < 8; ++i)
-		{
-			//outHash[i] = cuda_swab64( h[i] );
-			outHash[i] = SWAP64(h[i]);		
-		}
-#endif
-	}
-}
-
-__global__ void quark_blake512_gpu_hash_80(int threads, uint32_t startNounce, void *outputHash)
-{
-	int thread = (blockDim.x * blockIdx.x + threadIdx.x);
-	if (thread < threads)
-	{
-		// bestimme den aktuellen Zähler
-		uint32_t nounce = startNounce + thread;
-
-		// State vorbereiten
-		uint64_t h[8];
-		/*
-		h[0] = 0x6a09e667f3bcc908ULL;
-		h[1] = 0xbb67ae8584caa73bULL;
-		h[2] = 0x3c6ef372fe94f82bULL;
-		h[3] = 0xa54ff53a5f1d36f1ULL;
-		h[4] = 0x510e527fade682d1ULL;
-		h[5] = 0x9b05688c2b3e6c1fULL;
-		h[6] = 0x1f83d9abfb41bd6bULL;
-		h[7] = 0x5be0cd19137e2179ULL;
-		*/
-#pragma unroll 8
-		for(int i=0;i<8;i++)
-			h[i] = d_constMem[i];
-		// 128 Byte für die Message
-		uint64_t buf[16];
-
-		// Message für die erste Runde in Register holen
-#pragma unroll 16
-		for (int i=0; i < 16; ++i) buf[i] = c_PaddedMessage80[i];
-
-		// die Nounce durch die thread-spezifische ersetzen
-		buf[9] = REPLACE_HIWORD(buf[9], cuda_swab32(nounce));
-
-		// die einzige Hashing-Runde
-		quark_blake512_compress( h, buf, c_sigma, c_u512, 640 );
-
-		// Hash rauslassen
-#if __CUDA_ARCH__ >= 130
-		// ausschliesslich 32 bit Operationen sofern die SM1.3 double intrinsics verfügbar sind
-		uint32_t *outHash = (uint32_t *)outputHash + 16 * thread;
-#pragma unroll 8
-		for (int i=0; i < 8; ++i) {
-			outHash[2*i+0] = cuda_swab32( HIWORD(h[i]) );
-			outHash[2*i+1] = cuda_swab32( LOWORD(h[i]) );
-		}
-#else
-		// in dieser Version passieren auch ein paar 64 Bit Shifts
-		uint64_t *outHash = (uint64_t *)outputHash + 8 * thread;
-#pragma unroll 8
-		for (int i=0; i < 8; ++i)
-		{
-			//outHash[i] = cuda_swab64( h[i] );
-			outHash[i] = SWAP64(h[i]);			
-		}
-#endif
-	}
-}
-
-
-// ---------------------------- END CUDA quark_blake512 functions ------------------------------------
-
-// Setup-Funktionen
-__host__ void quark_blake512_cpu_init(int thr_id, int threads)
-{
-	// Kopiere die Hash-Tabellen in den GPU-Speicher
-	cudaMemcpyToSymbol( c_sigma,
-						host_sigma,
-						sizeof(host_sigma),
-						0, cudaMemcpyHostToDevice);
-
-	cudaMemcpyToSymbol( c_u512,
-						host_u512,
-						sizeof(host_u512),
-						0, cudaMemcpyHostToDevice);
-
-	cudaMemcpyToSymbol( d_constMem,
-                        h_constMem,
-                        sizeof(h_constMem),
-                        0, cudaMemcpyHostToDevice);
-
-	cudaMemcpyToSymbol( d_constHashPadding,
-                        h_constHashPadding,
-                        sizeof(h_constHashPadding),
-                        0, cudaMemcpyHostToDevice);
-}
-
-// Blake512 für 80 Byte grosse Eingangsdaten
-__host__ void quark_blake512_cpu_setBlock_80(void *pdata)
-{
-	// Message mit Padding bereitstellen
-	// lediglich die korrekte Nonce ist noch ab Byte 76 einzusetzen.
-	unsigned char PaddedMessage[128];
-	memcpy(PaddedMessage, pdata, 80);
-	memset(PaddedMessage+80, 0, 48);
-	PaddedMessage[80] = 0x80;
-	PaddedMessage[111] = 1;
-	PaddedMessage[126] = 0x02;
-	PaddedMessage[127] = 0x80;
-
-	// die Message zur Berechnung auf der GPU
-	cudaMemcpyToSymbol( c_PaddedMessage80, PaddedMessage, 16*sizeof(uint64_t), 0, cudaMemcpyHostToDevice);
-}
-
-#if 0
-// Blake512 für 64 Byte grosse Eingangsdaten
-// evtl. macht es gar keinen Sinn, das alles ins Constant Memory to schicken. Es sind hier sowieso
-// nur die letzten 64 Bytes des Blocks konstant, und die meisten Bytes davon sind 0. Das kann mnan
-// auch im Kernel initialisieren.
-__host__ void quark_blake512_cpu_setBlock_64(void *pdata)
-{
-	// Message mit Padding bereitstellen
-	unsigned char PaddedMessage[128];
-	memcpy(PaddedMessage, pdata, 64);  // Hinweis: diese 64 Bytes sind nonce-spezifisch und ändern sich KOMPLETT für jede Nonce!
-	memset(PaddedMessage+64, 0, 64);
-	PaddedMessage[64] = 0x80;
-	PaddedMessage[111] = 1;
-	PaddedMessage[126] = 0x02;
-	PaddedMessage[127] = 0x00;
-
-	// die Message zur Berechnung auf der GPU
-	cudaMemcpyToSymbol( c_PaddedMessage80, PaddedMessage, 16*sizeof(uint64_t), 0, cudaMemcpyHostToDevice);
-}
-#endif
-
-__host__ void quark_blake512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_outputHash, int order)
-{
-	const int threadsperblock = 256;
-
-	// berechne wie viele Thread Blocks wir brauchen
-	dim3 grid((threads + threadsperblock-1)/threadsperblock);
-	dim3 block(threadsperblock);
-
-	// Größe des dynamischen Shared Memory Bereichs
-	size_t shared_size = 0;
-
-//	fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
-
-	quark_blake512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_nonceVector, (uint64_t*)d_outputHash);
-
-	// Strategisches Sleep Kommando zur Senkung der CPU Last
-	MyStreamSynchronize(NULL, order, thr_id);
-}
-
-__host__ void quark_blake512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_outputHash, int order)
-{
-	const int threadsperblock = 256;
-
-	// berechne wie viele Thread Blocks wir brauchen
-	dim3 grid((threads + threadsperblock-1)/threadsperblock);
-	dim3 block(threadsperblock);
-
-	// Größe des dynamischen Shared Memory Bereichs
-	size_t shared_size = 0;
-
-//	fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
-
-	quark_blake512_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, d_outputHash);
-
-	// Strategisches Sleep Kommando zur Senkung der CPU Last
-	MyStreamSynchronize(NULL, order, thr_id);
-}
+#include <cuda.h>
+#include "cuda_runtime.h"
+#include "device_launch_parameters.h"
+
+#include <stdio.h>
+#include <memory.h>
+
+#define USE_SHUFFLE 0
+
+// Folgende Definitionen später durch header ersetzen
+typedef unsigned char uint8_t;
+typedef unsigned int uint32_t;
+typedef unsigned long long uint64_t;
+
+// aus heavy.cu
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+// die Message it Padding zur Berechnung auf der GPU
+__constant__ uint64_t c_PaddedMessage80[16]; // padded message (80 bytes + padding)
+
+// ---------------------------- BEGIN CUDA quark_blake512 functions ------------------------------------
+
+__constant__ uint8_t c_sigma[16][16];
+
+const uint8_t host_sigma[16][16] =
+{
+  { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
+  {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
+  {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
+  { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
+  { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
+  { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
+  {12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
+  {13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
+  { 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
+  {10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13 , 0 },
+  { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
+  {14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
+  {11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
+  { 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
+  { 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
+  { 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 }
+};
+
+// das Hi Word aus einem 64 Bit Typen extrahieren
+static __device__ uint32_t HIWORD(const uint64_t &x) {
+#if __CUDA_ARCH__ >= 130
+	return (uint32_t)__double2hiint(__longlong_as_double(x));
+#else
+	return (uint32_t)(x >> 32);
+#endif
+}
+
+// das Hi Word in einem 64 Bit Typen ersetzen
+static __device__ uint64_t REPLACE_HIWORD(const uint64_t &x, const uint32_t &y) {
+	return (x & 0xFFFFFFFFULL) | (((uint64_t)y) << 32ULL);
+}
+
+// das Lo Word aus einem 64 Bit Typen extrahieren
+static __device__ uint32_t LOWORD(const uint64_t &x) {
+#if __CUDA_ARCH__ >= 130
+	return (uint32_t)__double2loint(__longlong_as_double(x));
+#else
+	return (uint32_t)(x & 0xFFFFFFFFULL);
+#endif
+}
+
+// das Lo Word in einem 64 Bit Typen ersetzen
+static __device__ uint64_t REPLACE_LOWORD(const uint64_t &x, const uint32_t &y) {
+	return (x & 0xFFFFFFFF00000000ULL) | ((uint64_t)y);
+}
+
+__device__ __forceinline__ uint64_t SWAP64(uint64_t x)
+{
+	// Input:	77665544 33221100
+	// Output:	00112233 44556677
+	uint64_t temp[2];
+	temp[0] = __byte_perm(HIWORD(x), 0, 0x0123);
+	temp[1] = __byte_perm(LOWORD(x), 0, 0x0123);
+
+	return temp[0] | (temp[1]<<32);
+}
+
+__constant__ uint64_t c_u512[16];
+
+const uint64_t host_u512[16] =
+{
+  0x243f6a8885a308d3ULL, 0x13198a2e03707344ULL, 
+  0xa4093822299f31d0ULL, 0x082efa98ec4e6c89ULL,
+  0x452821e638d01377ULL, 0xbe5466cf34e90c6cULL, 
+  0xc0ac29b7c97c50ddULL, 0x3f84d5b5b5470917ULL,
+  0x9216d5d98979fb1bULL, 0xd1310ba698dfb5acULL, 
+  0x2ffd72dbd01adfb7ULL, 0xb8e1afed6a267e96ULL,
+  0xba7c9045f12c7f99ULL, 0x24a19947b3916cf7ULL, 
+  0x0801f2e2858efc16ULL, 0x636920d871574e69ULL
+};
+
+
+// diese 64 Bit Rotates werden unter Compute 3.5 (und besser) mit dem Funnel Shifter beschleunigt
+#if __CUDA_ARCH__ >= 350
+__forceinline__ __device__ uint64_t ROTR(const uint64_t value, const int offset) {
+    uint2 result;
+    if(offset < 32) {
+        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
+        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
+    } else {
+        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
+        asm("shf.r.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
+    }
+    return  __double_as_longlong(__hiloint2double(result.y, result.x));
+}
+#else
+#define ROTR(x, n)        (((x) >> (n)) | ((x) << (64 - (n))))
+#endif
+
+#define G(a,b,c,d,e)          \
+    v[a] += (m[sigma[i][e]] ^ u512[sigma[i][e+1]]) + v[b];\
+    v[d] = ROTR( v[d] ^ v[a],32);        \
+    v[c] += v[d];           \
+    v[b] = ROTR( v[b] ^ v[c],25);        \
+    v[a] += (m[sigma[i][e+1]] ^ u512[sigma[i][e]])+v[b];  \
+    v[d] = ROTR( v[d] ^ v[a],16);        \
+    v[c] += v[d];           \
+    v[b] = ROTR( v[b] ^ v[c],11);
+
+
+__device__ void quark_blake512_compress( uint64_t *h, const uint64_t *block, const uint8_t ((*sigma)[16]), const uint64_t *u512, const int bits )
+{
+    uint64_t v[16], m[16], i;
+
+#pragma unroll 16
+    for( i = 0; i < 16; ++i )
+    {
+        m[i] = SWAP64(block[i]);
+    }
+
+#pragma unroll 8
+    for( i = 0; i < 8; ++i )  v[i] = h[i];
+
+    v[ 8] = u512[0];
+    v[ 9] = u512[1];
+    v[10] = u512[2];
+    v[11] = u512[3];
+    v[12] = u512[4];
+    v[13] = u512[5];
+    v[14] = u512[6];
+    v[15] = u512[7];
+
+    v[12] ^= bits;
+    v[13] ^= bits;
+
+//#pragma unroll 16
+    for( i = 0; i < 16; ++i )
+    {
+        /* column step */
+        G( 0, 4, 8, 12, 0 );
+        G( 1, 5, 9, 13, 2 );
+        G( 2, 6, 10, 14, 4 );
+        G( 3, 7, 11, 15, 6 );
+        /* diagonal step */
+        G( 0, 5, 10, 15, 8 );
+        G( 1, 6, 11, 12, 10 );
+        G( 2, 7, 8, 13, 12 );
+        G( 3, 4, 9, 14, 14 );
+    }
+
+#pragma unroll 16
+    for( i = 0; i < 16; ++i )  h[i % 8] ^= v[i];
+}
+
+// Endian Drehung für 32 Bit Typen
+
+static __device__ uint32_t cuda_swab32(uint32_t x)
+{
+	return __byte_perm(x, 0, 0x0123);
+}
+
+/*
+// Endian Drehung für 64 Bit Typen
+static __device__ uint64_t cuda_swab64(uint64_t x) {
+    uint32_t h = (x >> 32);
+    uint32_t l = (x & 0xFFFFFFFFULL);
+    return (((uint64_t)cuda_swab32(l)) << 32) | ((uint64_t)cuda_swab32(h));
+}
+*/
+
+static __constant__ uint64_t d_constMem[8];
+static const uint64_t h_constMem[8] = {
+	0x6a09e667f3bcc908ULL,
+	0xbb67ae8584caa73bULL,
+	0x3c6ef372fe94f82bULL,
+	0xa54ff53a5f1d36f1ULL,
+	0x510e527fade682d1ULL,
+	0x9b05688c2b3e6c1fULL,
+	0x1f83d9abfb41bd6bULL,
+	0x5be0cd19137e2179ULL };
+
+// Hash-Padding
+static __constant__ uint64_t d_constHashPadding[8];
+static const uint64_t h_constHashPadding[8] = {
+	0x0000000000000080ull,
+	0,
+	0,
+	0,
+	0,
+	0x0100000000000000ull,
+	0,
+	0x0002000000000000ull };
+
+__global__ __launch_bounds__(256, 2) void quark_blake512_gpu_hash_64(int threads, uint32_t startNounce, uint32_t *g_nonceVector, uint64_t *g_hash)
+{
+	int thread = (blockDim.x * blockIdx.x + threadIdx.x);
+
+#if USE_SHUFFLE
+	const int warpID = threadIdx.x & 0x0F; // 16 warps
+	const int warpBlockID = (thread + 15)>>4; // aufrunden auf volle Warp-Blöcke
+	const int maxHashPosition = thread<<3;
+#endif
+
+#if USE_SHUFFLE
+	if (warpBlockID < ( (threads+15)>>4 ))
+#else
+	if (thread < threads)
+#endif
+	{
+		// bestimme den aktuellen Zähler
+		uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
+
+		int hashPosition = nounce - startNounce;
+		//uint64_t *inpHash = &g_hash[8 * hashPosition];
+		uint64_t *inpHash = &g_hash[hashPosition<<3];
+
+		// State vorbereiten
+		uint64_t h[8];
+		/*
+		h[0] = 0x6a09e667f3bcc908ULL;
+		h[1] = 0xbb67ae8584caa73bULL;
+		h[2] = 0x3c6ef372fe94f82bULL;
+		h[3] = 0xa54ff53a5f1d36f1ULL;
+		h[4] = 0x510e527fade682d1ULL;
+		h[5] = 0x9b05688c2b3e6c1fULL;
+		h[6] = 0x1f83d9abfb41bd6bULL;
+		h[7] = 0x5be0cd19137e2179ULL;
+		*/
+#pragma unroll 8
+		for(int i=0;i<8;i++)
+			h[i] = d_constMem[i];
+
+		// 128 Byte für die Message
+		uint64_t buf[16];
+
+		// Message für die erste Runde in Register holen
+#pragma unroll 8
+		for (int i=0; i < 8; ++i) buf[i] = inpHash[i];
+
+		/*
+		buf[ 8] = 0x0000000000000080ull;
+		buf[ 9] = 0;
+		buf[10] = 0;
+		buf[11] = 0;
+		buf[12] = 0;
+		buf[13] = 0x0100000000000000ull;
+		buf[14] = 0;
+		buf[15] = 0x0002000000000000ull;
+		*/
+#pragma unroll 8
+		for(int i=0;i<8;i++)
+			buf[i+8] = d_constHashPadding[i];
+
+		// die einzige Hashing-Runde
+		quark_blake512_compress( h, buf, c_sigma, c_u512, 512 );
+
+		// Hash rauslassen
+#if __CUDA_ARCH__ >= 130
+		// ausschliesslich 32 bit Operationen sofern die SM1.3 double intrinsics verfügbar sind
+		uint32_t *outHash = (uint32_t*)&g_hash[8 * hashPosition];
+#pragma unroll 8
+		for (int i=0; i < 8; ++i) {
+			outHash[2*i+0] = cuda_swab32( HIWORD(h[i]) );
+			outHash[2*i+1] = cuda_swab32( LOWORD(h[i]) );
+		}
+#else
+		// in dieser Version passieren auch ein paar 64 Bit Shifts
+		uint64_t *outHash = &g_hash[8 * hashPosition];
+#pragma unroll 8
+		for (int i=0; i < 8; ++i)
+		{
+			//outHash[i] = cuda_swab64( h[i] );
+			outHash[i] = SWAP64(h[i]);
+		}
+#endif
+	}
+}
+
+__global__ void quark_blake512_gpu_hash_80(int threads, uint32_t startNounce, void *outputHash)
+{
+	int thread = (blockDim.x * blockIdx.x + threadIdx.x);
+	if (thread < threads)
+	{
+		// bestimme den aktuellen Zähler
+		uint32_t nounce = startNounce + thread;
+
+		// State vorbereiten
+		uint64_t h[8];
+		/*
+		h[0] = 0x6a09e667f3bcc908ULL;
+		h[1] = 0xbb67ae8584caa73bULL;
+		h[2] = 0x3c6ef372fe94f82bULL;
+		h[3] = 0xa54ff53a5f1d36f1ULL;
+		h[4] = 0x510e527fade682d1ULL;
+		h[5] = 0x9b05688c2b3e6c1fULL;
+		h[6] = 0x1f83d9abfb41bd6bULL;
+		h[7] = 0x5be0cd19137e2179ULL;
+		*/
+#pragma unroll 8
+		for(int i=0;i<8;i++)
+			h[i] = d_constMem[i];
+		// 128 Byte für die Message
+		uint64_t buf[16];
+
+		// Message für die erste Runde in Register holen
+#pragma unroll 16
+		for (int i=0; i < 16; ++i) buf[i] = c_PaddedMessage80[i];
+
+		// die Nounce durch die thread-spezifische ersetzen
+		buf[9] = REPLACE_HIWORD(buf[9], cuda_swab32(nounce));
+
+		// die einzige Hashing-Runde
+		quark_blake512_compress( h, buf, c_sigma, c_u512, 640 );
+
+		// Hash rauslassen
+#if __CUDA_ARCH__ >= 130
+		// ausschliesslich 32 bit Operationen sofern die SM1.3 double intrinsics verfügbar sind
+		uint32_t *outHash = (uint32_t *)outputHash + 16 * thread;
+#pragma unroll 8
+		for (int i=0; i < 8; ++i) {
+			outHash[2*i+0] = cuda_swab32( HIWORD(h[i]) );
+			outHash[2*i+1] = cuda_swab32( LOWORD(h[i]) );
+		}
+#else
+		// in dieser Version passieren auch ein paar 64 Bit Shifts
+		uint64_t *outHash = (uint64_t *)outputHash + 8 * thread;
+#pragma unroll 8
+		for (int i=0; i < 8; ++i)
+		{
+			//outHash[i] = cuda_swab64( h[i] );
+			outHash[i] = SWAP64(h[i]);
+		}
+#endif
+	}
+}
+
+
+// ---------------------------- END CUDA quark_blake512 functions ------------------------------------
+
+// Setup-Funktionen
+__host__ void quark_blake512_cpu_init(int thr_id, int threads)
+{
+	// Kopiere die Hash-Tabellen in den GPU-Speicher
+	cudaMemcpyToSymbol( c_sigma,
+						host_sigma,
+						sizeof(host_sigma),
+						0, cudaMemcpyHostToDevice);
+
+	cudaMemcpyToSymbol( c_u512,
+						host_u512,
+						sizeof(host_u512),
+						0, cudaMemcpyHostToDevice);
+
+	cudaMemcpyToSymbol( d_constMem,
+						h_constMem,
+						sizeof(h_constMem),
+						0, cudaMemcpyHostToDevice);
+
+	cudaMemcpyToSymbol( d_constHashPadding,
+						h_constHashPadding,
+						sizeof(h_constHashPadding),
+						0, cudaMemcpyHostToDevice);
+}
+
+// Blake512 für 80 Byte grosse Eingangsdaten
+__host__ void quark_blake512_cpu_setBlock_80(void *pdata)
+{
+	// Message mit Padding bereitstellen
+	// lediglich die korrekte Nonce ist noch ab Byte 76 einzusetzen.
+	unsigned char PaddedMessage[128];
+	memcpy(PaddedMessage, pdata, 80);
+	memset(PaddedMessage+80, 0, 48);
+	PaddedMessage[80] = 0x80;
+	PaddedMessage[111] = 1;
+	PaddedMessage[126] = 0x02;
+	PaddedMessage[127] = 0x80;
+
+	// die Message zur Berechnung auf der GPU
+	cudaMemcpyToSymbol( c_PaddedMessage80, PaddedMessage, 16*sizeof(uint64_t), 0, cudaMemcpyHostToDevice);
+}
+
+__host__ void quark_blake512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_outputHash, int order)
+{
+	const int threadsperblock = 256;
+
+	// berechne wie viele Thread Blocks wir brauchen
+	dim3 grid((threads + threadsperblock-1)/threadsperblock);
+	dim3 block(threadsperblock);
+
+	// Größe des dynamischen Shared Memory Bereichs
+	size_t shared_size = 0;
+
+//	fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
+
+	quark_blake512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, d_nonceVector, (uint64_t*)d_outputHash);
+
+	// Strategisches Sleep Kommando zur Senkung der CPU Last
+	MyStreamSynchronize(NULL, order, thr_id);
+}
+
+__host__ void quark_blake512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_outputHash, int order)
+{
+	const int threadsperblock = 256;
+
+	// berechne wie viele Thread Blocks wir brauchen
+	dim3 grid((threads + threadsperblock-1)/threadsperblock);
+	dim3 block(threadsperblock);
+
+	// Größe des dynamischen Shared Memory Bereichs
+	size_t shared_size = 0;
+
+//	fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
+
+	quark_blake512_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, d_outputHash);
+
+	// Strategisches Sleep Kommando zur Senkung der CPU Last
+	MyStreamSynchronize(NULL, order, thr_id);
+}
diff --git a/quark/cuda_quark_compactionTest.cu b/quark/cuda_quark_compactionTest.cu
new file mode 100644
index 0000000..c21169c
--- /dev/null
+++ b/quark/cuda_quark_compactionTest.cu
@@ -0,0 +1,363 @@
+#include <cuda.h>
+#include "cuda_runtime.h"
+#include "device_launch_parameters.h"
+#include "sm_30_intrinsics.h"
+
+#include <stdio.h>
+#include <memory.h>
+#include <stdint.h>
+
+// aus cpu-miner.c
+extern "C" int device_map[8];
+
+// diese Struktur wird in der Init Funktion angefordert
+static cudaDeviceProp props[8];
+
+static uint32_t *d_tempBranch1Nonces[8];
+static uint32_t *d_numValid[8];
+static uint32_t *h_numValid[8];
+
+static uint32_t *d_partSum[2][8]; // für bis zu vier partielle Summen
+
+// aus heavy.cu
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+// True/False tester
+typedef uint32_t(*cuda_compactTestFunction_t)(uint32_t *inpHash);
+
+__device__ uint32_t QuarkTrueTest(uint32_t *inpHash)
+{
+	return ((inpHash[0] & 0x08) == 0x08);
+}
+
+__device__ uint32_t QuarkFalseTest(uint32_t *inpHash)
+{
+	return ((inpHash[0] & 0x08) == 0);
+}
+
+__device__ cuda_compactTestFunction_t d_QuarkTrueFunction = QuarkTrueTest, d_QuarkFalseFunction = QuarkFalseTest;
+cuda_compactTestFunction_t h_QuarkTrueFunction[8], h_QuarkFalseFunction[8];
+
+// Setup-Funktionen
+__host__ void quark_compactTest_cpu_init(int thr_id, int threads)
+{
+	cudaGetDeviceProperties(&props[thr_id], device_map[thr_id]);
+
+	cudaMemcpyFromSymbol(&h_QuarkTrueFunction[thr_id], d_QuarkTrueFunction, sizeof(cuda_compactTestFunction_t));
+	cudaMemcpyFromSymbol(&h_QuarkFalseFunction[thr_id], d_QuarkFalseFunction, sizeof(cuda_compactTestFunction_t));
+
+	// wir brauchen auch Speicherplatz auf dem Device
+	cudaMalloc(&d_tempBranch1Nonces[thr_id], sizeof(uint32_t) * threads * 2);	
+	cudaMalloc(&d_numValid[thr_id], 2*sizeof(uint32_t));
+	cudaMallocHost(&h_numValid[thr_id], 2*sizeof(uint32_t));
+
+	uint32_t s1;
+	s1 = (threads / 256) * 2;
+
+	cudaMalloc(&d_partSum[0][thr_id], sizeof(uint32_t) * s1); // BLOCKSIZE (Threads/Block)
+	cudaMalloc(&d_partSum[1][thr_id], sizeof(uint32_t) * s1); // BLOCKSIZE (Threads/Block)
+}
+
+// Die Summenfunktion (vom NVIDIA SDK)
+__global__ void quark_compactTest_gpu_SCAN(uint32_t *data, int width, uint32_t *partial_sums=NULL, cuda_compactTestFunction_t testFunc=NULL, int threads=0, uint32_t startNounce=0, uint32_t *inpHashes=NULL, uint32_t *d_validNonceTable=NULL)
+{
+	extern __shared__ uint32_t sums[];
+	int id = ((blockIdx.x * blockDim.x) + threadIdx.x);
+	//int lane_id = id % warpSize;
+	int lane_id = id % width;
+	// determine a warp_id within a block
+	 //int warp_id = threadIdx.x / warpSize;
+	int warp_id = threadIdx.x / width;
+
+	sums[lane_id] = 0;
+
+	// Below is the basic structure of using a shfl instruction
+	// for a scan.
+	// Record "value" as a variable - we accumulate it along the way
+	uint32_t value;
+	if(testFunc != NULL)
+	{
+		if (id < threads)
+		{
+			uint32_t *inpHash;
+			if(d_validNonceTable == NULL)
+			{
+				// keine Nonce-Liste
+				inpHash = &inpHashes[id<<4];
+			}else
+			{
+				// Nonce-Liste verfügbar
+				int nonce = d_validNonceTable[id] - startNounce;
+				inpHash = &inpHashes[nonce<<4];
+			}			
+			value = (*testFunc)(inpHash);
+		}else
+		{
+			value = 0;
+		}
+	}else
+	{
+		value = data[id];
+	}
+
+	__syncthreads();
+
+	// Now accumulate in log steps up the chain
+	// compute sums, with another thread's value who is
+	// distance delta away (i).  Note
+	// those threads where the thread 'i' away would have
+	// been out of bounds of the warp are unaffected.  This
+	// creates the scan sum.
+#pragma unroll
+
+	for (int i=1; i<=width; i*=2)
+	{
+		uint32_t n = __shfl_up((int)value, i, width);
+
+		if (lane_id >= i) value += n;
+	}
+
+	// value now holds the scan value for the individual thread
+	// next sum the largest values for each warp
+
+	// write the sum of the warp to smem
+	//if (threadIdx.x % warpSize == warpSize-1)
+	if (threadIdx.x % width == width-1)
+	{
+		sums[warp_id] = value;
+	}
+
+	__syncthreads();
+
+	//
+	// scan sum the warp sums
+	// the same shfl scan operation, but performed on warp sums
+	//
+	if (warp_id == 0)
+	{
+		uint32_t warp_sum = sums[lane_id];
+
+		for (int i=1; i<=width; i*=2)
+		{
+			uint32_t n = __shfl_up((int)warp_sum, i, width);
+
+		if (lane_id >= i) warp_sum += n;
+		}
+
+		sums[lane_id] = warp_sum;
+	}
+
+	__syncthreads();
+
+	// perform a uniform add across warps in the block
+	// read neighbouring warp's sum and add it to threads value
+	uint32_t blockSum = 0;
+
+	if (warp_id > 0)
+	{
+		blockSum = sums[warp_id-1];
+	}
+
+	value += blockSum;
+
+	// Now write out our result
+	data[id] = value;
+
+	// last thread has sum, write write out the block's sum
+	if (partial_sums != NULL && threadIdx.x == blockDim.x-1)
+	{
+		partial_sums[blockIdx.x] = value;
+	}
+}
+
+// Uniform add: add partial sums array
+__global__ void quark_compactTest_gpu_ADD(uint32_t *data, uint32_t *partial_sums, int len)
+{
+	__shared__ uint32_t buf;
+	int id = ((blockIdx.x * blockDim.x) + threadIdx.x);
+
+	if (id > len) return;
+
+	if (threadIdx.x == 0)
+	{
+		buf = partial_sums[blockIdx.x];
+	}
+
+	__syncthreads();
+	data[id] += buf;
+}
+
+// Der Scatter
+__global__ void quark_compactTest_gpu_SCATTER(uint32_t *sum, uint32_t *outp, cuda_compactTestFunction_t testFunc, int threads=0, uint32_t startNounce=0, uint32_t *inpHashes=NULL, uint32_t *d_validNonceTable=NULL)
+{
+	int id = ((blockIdx.x * blockDim.x) + threadIdx.x);
+	uint32_t actNounce = id;
+	uint32_t value;
+	if (id < threads)
+	{
+//		uint32_t nounce = startNounce + id;
+		uint32_t *inpHash;
+		if(d_validNonceTable == NULL)
+		{
+			// keine Nonce-Liste
+			inpHash = &inpHashes[id<<4];
+		}else
+		{
+			// Nonce-Liste verfügbar
+			int nonce = d_validNonceTable[id] - startNounce;
+			actNounce = nonce;
+			inpHash = &inpHashes[nonce<<4];
+		}
+
+		value = (*testFunc)(inpHash);
+	}else
+	{
+		value = 0;
+	}
+
+	if( value )
+	{
+		int idx = sum[id];
+		if(idx > 0)
+			outp[idx-1] = startNounce + actNounce;
+	}
+}
+
+__host__ static uint32_t quark_compactTest_roundUpExp(uint32_t val)
+{
+	if(val == 0)
+		return 0;
+
+	uint32_t mask = 0x80000000;
+	while( (val & mask) == 0 ) mask = mask >> 1;
+
+	if( (val & (~mask)) != 0 )
+		return mask << 1;
+
+	return mask;
+}
+
+__host__ void quark_compactTest_cpu_singleCompaction(int thr_id, int threads, uint32_t *nrm,
+														uint32_t *d_nonces1, cuda_compactTestFunction_t function,
+														uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable)
+{
+	int orgThreads = threads;
+	threads = (int)quark_compactTest_roundUpExp((uint32_t)threads);
+	// threadsPerBlock ausrechnen
+	int blockSize = 256;
+	int nSummen = threads / blockSize;
+
+	int thr1 = (threads+blockSize-1) / blockSize;
+	int thr2 = threads / (blockSize*blockSize);
+	int blockSize2 = (nSummen < blockSize) ? nSummen : blockSize;
+	int thr3 = (nSummen + blockSize2-1) / blockSize2;
+
+	bool callThrid = (thr2 > 0) ? true : false;
+
+	// Erster Initialscan
+	quark_compactTest_gpu_SCAN<<<thr1,blockSize, 32*sizeof(uint32_t)>>>(
+		d_tempBranch1Nonces[thr_id], 32, d_partSum[0][thr_id], function, orgThreads, startNounce, inpHashes, d_validNonceTable);	
+
+	// weitere Scans
+	if(callThrid)
+	{		
+		quark_compactTest_gpu_SCAN<<<thr2,blockSize, 32*sizeof(uint32_t)>>>(d_partSum[0][thr_id], 32, d_partSum[1][thr_id]);
+		quark_compactTest_gpu_SCAN<<<1, thr2, 32*sizeof(uint32_t)>>>(d_partSum[1][thr_id], (thr2>32) ? 32 : thr2);
+	}else
+	{
+		quark_compactTest_gpu_SCAN<<<thr3,blockSize2, 32*sizeof(uint32_t)>>>(d_partSum[0][thr_id], (blockSize2>32) ? 32 : blockSize2);
+	}
+
+	// Sync + Anzahl merken
+	cudaStreamSynchronize(NULL);
+
+	if(callThrid)
+		cudaMemcpy(nrm, &(d_partSum[1][thr_id])[thr2-1], sizeof(uint32_t), cudaMemcpyDeviceToHost);
+	else
+		cudaMemcpy(nrm, &(d_partSum[0][thr_id])[nSummen-1], sizeof(uint32_t), cudaMemcpyDeviceToHost);
+
+	
+	// Addieren
+	if(callThrid)
+	{
+		quark_compactTest_gpu_ADD<<<thr2-1, blockSize>>>(d_partSum[0][thr_id]+blockSize, d_partSum[1][thr_id], blockSize*thr2);
+	}
+	quark_compactTest_gpu_ADD<<<thr1-1, blockSize>>>(d_tempBranch1Nonces[thr_id]+blockSize, d_partSum[0][thr_id], threads);
+	
+	// Scatter
+	quark_compactTest_gpu_SCATTER<<<thr1,blockSize,0>>>(d_tempBranch1Nonces[thr_id], d_nonces1, 
+		function, orgThreads, startNounce, inpHashes, d_validNonceTable);
+
+	// Sync
+	cudaStreamSynchronize(NULL);
+}
+
+////// ACHTUNG: Diese funktion geht aktuell nur mit threads > 65536 (Am besten 256 * 1024 oder 256*2048)
+__host__ void quark_compactTest_cpu_dualCompaction(int thr_id, int threads, uint32_t *nrm,
+													 uint32_t *d_nonces1, uint32_t *d_nonces2,
+													 uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable)
+{
+	quark_compactTest_cpu_singleCompaction(thr_id, threads, &nrm[0], d_nonces1, h_QuarkTrueFunction[thr_id], startNounce, inpHashes, d_validNonceTable);
+	quark_compactTest_cpu_singleCompaction(thr_id, threads, &nrm[1], d_nonces2, h_QuarkFalseFunction[thr_id], startNounce, inpHashes, d_validNonceTable);
+
+	/*
+	// threadsPerBlock ausrechnen
+	int blockSize = 256;
+	int thr1 = threads / blockSize;
+	int thr2 = threads / (blockSize*blockSize);
+
+	// 1
+	quark_compactTest_gpu_SCAN<<<thr1,blockSize, 32*sizeof(uint32_t)>>>(d_tempBranch1Nonces[thr_id], 32, d_partSum1[thr_id], h_QuarkTrueFunction[thr_id], threads, startNounce, inpHashes);
+	quark_compactTest_gpu_SCAN<<<thr2,blockSize, 32*sizeof(uint32_t)>>>(d_partSum1[thr_id], 32, d_partSum2[thr_id]);
+	quark_compactTest_gpu_SCAN<<<1, thr2, 32*sizeof(uint32_t)>>>(d_partSum2[thr_id], (thr2>32) ? 32 : thr2);
+	cudaStreamSynchronize(NULL);
+	cudaMemcpy(&nrm[0], &(d_partSum2[thr_id])[thr2-1], sizeof(uint32_t), cudaMemcpyDeviceToHost);
+	quark_compactTest_gpu_ADD<<<thr2-1, blockSize>>>(d_partSum1[thr_id]+blockSize, d_partSum2[thr_id], blockSize*thr2);
+	quark_compactTest_gpu_ADD<<<thr1-1, blockSize>>>(d_tempBranch1Nonces[thr_id]+blockSize, d_partSum1[thr_id], threads);
+
+	// 2
+	quark_compactTest_gpu_SCAN<<<thr1,blockSize, 32*sizeof(uint32_t)>>>(d_tempBranch2Nonces[thr_id], 32, d_partSum1[thr_id], h_QuarkFalseFunction[thr_id], threads, startNounce, inpHashes);
+	quark_compactTest_gpu_SCAN<<<thr2,blockSize, 32*sizeof(uint32_t)>>>(d_partSum1[thr_id], 32, d_partSum2[thr_id]);
+	quark_compactTest_gpu_SCAN<<<1, thr2, 32*sizeof(uint32_t)>>>(d_partSum2[thr_id], (thr2>32) ? 32 : thr2);
+	cudaStreamSynchronize(NULL);
+	cudaMemcpy(&nrm[1], &(d_partSum2[thr_id])[thr2-1], sizeof(uint32_t), cudaMemcpyDeviceToHost);	
+	quark_compactTest_gpu_ADD<<<thr2-1, blockSize>>>(d_partSum1[thr_id]+blockSize, d_partSum2[thr_id], blockSize*thr2);
+	quark_compactTest_gpu_ADD<<<thr1-1, blockSize>>>(d_tempBranch2Nonces[thr_id]+blockSize, d_partSum1[thr_id], threads);
+	
+	// Hier ist noch eine Besonderheit: in d_tempBranch1Nonces sind die element von 1...nrm1 die Interessanten
+	// Schritt 3: Scatter
+	quark_compactTest_gpu_SCATTER<<<thr1,blockSize,0>>>(d_tempBranch1Nonces[thr_id], d_nonces1, h_QuarkTrueFunction[thr_id], threads, startNounce, inpHashes);
+	quark_compactTest_gpu_SCATTER<<<thr1,blockSize,0>>>(d_tempBranch2Nonces[thr_id], d_nonces2, h_QuarkFalseFunction[thr_id], threads, startNounce, inpHashes);
+	cudaStreamSynchronize(NULL);
+	*/
+}
+
+__host__ void quark_compactTest_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
+											uint32_t *d_nonces1, size_t *nrm1,
+											uint32_t *d_nonces2, size_t *nrm2,
+											int order)
+{
+	// Wenn validNonceTable genutzt wird, dann werden auch nur die Nonces betrachtet, die dort enthalten sind
+	// "threads" ist in diesem Fall auf die Länge dieses Array's zu setzen!
+	
+	quark_compactTest_cpu_dualCompaction(thr_id, threads,
+		h_numValid[thr_id], d_nonces1, d_nonces2,
+		startNounce, inpHashes, d_validNonceTable);
+
+	cudaStreamSynchronize(NULL); // Das original braucht zwar etwas CPU-Last, ist an dieser Stelle aber evtl besser
+	*nrm1 = (size_t)h_numValid[thr_id][0];
+	*nrm2 = (size_t)h_numValid[thr_id][1];
+}
+
+__host__ void quark_compactTest_single_false_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
+											uint32_t *d_nonces1, size_t *nrm1,
+											int order)
+{
+	// Wenn validNonceTable genutzt wird, dann werden auch nur die Nonces betrachtet, die dort enthalten sind
+	// "threads" ist in diesem Fall auf die Länge dieses Array's zu setzen!
+
+	quark_compactTest_cpu_singleCompaction(thr_id, threads, h_numValid[thr_id], d_nonces1, h_QuarkFalseFunction[thr_id], startNounce, inpHashes, d_validNonceTable);
+
+	cudaStreamSynchronize(NULL); // Das original braucht zwar etwas CPU-Last, ist an dieser Stelle aber evtl besser
+	*nrm1 = (size_t)h_numValid[thr_id][0];
+}
diff --git a/quark/cuda_quark_keccak512.cu b/quark/cuda_quark_keccak512.cu
new file mode 100644
index 0000000..c55a7a1
--- /dev/null
+++ b/quark/cuda_quark_keccak512.cu
@@ -0,0 +1,182 @@
+#include <cuda.h>
+#include "cuda_runtime.h"
+#include "device_launch_parameters.h"
+
+#include <stdio.h>
+#include <memory.h>
+
+// Folgende Definitionen später durch header ersetzen
+typedef unsigned char uint8_t;
+typedef unsigned int uint32_t;
+typedef unsigned long long uint64_t;
+
+// aus heavy.cu
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+#include "cuda_helper.h"
+
+#define U32TO64_LE(p) \
+    (((uint64_t)(*p)) | (((uint64_t)(*(p + 1))) << 32))
+
+#define U64TO32_LE(p, v) \
+    *p = (uint32_t)((v)); *(p+1) = (uint32_t)((v) >> 32);
+
+static const uint64_t host_keccak_round_constants[24] = {
+    0x0000000000000001ull, 0x0000000000008082ull,
+    0x800000000000808aull, 0x8000000080008000ull,
+    0x000000000000808bull, 0x0000000080000001ull,
+    0x8000000080008081ull, 0x8000000000008009ull,
+    0x000000000000008aull, 0x0000000000000088ull,
+    0x0000000080008009ull, 0x000000008000000aull,
+    0x000000008000808bull, 0x800000000000008bull,
+    0x8000000000008089ull, 0x8000000000008003ull,
+    0x8000000000008002ull, 0x8000000000000080ull,
+    0x000000000000800aull, 0x800000008000000aull,
+    0x8000000080008081ull, 0x8000000000008080ull,
+    0x0000000080000001ull, 0x8000000080008008ull
+};
+
+__constant__ uint64_t c_keccak_round_constants[24];
+
+static __device__ __forceinline__ void
+keccak_block(uint64_t *s, const uint32_t *in, const uint64_t *keccak_round_constants) {
+    size_t i;
+    uint64_t t[5], u[5], v, w;
+
+    /* absorb input */
+#pragma unroll 9
+    for (i = 0; i < 72 / 8; i++, in += 2)
+        s[i] ^= U32TO64_LE(in);
+    
+    for (i = 0; i < 24; i++) {
+        /* theta: c = a[0,i] ^ a[1,i] ^ .. a[4,i] */
+        t[0] = s[0] ^ s[5] ^ s[10] ^ s[15] ^ s[20];
+        t[1] = s[1] ^ s[6] ^ s[11] ^ s[16] ^ s[21];
+        t[2] = s[2] ^ s[7] ^ s[12] ^ s[17] ^ s[22];
+        t[3] = s[3] ^ s[8] ^ s[13] ^ s[18] ^ s[23];
+        t[4] = s[4] ^ s[9] ^ s[14] ^ s[19] ^ s[24];
+
+        /* theta: d[i] = c[i+4] ^ rotl(c[i+1],1) */
+        u[0] = t[4] ^ ROTL64(t[1], 1);
+        u[1] = t[0] ^ ROTL64(t[2], 1);
+        u[2] = t[1] ^ ROTL64(t[3], 1);
+        u[3] = t[2] ^ ROTL64(t[4], 1);
+        u[4] = t[3] ^ ROTL64(t[0], 1);
+
+        /* theta: a[0,i], a[1,i], .. a[4,i] ^= d[i] */
+        s[0] ^= u[0]; s[5] ^= u[0]; s[10] ^= u[0]; s[15] ^= u[0]; s[20] ^= u[0];
+        s[1] ^= u[1]; s[6] ^= u[1]; s[11] ^= u[1]; s[16] ^= u[1]; s[21] ^= u[1];
+        s[2] ^= u[2]; s[7] ^= u[2]; s[12] ^= u[2]; s[17] ^= u[2]; s[22] ^= u[2];
+        s[3] ^= u[3]; s[8] ^= u[3]; s[13] ^= u[3]; s[18] ^= u[3]; s[23] ^= u[3];
+        s[4] ^= u[4]; s[9] ^= u[4]; s[14] ^= u[4]; s[19] ^= u[4]; s[24] ^= u[4];
+
+        /* rho pi: b[..] = rotl(a[..], ..) */
+        v = s[ 1];
+        s[ 1] = ROTL64(s[ 6], 44);
+        s[ 6] = ROTL64(s[ 9], 20);
+        s[ 9] = ROTL64(s[22], 61);
+        s[22] = ROTL64(s[14], 39);
+        s[14] = ROTL64(s[20], 18);
+        s[20] = ROTL64(s[ 2], 62);
+        s[ 2] = ROTL64(s[12], 43);
+        s[12] = ROTL64(s[13], 25);
+        s[13] = ROTL64(s[19],  8);
+        s[19] = ROTL64(s[23], 56);
+        s[23] = ROTL64(s[15], 41);
+        s[15] = ROTL64(s[ 4], 27);
+        s[ 4] = ROTL64(s[24], 14);
+        s[24] = ROTL64(s[21],  2);
+        s[21] = ROTL64(s[ 8], 55);
+        s[ 8] = ROTL64(s[16], 45);
+        s[16] = ROTL64(s[ 5], 36);
+        s[ 5] = ROTL64(s[ 3], 28);
+        s[ 3] = ROTL64(s[18], 21);
+        s[18] = ROTL64(s[17], 15);
+        s[17] = ROTL64(s[11], 10);
+        s[11] = ROTL64(s[ 7],  6);
+        s[ 7] = ROTL64(s[10],  3);
+        s[10] = ROTL64(    v,  1);
+
+        /* chi: a[i,j] ^= ~b[i,j+1] & b[i,j+2] */
+        v = s[ 0]; w = s[ 1]; s[ 0] ^= (~w) & s[ 2]; s[ 1] ^= (~s[ 2]) & s[ 3]; s[ 2] ^= (~s[ 3]) & s[ 4]; s[ 3] ^= (~s[ 4]) & v; s[ 4] ^= (~v) & w;
+        v = s[ 5]; w = s[ 6]; s[ 5] ^= (~w) & s[ 7]; s[ 6] ^= (~s[ 7]) & s[ 8]; s[ 7] ^= (~s[ 8]) & s[ 9]; s[ 8] ^= (~s[ 9]) & v; s[ 9] ^= (~v) & w;
+        v = s[10]; w = s[11]; s[10] ^= (~w) & s[12]; s[11] ^= (~s[12]) & s[13]; s[12] ^= (~s[13]) & s[14]; s[13] ^= (~s[14]) & v; s[14] ^= (~v) & w;
+        v = s[15]; w = s[16]; s[15] ^= (~w) & s[17]; s[16] ^= (~s[17]) & s[18]; s[17] ^= (~s[18]) & s[19]; s[18] ^= (~s[19]) & v; s[19] ^= (~v) & w;
+        v = s[20]; w = s[21]; s[20] ^= (~w) & s[22]; s[21] ^= (~s[22]) & s[23]; s[22] ^= (~s[23]) & s[24]; s[23] ^= (~s[24]) & v; s[24] ^= (~v) & w;
+
+        /* iota: a[0,0] ^= round constant */
+        s[0] ^= keccak_round_constants[i];
+    }
+}
+
+__global__ void quark_keccak512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
+{
+    int thread = (blockDim.x * blockIdx.x + threadIdx.x);
+    if (thread < threads)
+    {
+        uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
+
+        int hashPosition = nounce - startNounce;
+        uint32_t *inpHash = (uint32_t*)&g_hash[8 * hashPosition];
+
+        // Nachricht kopieren
+        uint32_t message[18];
+#pragma unroll 16
+        for(int i=0;i<16;i++)
+            message[i] = inpHash[i];
+
+        message[16] = 0x01;
+        message[17] = 0x80000000;
+
+        // State initialisieren
+        uint64_t keccak_gpu_state[25];
+#pragma unroll 25
+        for (int i=0; i<25; i++)
+            keccak_gpu_state[i] = 0;
+
+        // den Block einmal gut durchschütteln
+        keccak_block(keccak_gpu_state, message, c_keccak_round_constants);
+
+        // das Hash erzeugen
+        uint32_t hash[16];
+
+#pragma unroll 8
+        for (size_t i = 0; i < 64; i += 8) {
+            U64TO32_LE((&hash[i/4]), keccak_gpu_state[i / 8]);
+        }
+
+        // fertig
+        uint32_t *outpHash = (uint32_t*)&g_hash[8 * hashPosition];
+
+#pragma unroll 16
+        for(int i=0;i<16;i++)
+            outpHash[i] = hash[i];
+    }
+}
+
+// Setup-Funktionen
+__host__ void quark_keccak512_cpu_init(int thr_id, int threads)
+{
+    // Kopiere die Hash-Tabellen in den GPU-Speicher
+    cudaMemcpyToSymbol( c_keccak_round_constants,
+                        host_keccak_round_constants,
+                        sizeof(host_keccak_round_constants),
+                        0, cudaMemcpyHostToDevice);
+}
+
+__host__ void quark_keccak512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
+{
+    const int threadsperblock = 256;
+
+    // berechne wie viele Thread Blocks wir brauchen
+    dim3 grid((threads + threadsperblock-1)/threadsperblock);
+    dim3 block(threadsperblock);
+
+    // Größe des dynamischen Shared Memory Bereichs
+    size_t shared_size = 0;
+
+//    fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
+
+    quark_keccak512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
+    MyStreamSynchronize(NULL, order, thr_id);
+}
diff --git a/quark/cuda_skein512.cu b/quark/cuda_skein512.cu
index 7e807ea..dc4030a 100644
--- a/quark/cuda_skein512.cu
+++ b/quark/cuda_skein512.cu
@@ -1,450 +1,450 @@
-#include <cuda.h>
-#include "cuda_runtime.h"
-#include "device_launch_parameters.h"
-
-#include <stdio.h>
-#include <memory.h>
-
-// Folgende Definitionen später durch header ersetzen
-typedef unsigned char uint8_t;
-typedef unsigned int uint32_t;
-typedef unsigned long long uint64_t;
-
-#define SPH_C64(x)    ((uint64_t)(x ## ULL))
-
-// aus cpu-miner.c
-extern "C" extern int device_map[8];
-// aus heavy.cu
-extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
-
-// Take a look at: https://www.schneier.com/skein1.3.pdf
-
-#if __CUDA_ARCH__ >= 350
-__forceinline__ __device__ uint64_t ROTL64(const uint64_t value, const int offset) {
-    uint2 result;
-    if(offset >= 32) {
-        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
-        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
-    } else {
-        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
-        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
-    }
-    return  __double_as_longlong(__hiloint2double(result.y, result.x));
-}
-#else
-#define ROTL64(x, n)        (((x) << (n)) | ((x) >> (64 - (n))))
-#endif
-#define SHL(x, n)			((x) << (n))
-#define SHR(x, n)			((x) >> (n))
-
-// Zum testen Hostcode...
-/*	Hier erstmal die Tabelle mit den Konstanten für die Mix-Funktion. Kann später vll.
-	mal direkt in den Code eingesetzt werden
-*/
-
-/*
- * M9_ ## s ## _ ## i  evaluates to s+i mod 9 (0 <= s <= 18, 0 <= i <= 7).
- */
-
-#define M9_0_0    0
-#define M9_0_1    1
-#define M9_0_2    2
-#define M9_0_3    3
-#define M9_0_4    4
-#define M9_0_5    5
-#define M9_0_6    6
-#define M9_0_7    7
-
-#define M9_1_0    1
-#define M9_1_1    2
-#define M9_1_2    3
-#define M9_1_3    4
-#define M9_1_4    5
-#define M9_1_5    6
-#define M9_1_6    7
-#define M9_1_7    8
-
-#define M9_2_0    2
-#define M9_2_1    3
-#define M9_2_2    4
-#define M9_2_3    5
-#define M9_2_4    6
-#define M9_2_5    7
-#define M9_2_6    8
-#define M9_2_7    0
-
-#define M9_3_0    3
-#define M9_3_1    4
-#define M9_3_2    5
-#define M9_3_3    6
-#define M9_3_4    7
-#define M9_3_5    8
-#define M9_3_6    0
-#define M9_3_7    1
-
-#define M9_4_0    4
-#define M9_4_1    5
-#define M9_4_2    6
-#define M9_4_3    7
-#define M9_4_4    8
-#define M9_4_5    0
-#define M9_4_6    1
-#define M9_4_7    2
-
-#define M9_5_0    5
-#define M9_5_1    6
-#define M9_5_2    7
-#define M9_5_3    8
-#define M9_5_4    0
-#define M9_5_5    1
-#define M9_5_6    2
-#define M9_5_7    3
-
-#define M9_6_0    6
-#define M9_6_1    7
-#define M9_6_2    8
-#define M9_6_3    0
-#define M9_6_4    1
-#define M9_6_5    2
-#define M9_6_6    3
-#define M9_6_7    4
-
-#define M9_7_0    7
-#define M9_7_1    8
-#define M9_7_2    0
-#define M9_7_3    1
-#define M9_7_4    2
-#define M9_7_5    3
-#define M9_7_6    4
-#define M9_7_7    5
-
-#define M9_8_0    8
-#define M9_8_1    0
-#define M9_8_2    1
-#define M9_8_3    2
-#define M9_8_4    3
-#define M9_8_5    4
-#define M9_8_6    5
-#define M9_8_7    6
-
-#define M9_9_0    0
-#define M9_9_1    1
-#define M9_9_2    2
-#define M9_9_3    3
-#define M9_9_4    4
-#define M9_9_5    5
-#define M9_9_6    6
-#define M9_9_7    7
-
-#define M9_10_0   1
-#define M9_10_1   2
-#define M9_10_2   3
-#define M9_10_3   4
-#define M9_10_4   5
-#define M9_10_5   6
-#define M9_10_6   7
-#define M9_10_7   8
-
-#define M9_11_0   2
-#define M9_11_1   3
-#define M9_11_2   4
-#define M9_11_3   5
-#define M9_11_4   6
-#define M9_11_5   7
-#define M9_11_6   8
-#define M9_11_7   0
-
-#define M9_12_0   3
-#define M9_12_1   4
-#define M9_12_2   5
-#define M9_12_3   6
-#define M9_12_4   7
-#define M9_12_5   8
-#define M9_12_6   0
-#define M9_12_7   1
-
-#define M9_13_0   4
-#define M9_13_1   5
-#define M9_13_2   6
-#define M9_13_3   7
-#define M9_13_4   8
-#define M9_13_5   0
-#define M9_13_6   1
-#define M9_13_7   2
-
-#define M9_14_0   5
-#define M9_14_1   6
-#define M9_14_2   7
-#define M9_14_3   8
-#define M9_14_4   0
-#define M9_14_5   1
-#define M9_14_6   2
-#define M9_14_7   3
-
-#define M9_15_0   6
-#define M9_15_1   7
-#define M9_15_2   8
-#define M9_15_3   0
-#define M9_15_4   1
-#define M9_15_5   2
-#define M9_15_6   3
-#define M9_15_7   4
-
-#define M9_16_0   7
-#define M9_16_1   8
-#define M9_16_2   0
-#define M9_16_3   1
-#define M9_16_4   2
-#define M9_16_5   3
-#define M9_16_6   4
-#define M9_16_7   5
-
-#define M9_17_0   8
-#define M9_17_1   0
-#define M9_17_2   1
-#define M9_17_3   2
-#define M9_17_4   3
-#define M9_17_5   4
-#define M9_17_6   5
-#define M9_17_7   6
-
-#define M9_18_0   0
-#define M9_18_1   1
-#define M9_18_2   2
-#define M9_18_3   3
-#define M9_18_4   4
-#define M9_18_5   5
-#define M9_18_6   6
-#define M9_18_7   7
-
-/*
- * M3_ ## s ## _ ## i  evaluates to s+i mod 3 (0 <= s <= 18, 0 <= i <= 1).
- */
-
-#define M3_0_0    0
-#define M3_0_1    1
-#define M3_1_0    1
-#define M3_1_1    2
-#define M3_2_0    2
-#define M3_2_1    0
-#define M3_3_0    0
-#define M3_3_1    1
-#define M3_4_0    1
-#define M3_4_1    2
-#define M3_5_0    2
-#define M3_5_1    0
-#define M3_6_0    0
-#define M3_6_1    1
-#define M3_7_0    1
-#define M3_7_1    2
-#define M3_8_0    2
-#define M3_8_1    0
-#define M3_9_0    0
-#define M3_9_1    1
-#define M3_10_0   1
-#define M3_10_1   2
-#define M3_11_0   2
-#define M3_11_1   0
-#define M3_12_0   0
-#define M3_12_1   1
-#define M3_13_0   1
-#define M3_13_1   2
-#define M3_14_0   2
-#define M3_14_1   0
-#define M3_15_0   0
-#define M3_15_1   1
-#define M3_16_0   1
-#define M3_16_1   2
-#define M3_17_0   2
-#define M3_17_1   0
-#define M3_18_0   0
-#define M3_18_1   1
-
-#define XCAT(x, y)     XCAT_(x, y)
-#define XCAT_(x, y)    x ## y
-
-#define SKBI(k, s, i)   XCAT(k, XCAT(XCAT(XCAT(M9_, s), _), i))
-#define SKBT(t, s, v)   XCAT(t, XCAT(XCAT(XCAT(M3_, s), _), v))
-
-#define TFBIG_KINIT(k0, k1, k2, k3, k4, k5, k6, k7, k8, t0, t1, t2) { \
-		k8 = ((k0 ^ k1) ^ (k2 ^ k3)) ^ ((k4 ^ k5) ^ (k6 ^ k7)) \
-			^ SPH_C64(0x1BD11BDAA9FC1A22); \
-		t2 = t0 ^ t1; \
-	}
-
-#define TFBIG_ADDKEY(w0, w1, w2, w3, w4, w5, w6, w7, k, t, s) { \
-		w0 = (w0 + SKBI(k, s, 0)); \
-		w1 = (w1 + SKBI(k, s, 1)); \
-		w2 = (w2 + SKBI(k, s, 2)); \
-		w3 = (w3 + SKBI(k, s, 3)); \
-		w4 = (w4 + SKBI(k, s, 4)); \
-		w5 = (w5 + SKBI(k, s, 5) + SKBT(t, s, 0)); \
-		w6 = (w6 + SKBI(k, s, 6) + SKBT(t, s, 1)); \
-		w7 = (w7 + SKBI(k, s, 7) + (uint64_t)s); \
-	}
-
-#define TFBIG_MIX(x0, x1, rc) { \
-		x0 = x0 + x1; \
-		x1 = ROTL64(x1, rc) ^ x0; \
-	}
-
-#define TFBIG_MIX8(w0, w1, w2, w3, w4, w5, w6, w7, rc0, rc1, rc2, rc3) { \
-		TFBIG_MIX(w0, w1, rc0); \
-		TFBIG_MIX(w2, w3, rc1); \
-		TFBIG_MIX(w4, w5, rc2); \
-		TFBIG_MIX(w6, w7, rc3); \
-	}
-
-#define TFBIG_4e(s)  { \
-		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \
-		TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 46, 36, 19, 37); \
-		TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 33, 27, 14, 42); \
-		TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 17, 49, 36, 39); \
-		TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 44,  9, 54, 56); \
-	}
-
-#define TFBIG_4o(s)  { \
-		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \
-		TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 39, 30, 34, 24); \
-		TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 13, 50, 10, 17); \
-		TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 25, 29, 39, 43); \
-		TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3],  8, 35, 56, 22); \
-	}
-
-static __constant__ uint64_t d_constMem[8];
-static uint64_t h_constMem[8] = {
-	SPH_C64(0x4903ADFF749C51CE),
-	SPH_C64(0x0D95DE399746DF03),
-	SPH_C64(0x8FD1934127C79BCE),
-	SPH_C64(0x9A255629FF352CB1),
-	SPH_C64(0x5DB62599DF6CA7B0),
-	SPH_C64(0xEABE394CA9D5C3F4),
-	SPH_C64(0x991112C71A75B523),
-	SPH_C64(0xAE18A40B660FCC33) };
-
-__global__ void quark_skein512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
-{
-	int thread = (blockDim.x * blockIdx.x + threadIdx.x);
-	if (thread < threads)
-	{
-		// Skein
-		uint64_t p[8];
-		uint64_t h0, h1, h2, h3, h4, h5, h6, h7, h8;
-		uint64_t t0, t1, t2;
-
-		uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
-
-		int hashPosition = nounce - startNounce;
-		uint64_t *inpHash = &g_hash[8 * hashPosition];
-
-		// Initialisierung
-		h0 = d_constMem[0];
-		h1 = d_constMem[1];
-		h2 = d_constMem[2];
-		h3 = d_constMem[3];
-		h4 = d_constMem[4];
-		h5 = d_constMem[5];
-		h6 = d_constMem[6];
-		h7 = d_constMem[7];
-
-		// 1. Runde -> etype = 480, ptr = 64, bcount = 0, data = msg		
-#pragma unroll 8
-		for(int i=0;i<8;i++)
-			p[i] = inpHash[i];
-
-		t0 = 64; // ptr
-		t1 = 480ull << 55; // etype
-		TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
-		TFBIG_4e(0);
-		TFBIG_4o(1);
-		TFBIG_4e(2);
-		TFBIG_4o(3);
-		TFBIG_4e(4);
-		TFBIG_4o(5);
-		TFBIG_4e(6);
-		TFBIG_4o(7);
-		TFBIG_4e(8);
-		TFBIG_4o(9);
-		TFBIG_4e(10);
-		TFBIG_4o(11);
-		TFBIG_4e(12);
-		TFBIG_4o(13);
-		TFBIG_4e(14);
-		TFBIG_4o(15);
-		TFBIG_4e(16);
-		TFBIG_4o(17);
-		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18);
-
-		h0 = inpHash[0] ^ p[0];
-		h1 = inpHash[1] ^ p[1];
-		h2 = inpHash[2] ^ p[2];
-		h3 = inpHash[3] ^ p[3];
-		h4 = inpHash[4] ^ p[4];
-		h5 = inpHash[5] ^ p[5];
-		h6 = inpHash[6] ^ p[6];
-		h7 = inpHash[7] ^ p[7];
-
-		// 2. Runde -> etype = 510, ptr = 8, bcount = 0, data = 0
-#pragma unroll 8
-		for(int i=0;i<8;i++)
-			p[i] = 0;
-
-		t0 = 8; // ptr
-		t1 = 510ull << 55; // etype
-		TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
-		TFBIG_4e(0);
-		TFBIG_4o(1);
-		TFBIG_4e(2);
-		TFBIG_4o(3);
-		TFBIG_4e(4);
-		TFBIG_4o(5);
-		TFBIG_4e(6);
-		TFBIG_4o(7);
-		TFBIG_4e(8);
-		TFBIG_4o(9);
-		TFBIG_4e(10);
-		TFBIG_4o(11);
-		TFBIG_4e(12);
-		TFBIG_4o(13);
-		TFBIG_4e(14);
-		TFBIG_4o(15);
-		TFBIG_4e(16);
-		TFBIG_4o(17);
-		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18);
-
-		// fertig
-		uint64_t *outpHash = &g_hash[8 * hashPosition];
-
-#pragma unroll 8
-		for(int i=0;i<8;i++)
-			outpHash[i] = p[i];
-	}
-}
-
-// Setup-Funktionen
-__host__ void quark_skein512_cpu_init(int thr_id, int threads)
-{
-	// nix zu tun ;-)
-	cudaMemcpyToSymbol( d_constMem,
-                        h_constMem,
-                        sizeof(h_constMem),
-                        0, cudaMemcpyHostToDevice);
-}
-
-__host__ void quark_skein512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
-{
-	const int threadsperblock = 256;
-
-	// berechne wie viele Thread Blocks wir brauchen
-	dim3 grid((threads + threadsperblock-1)/threadsperblock);
-	dim3 block(threadsperblock);
-
-	// Größe des dynamischen Shared Memory Bereichs
-	size_t shared_size = 0;
-
-//	fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
-	quark_skein512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
-
-	// Strategisches Sleep Kommando zur Senkung der CPU Last
-	MyStreamSynchronize(NULL, order, thr_id);
-}
+#include <cuda.h>
+#include "cuda_runtime.h"
+#include "device_launch_parameters.h"
+
+#include <stdio.h>
+#include <memory.h>
+
+// Folgende Definitionen später durch header ersetzen
+typedef unsigned char uint8_t;
+typedef unsigned int uint32_t;
+typedef unsigned long long uint64_t;
+
+#define SPH_C64(x)    ((uint64_t)(x ## ULL))
+
+// aus cpu-miner.c
+extern "C" extern int device_map[8];
+// aus heavy.cu
+extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
+
+// Take a look at: https://www.schneier.com/skein1.3.pdf
+
+#if __CUDA_ARCH__ >= 350
+__forceinline__ __device__ uint64_t ROTL64(const uint64_t value, const int offset) {
+    uint2 result;
+    if(offset >= 32) {
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
+    } else {
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(__double2hiint(__longlong_as_double(value))), "r"(__double2loint(__longlong_as_double(value))), "r"(offset));
+        asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(__double2loint(__longlong_as_double(value))), "r"(__double2hiint(__longlong_as_double(value))), "r"(offset));
+    }
+    return  __double_as_longlong(__hiloint2double(result.y, result.x));
+}
+#else
+#define ROTL64(x, n)        (((x) << (n)) | ((x) >> (64 - (n))))
+#endif
+#define SHL(x, n)			((x) << (n))
+#define SHR(x, n)			((x) >> (n))
+
+// Zum testen Hostcode...
+/*	Hier erstmal die Tabelle mit den Konstanten für die Mix-Funktion. Kann später vll.
+	mal direkt in den Code eingesetzt werden
+*/
+
+/*
+ * M9_ ## s ## _ ## i  evaluates to s+i mod 9 (0 <= s <= 18, 0 <= i <= 7).
+ */
+
+#define M9_0_0    0
+#define M9_0_1    1
+#define M9_0_2    2
+#define M9_0_3    3
+#define M9_0_4    4
+#define M9_0_5    5
+#define M9_0_6    6
+#define M9_0_7    7
+
+#define M9_1_0    1
+#define M9_1_1    2
+#define M9_1_2    3
+#define M9_1_3    4
+#define M9_1_4    5
+#define M9_1_5    6
+#define M9_1_6    7
+#define M9_1_7    8
+
+#define M9_2_0    2
+#define M9_2_1    3
+#define M9_2_2    4
+#define M9_2_3    5
+#define M9_2_4    6
+#define M9_2_5    7
+#define M9_2_6    8
+#define M9_2_7    0
+
+#define M9_3_0    3
+#define M9_3_1    4
+#define M9_3_2    5
+#define M9_3_3    6
+#define M9_3_4    7
+#define M9_3_5    8
+#define M9_3_6    0
+#define M9_3_7    1
+
+#define M9_4_0    4
+#define M9_4_1    5
+#define M9_4_2    6
+#define M9_4_3    7
+#define M9_4_4    8
+#define M9_4_5    0
+#define M9_4_6    1
+#define M9_4_7    2
+
+#define M9_5_0    5
+#define M9_5_1    6
+#define M9_5_2    7
+#define M9_5_3    8
+#define M9_5_4    0
+#define M9_5_5    1
+#define M9_5_6    2
+#define M9_5_7    3
+
+#define M9_6_0    6
+#define M9_6_1    7
+#define M9_6_2    8
+#define M9_6_3    0
+#define M9_6_4    1
+#define M9_6_5    2
+#define M9_6_6    3
+#define M9_6_7    4
+
+#define M9_7_0    7
+#define M9_7_1    8
+#define M9_7_2    0
+#define M9_7_3    1
+#define M9_7_4    2
+#define M9_7_5    3
+#define M9_7_6    4
+#define M9_7_7    5
+
+#define M9_8_0    8
+#define M9_8_1    0
+#define M9_8_2    1
+#define M9_8_3    2
+#define M9_8_4    3
+#define M9_8_5    4
+#define M9_8_6    5
+#define M9_8_7    6
+
+#define M9_9_0    0
+#define M9_9_1    1
+#define M9_9_2    2
+#define M9_9_3    3
+#define M9_9_4    4
+#define M9_9_5    5
+#define M9_9_6    6
+#define M9_9_7    7
+
+#define M9_10_0   1
+#define M9_10_1   2
+#define M9_10_2   3
+#define M9_10_3   4
+#define M9_10_4   5
+#define M9_10_5   6
+#define M9_10_6   7
+#define M9_10_7   8
+
+#define M9_11_0   2
+#define M9_11_1   3
+#define M9_11_2   4
+#define M9_11_3   5
+#define M9_11_4   6
+#define M9_11_5   7
+#define M9_11_6   8
+#define M9_11_7   0
+
+#define M9_12_0   3
+#define M9_12_1   4
+#define M9_12_2   5
+#define M9_12_3   6
+#define M9_12_4   7
+#define M9_12_5   8
+#define M9_12_6   0
+#define M9_12_7   1
+
+#define M9_13_0   4
+#define M9_13_1   5
+#define M9_13_2   6
+#define M9_13_3   7
+#define M9_13_4   8
+#define M9_13_5   0
+#define M9_13_6   1
+#define M9_13_7   2
+
+#define M9_14_0   5
+#define M9_14_1   6
+#define M9_14_2   7
+#define M9_14_3   8
+#define M9_14_4   0
+#define M9_14_5   1
+#define M9_14_6   2
+#define M9_14_7   3
+
+#define M9_15_0   6
+#define M9_15_1   7
+#define M9_15_2   8
+#define M9_15_3   0
+#define M9_15_4   1
+#define M9_15_5   2
+#define M9_15_6   3
+#define M9_15_7   4
+
+#define M9_16_0   7
+#define M9_16_1   8
+#define M9_16_2   0
+#define M9_16_3   1
+#define M9_16_4   2
+#define M9_16_5   3
+#define M9_16_6   4
+#define M9_16_7   5
+
+#define M9_17_0   8
+#define M9_17_1   0
+#define M9_17_2   1
+#define M9_17_3   2
+#define M9_17_4   3
+#define M9_17_5   4
+#define M9_17_6   5
+#define M9_17_7   6
+
+#define M9_18_0   0
+#define M9_18_1   1
+#define M9_18_2   2
+#define M9_18_3   3
+#define M9_18_4   4
+#define M9_18_5   5
+#define M9_18_6   6
+#define M9_18_7   7
+
+/*
+ * M3_ ## s ## _ ## i  evaluates to s+i mod 3 (0 <= s <= 18, 0 <= i <= 1).
+ */
+
+#define M3_0_0    0
+#define M3_0_1    1
+#define M3_1_0    1
+#define M3_1_1    2
+#define M3_2_0    2
+#define M3_2_1    0
+#define M3_3_0    0
+#define M3_3_1    1
+#define M3_4_0    1
+#define M3_4_1    2
+#define M3_5_0    2
+#define M3_5_1    0
+#define M3_6_0    0
+#define M3_6_1    1
+#define M3_7_0    1
+#define M3_7_1    2
+#define M3_8_0    2
+#define M3_8_1    0
+#define M3_9_0    0
+#define M3_9_1    1
+#define M3_10_0   1
+#define M3_10_1   2
+#define M3_11_0   2
+#define M3_11_1   0
+#define M3_12_0   0
+#define M3_12_1   1
+#define M3_13_0   1
+#define M3_13_1   2
+#define M3_14_0   2
+#define M3_14_1   0
+#define M3_15_0   0
+#define M3_15_1   1
+#define M3_16_0   1
+#define M3_16_1   2
+#define M3_17_0   2
+#define M3_17_1   0
+#define M3_18_0   0
+#define M3_18_1   1
+
+#define XCAT(x, y)     XCAT_(x, y)
+#define XCAT_(x, y)    x ## y
+
+#define SKBI(k, s, i)   XCAT(k, XCAT(XCAT(XCAT(M9_, s), _), i))
+#define SKBT(t, s, v)   XCAT(t, XCAT(XCAT(XCAT(M3_, s), _), v))
+
+#define TFBIG_KINIT(k0, k1, k2, k3, k4, k5, k6, k7, k8, t0, t1, t2) { \
+		k8 = ((k0 ^ k1) ^ (k2 ^ k3)) ^ ((k4 ^ k5) ^ (k6 ^ k7)) \
+			^ SPH_C64(0x1BD11BDAA9FC1A22); \
+		t2 = t0 ^ t1; \
+	}
+
+#define TFBIG_ADDKEY(w0, w1, w2, w3, w4, w5, w6, w7, k, t, s) { \
+		w0 = (w0 + SKBI(k, s, 0)); \
+		w1 = (w1 + SKBI(k, s, 1)); \
+		w2 = (w2 + SKBI(k, s, 2)); \
+		w3 = (w3 + SKBI(k, s, 3)); \
+		w4 = (w4 + SKBI(k, s, 4)); \
+		w5 = (w5 + SKBI(k, s, 5) + SKBT(t, s, 0)); \
+		w6 = (w6 + SKBI(k, s, 6) + SKBT(t, s, 1)); \
+		w7 = (w7 + SKBI(k, s, 7) + (uint64_t)s); \
+	}
+
+#define TFBIG_MIX(x0, x1, rc) { \
+		x0 = x0 + x1; \
+		x1 = ROTL64(x1, rc) ^ x0; \
+	}
+
+#define TFBIG_MIX8(w0, w1, w2, w3, w4, w5, w6, w7, rc0, rc1, rc2, rc3) { \
+		TFBIG_MIX(w0, w1, rc0); \
+		TFBIG_MIX(w2, w3, rc1); \
+		TFBIG_MIX(w4, w5, rc2); \
+		TFBIG_MIX(w6, w7, rc3); \
+	}
+
+#define TFBIG_4e(s)  { \
+		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \
+		TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 46, 36, 19, 37); \
+		TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 33, 27, 14, 42); \
+		TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 17, 49, 36, 39); \
+		TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3], 44,  9, 54, 56); \
+	}
+
+#define TFBIG_4o(s)  { \
+		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, s); \
+		TFBIG_MIX8(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], 39, 30, 34, 24); \
+		TFBIG_MIX8(p[2], p[1], p[4], p[7], p[6], p[5], p[0], p[3], 13, 50, 10, 17); \
+		TFBIG_MIX8(p[4], p[1], p[6], p[3], p[0], p[5], p[2], p[7], 25, 29, 39, 43); \
+		TFBIG_MIX8(p[6], p[1], p[0], p[7], p[2], p[5], p[4], p[3],  8, 35, 56, 22); \
+	}
+
+static __constant__ uint64_t d_constMem[8];
+static uint64_t h_constMem[8] = {
+	SPH_C64(0x4903ADFF749C51CE),
+	SPH_C64(0x0D95DE399746DF03),
+	SPH_C64(0x8FD1934127C79BCE),
+	SPH_C64(0x9A255629FF352CB1),
+	SPH_C64(0x5DB62599DF6CA7B0),
+	SPH_C64(0xEABE394CA9D5C3F4),
+	SPH_C64(0x991112C71A75B523),
+	SPH_C64(0xAE18A40B660FCC33) };
+
+__global__ void quark_skein512_gpu_hash_64(int threads, uint32_t startNounce, uint64_t *g_hash, uint32_t *g_nonceVector)
+{
+	int thread = (blockDim.x * blockIdx.x + threadIdx.x);
+	if (thread < threads)
+	{
+		// Skein
+		uint64_t p[8];
+		uint64_t h0, h1, h2, h3, h4, h5, h6, h7, h8;
+		uint64_t t0, t1, t2;
+
+		uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
+
+		int hashPosition = nounce - startNounce;
+		uint64_t *inpHash = &g_hash[8 * hashPosition];
+
+		// Initialisierung
+		h0 = d_constMem[0];
+		h1 = d_constMem[1];
+		h2 = d_constMem[2];
+		h3 = d_constMem[3];
+		h4 = d_constMem[4];
+		h5 = d_constMem[5];
+		h6 = d_constMem[6];
+		h7 = d_constMem[7];
+
+		// 1. Runde -> etype = 480, ptr = 64, bcount = 0, data = msg		
+#pragma unroll 8
+		for(int i=0;i<8;i++)
+			p[i] = inpHash[i];
+
+		t0 = 64; // ptr
+		t1 = 480ull << 55; // etype
+		TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
+		TFBIG_4e(0);
+		TFBIG_4o(1);
+		TFBIG_4e(2);
+		TFBIG_4o(3);
+		TFBIG_4e(4);
+		TFBIG_4o(5);
+		TFBIG_4e(6);
+		TFBIG_4o(7);
+		TFBIG_4e(8);
+		TFBIG_4o(9);
+		TFBIG_4e(10);
+		TFBIG_4o(11);
+		TFBIG_4e(12);
+		TFBIG_4o(13);
+		TFBIG_4e(14);
+		TFBIG_4o(15);
+		TFBIG_4e(16);
+		TFBIG_4o(17);
+		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18);
+
+		h0 = inpHash[0] ^ p[0];
+		h1 = inpHash[1] ^ p[1];
+		h2 = inpHash[2] ^ p[2];
+		h3 = inpHash[3] ^ p[3];
+		h4 = inpHash[4] ^ p[4];
+		h5 = inpHash[5] ^ p[5];
+		h6 = inpHash[6] ^ p[6];
+		h7 = inpHash[7] ^ p[7];
+
+		// 2. Runde -> etype = 510, ptr = 8, bcount = 0, data = 0
+#pragma unroll 8
+		for(int i=0;i<8;i++)
+			p[i] = 0;
+
+		t0 = 8; // ptr
+		t1 = 510ull << 55; // etype
+		TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2);
+		TFBIG_4e(0);
+		TFBIG_4o(1);
+		TFBIG_4e(2);
+		TFBIG_4o(3);
+		TFBIG_4e(4);
+		TFBIG_4o(5);
+		TFBIG_4e(6);
+		TFBIG_4o(7);
+		TFBIG_4e(8);
+		TFBIG_4o(9);
+		TFBIG_4e(10);
+		TFBIG_4o(11);
+		TFBIG_4e(12);
+		TFBIG_4o(13);
+		TFBIG_4e(14);
+		TFBIG_4o(15);
+		TFBIG_4e(16);
+		TFBIG_4o(17);
+		TFBIG_ADDKEY(p[0], p[1], p[2], p[3], p[4], p[5], p[6], p[7], h, t, 18);
+
+		// fertig
+		uint64_t *outpHash = &g_hash[8 * hashPosition];
+
+#pragma unroll 8
+		for(int i=0;i<8;i++)
+			outpHash[i] = p[i];
+	}
+}
+
+// Setup-Funktionen
+__host__ void quark_skein512_cpu_init(int thr_id, int threads)
+{
+	// nix zu tun ;-)
+	cudaMemcpyToSymbol( d_constMem,
+                        h_constMem,
+                        sizeof(h_constMem),
+                        0, cudaMemcpyHostToDevice);
+}
+
+__host__ void quark_skein512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order)
+{
+	const int threadsperblock = 256;
+
+	// berechne wie viele Thread Blocks wir brauchen
+	dim3 grid((threads + threadsperblock-1)/threadsperblock);
+	dim3 block(threadsperblock);
+
+	// Größe des dynamischen Shared Memory Bereichs
+	size_t shared_size = 0;
+
+//	fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
+	quark_skein512_gpu_hash_64<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash, d_nonceVector);
+
+	// Strategisches Sleep Kommando zur Senkung der CPU Last
+	MyStreamSynchronize(NULL, order, thr_id);
+}
diff --git a/quark/quarkcoin.cu b/quark/quarkcoin.cu
new file mode 100644
index 0000000..acfe731
--- /dev/null
+++ b/quark/quarkcoin.cu
@@ -0,0 +1,274 @@
+
+extern "C"
+{
+#include "sph/sph_blake.h"
+#include "sph/sph_bmw.h"
+#include "sph/sph_groestl.h"
+#include "sph/sph_skein.h"
+#include "sph/sph_jh.h"
+#include "sph/sph_keccak.h"
+#include "miner.h"
+}
+
+#include <stdint.h>
+
+// aus cpu-miner.c
+extern int device_map[8];
+
+// Speicher für Input/Output der verketteten Hashfunktionen
+static uint32_t *d_hash[8];
+
+// Speicher zur Generierung der Noncevektoren für die bedingten Hashes
+static uint32_t *d_quarkNonces[8];
+static uint32_t *d_branch1Nonces[8];
+static uint32_t *d_branch2Nonces[8];
+static uint32_t *d_branch3Nonces[8];
+
+extern void quark_blake512_cpu_init(int thr_id, int threads);
+extern void quark_blake512_cpu_setBlock_80(void *pdata);
+extern void quark_blake512_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_hash, int order);
+extern void quark_blake512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_bmw512_cpu_init(int thr_id, int threads);
+extern void quark_bmw512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_groestl512_cpu_init(int thr_id, int threads);
+extern void quark_groestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+extern void quark_doublegroestl512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_skein512_cpu_init(int thr_id, int threads);
+extern void quark_skein512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_keccak512_cpu_init(int thr_id, int threads);
+extern void quark_keccak512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_jh512_cpu_init(int thr_id, int threads);
+extern void quark_jh512_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void quark_check_cpu_init(int thr_id, int threads);
+extern void quark_check_cpu_setTarget(const void *ptarget);
+extern uint32_t quark_check_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order);
+
+extern void quark_compactTest_cpu_init(int thr_id, int threads);
+extern void quark_compactTest_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
+											uint32_t *d_nonces1, size_t *nrm1,
+											uint32_t *d_nonces2, size_t *nrm2,
+											int order);
+extern void quark_compactTest_single_false_cpu_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *inpHashes, uint32_t *d_validNonceTable,
+											uint32_t *d_nonces1, size_t *nrm1,
+											int order);
+
+// Original Quarkhash Funktion aus einem miner Quelltext
+inline void quarkhash(void *state, const void *input)
+{
+    sph_blake512_context ctx_blake;
+    sph_bmw512_context ctx_bmw;
+    sph_groestl512_context ctx_groestl;
+    sph_jh512_context ctx_jh;
+    sph_keccak512_context ctx_keccak;
+    sph_skein512_context ctx_skein;
+    
+    unsigned char hash[64];
+
+    sph_blake512_init(&ctx_blake);
+    // ZBLAKE;
+    sph_blake512 (&ctx_blake, input, 80);
+    sph_blake512_close(&ctx_blake, (void*) hash);
+    
+    sph_bmw512_init(&ctx_bmw);
+    // ZBMW;
+    sph_bmw512 (&ctx_bmw, (const void*) hash, 64);
+    sph_bmw512_close(&ctx_bmw, (void*) hash);
+
+    if (hash[0] & 0x8)
+    {
+        sph_groestl512_init(&ctx_groestl);
+        // ZGROESTL;
+        sph_groestl512 (&ctx_groestl, (const void*) hash, 64);
+        sph_groestl512_close(&ctx_groestl, (void*) hash);
+    }
+    else
+    {
+        sph_skein512_init(&ctx_skein);
+        // ZSKEIN;
+        sph_skein512 (&ctx_skein, (const void*) hash, 64);
+        sph_skein512_close(&ctx_skein, (void*) hash);
+    }
+    
+    sph_groestl512_init(&ctx_groestl);
+    // ZGROESTL;
+    sph_groestl512 (&ctx_groestl, (const void*) hash, 64);
+    sph_groestl512_close(&ctx_groestl, (void*) hash);
+
+    sph_jh512_init(&ctx_jh);
+    // ZJH;
+    sph_jh512 (&ctx_jh, (const void*) hash, 64);
+    sph_jh512_close(&ctx_jh, (void*) hash);
+
+    if (hash[0] & 0x8)
+    {
+        sph_blake512_init(&ctx_blake);
+        // ZBLAKE;
+        sph_blake512 (&ctx_blake, (const void*) hash, 64);
+        sph_blake512_close(&ctx_blake, (void*) hash);
+    }
+    else
+    {
+        sph_bmw512_init(&ctx_bmw);
+        // ZBMW;
+        sph_bmw512 (&ctx_bmw, (const void*) hash, 64);
+        sph_bmw512_close(&ctx_bmw, (void*) hash);
+    }
+
+    sph_keccak512_init(&ctx_keccak);
+    // ZKECCAK;
+    sph_keccak512 (&ctx_keccak, (const void*) hash, 64);
+    sph_keccak512_close(&ctx_keccak, (void*) hash);
+
+    sph_skein512_init(&ctx_skein);
+    // SKEIN;
+    sph_skein512 (&ctx_skein, (const void*) hash, 64);
+    sph_skein512_close(&ctx_skein, (void*) hash);
+
+    if (hash[0] & 0x8)
+    {
+        sph_keccak512_init(&ctx_keccak);
+        // ZKECCAK;
+        sph_keccak512 (&ctx_keccak, (const void*) hash, 64);
+        sph_keccak512_close(&ctx_keccak, (void*) hash);
+    }
+    else
+    {
+        sph_jh512_init(&ctx_jh);
+        // ZJH;
+        sph_jh512 (&ctx_jh, (const void*) hash, 64);
+        sph_jh512_close(&ctx_jh, (void*) hash);
+    }
+
+    memcpy(state, hash, 32);
+}
+
+
+extern bool opt_benchmark;
+
+extern "C" int scanhash_quark(int thr_id, uint32_t *pdata,
+    const uint32_t *ptarget, uint32_t max_nonce,
+    unsigned long *hashes_done)
+{
+	const uint32_t first_nonce = pdata[19];
+
+	// TODO: entfernen für eine Release! Ist nur zum Testen!
+	if (opt_benchmark)
+		((uint32_t*)ptarget)[7] = 0x0000ff;
+
+	const uint32_t Htarg = ptarget[7];
+
+	const int throughput = 256*4096; // 100;
+
+	static bool init[8] = {0,0,0,0,0,0,0,0};
+	if (!init[thr_id])
+	{
+		cudaSetDevice(device_map[thr_id]);
+
+		// Konstanten kopieren, Speicher belegen
+		cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput);
+		quark_blake512_cpu_init(thr_id, throughput);
+		quark_groestl512_cpu_init(thr_id, throughput);
+		quark_skein512_cpu_init(thr_id, throughput);
+		quark_bmw512_cpu_init(thr_id, throughput);
+		quark_keccak512_cpu_init(thr_id, throughput);
+		quark_jh512_cpu_init(thr_id, throughput);
+		quark_check_cpu_init(thr_id, throughput);
+		quark_compactTest_cpu_init(thr_id, throughput);
+		cudaMalloc(&d_quarkNonces[thr_id], sizeof(uint32_t)*throughput);
+		cudaMalloc(&d_branch1Nonces[thr_id], sizeof(uint32_t)*throughput);
+		cudaMalloc(&d_branch2Nonces[thr_id], sizeof(uint32_t)*throughput);
+		cudaMalloc(&d_branch3Nonces[thr_id], sizeof(uint32_t)*throughput);
+		init[thr_id] = true;
+	}
+
+	uint32_t endiandata[20];
+	for (int k=0; k < 20; k++)
+		be32enc(&endiandata[k], ((uint32_t*)pdata)[k]);
+
+	quark_blake512_cpu_setBlock_80((void*)endiandata);
+	quark_check_cpu_setTarget(ptarget);
+
+	do {
+		int order = 0;
+		size_t nrm1=0, nrm2=0, nrm3=0;
+
+		// erstes Blake512 Hash mit CUDA
+		quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für BMW512
+		quark_bmw512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+
+		quark_compactTest_single_false_cpu_hash_64(thr_id, throughput, pdata[19], d_hash[thr_id], NULL,
+				d_branch3Nonces[thr_id], &nrm3,
+				order++);
+		
+		// nur den Skein Branch weiterverfolgen
+		quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für Groestl512
+		quark_groestl512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für JH512
+		quark_jh512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// quarkNonces in branch1 und branch2 aufsplitten gemäss if (hash[0] & 0x8)
+		quark_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id],
+			d_branch1Nonces[thr_id], &nrm1,
+			d_branch2Nonces[thr_id], &nrm2,
+			order++);
+
+		// das ist der bedingte Branch für Blake512
+		quark_blake512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der bedingte Branch für Bmw512
+		quark_bmw512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für Keccak512
+		quark_keccak512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der unbedingte Branch für Skein512
+		quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+
+		// quarkNonces in branch1 und branch2 aufsplitten gemäss if (hash[0] & 0x8)
+		quark_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id],
+			d_branch1Nonces[thr_id], &nrm1,
+			d_branch2Nonces[thr_id], &nrm2,
+			order++);
+
+		// das ist der bedingte Branch für Keccak512
+		quark_keccak512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id], order++);
+
+		// das ist der bedingte Branch für JH512
+		quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++);
+
+		// Scan nach Gewinner Hashes auf der GPU
+		uint32_t foundNonce = quark_check_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+		if  (foundNonce != 0xffffffff)
+		{
+			uint32_t vhash64[8];
+			be32enc(&endiandata[19], foundNonce);
+			quarkhash(vhash64, endiandata);
+
+			if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) {
+
+				pdata[19] = foundNonce;
+				*hashes_done = (foundNonce - first_nonce + 1)/2;
+				return 1;
+			} else {
+				applog(LOG_INFO, "GPU #%d: result for nonce $%08X does not validate on CPU!", thr_id, foundNonce);
+			}
+		}
+
+		pdata[19] += throughput;
+
+	} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
+
+	*hashes_done = (pdata[19] - first_nonce + 1)/2;
+	return 0;
+}