Merge branch 'keccak' into bfgminer

.travis.yml

@@ -10,59 +10,63 @@ matrix:
   include:
     - compiler: ": Full GCC"
       # Upgrade GCC to avoid false warnings; build the full project with -Werror
-      env: myCC='gcc' UBUNTU_DEPS='gcc libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-system-libbase58 --enable-tool'
+      env: myCC='gcc' UBUNTU_DEPS='gcc libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-system-libbase58 --enable-tool'
     - compiler: ": Full LLVM"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-system-libbase58 --enable-tool'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-system-libbase58 --enable-tool'
     - compiler: ": pkgconf"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev pkgconf' EXTRA_DEPS='libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-system-libbase58 --enable-tool'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev pkgconf' EXTRA_DEPS='libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-system-libbase58 --enable-tool'
     - compiler: ": MinGW64 ncurses"
-      env: UBUNTU_DEPS='gcc-mingw-w64-x86-64' EXTRA_DEPS='pkg-config yasm' CROSS_BINPKGS='x86_64-w64-mingw32+ncurses' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-system-libbase58 --host=x86_64-w64-mingw32 --disable-knc --disable-bfsb --disable-jingtian --disable-metabank --disable-minergate --disable-titan --disable-kncasic'
+      env: UBUNTU_DEPS='gcc-mingw-w64-x86-64' EXTRA_DEPS='pkg-config yasm' CROSS_BINPKGS='x86_64-w64-mingw32+ncurses' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-system-libbase58 --host=x86_64-w64-mingw32 --disable-knc --disable-bfsb --disable-jingtian --disable-metabank --disable-minergate --disable-titan --disable-kncasic'
     - compiler: ": MinGW64 pdcurses"
-      env: UBUNTU_DEPS='gcc-mingw-w64-x86-64' EXTRA_DEPS='pkg-config yasm' CROSS_BINPKGS='x86_64-w64-mingw32' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-system-libbase58 --host=x86_64-w64-mingw32 --disable-knc --disable-bfsb --disable-jingtian --disable-metabank --disable-minergate --disable-titan --disable-kncasic'
+      env: UBUNTU_DEPS='gcc-mingw-w64-x86-64' EXTRA_DEPS='pkg-config yasm' CROSS_BINPKGS='x86_64-w64-mingw32' BUILD_CFLAGS='-Werror' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-system-libbase58 --host=x86_64-w64-mingw32 --disable-knc --disable-bfsb --disable-jingtian --disable-metabank --disable-minergate --disable-titan --disable-kncasic'
+    - compiler: ": Standard"
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-keccak --enable-scrypt'
     - compiler: ": Std SHA2"
       env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS=''
+    - compiler: ": Std Keccak"
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-keccak'
     - compiler: ": Std scrypt"
       env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-scrypt'
     - compiler: ": No hidapi"
-      env: myCC='clang' UBUNTU_DEPS='linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt'
+      env: myCC='clang' UBUNTU_DEPS='linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt'
     - compiler: ": No VFIO"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-vfio'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-vfio'
     - compiler: ": No UIO"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-uio'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-uio'
     - compiler: ": No VFIO/UIO"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-uio --without-vfio'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-uio --without-vfio'
     - compiler: ": Non-wide ncurses"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncurses5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --with-curses=ncurses5'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncurses5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --with-curses=ncurses5'
     - compiler: ": No curses"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-curses'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-curses'
     - compiler: ": No libudev"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-libudev'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-libudev'
     - compiler: ": No libusb"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-libusb'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-libusb'
     - compiler: ": No libevent"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-libevent'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libmicrohttpd-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-libevent'
     - compiler: ": No libmicrohttpd"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-libmicrohttpd'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libi2c-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-libmicrohttpd'
     - compiler: ": No libi2c-dev"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --disable-knc --disable-titan --disable-kncasic'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev yasm libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --disable-knc --disable-titan --disable-kncasic'
     - compiler: ": No yasm"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev libsensors4-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt'
     - compiler: ": No libsensors"
-      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-sensors'
+      env: myCC='clang' UBUNTU_DEPS='libhidapi-dev linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev libudev-dev libusb-1.0-0-dev libevent-dev libmicrohttpd-dev libi2c-dev yasm' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-sensors'
     - compiler: ": No opt deps"
-      env: myCC='clang' EXTRA_DEPS='pkg-config' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-uio --without-vfio --without-sensors --without-libmicrohttpd --without-libevent --without-libusb --without-curses --without-libudev --disable-knc --disable-titan --disable-kncasic'
+      env: myCC='clang' EXTRA_DEPS='pkg-config' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-uio --without-vfio --without-sensors --without-libmicrohttpd --without-libevent --without-libusb --without-curses --without-libudev --disable-knc --disable-titan --disable-kncasic'
     - compiler: ": Only ncurses"
-      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-scrypt --without-uio --without-vfio --without-sensors --without-libmicrohttpd --without-libevent --without-libusb --without-libudev --disable-knc --disable-titan --disable-kncasic'
+      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev' CONFIGURE_ARGS='--enable-other-drivers --enable-keccak --enable-scrypt --without-uio --without-vfio --without-sensors --without-libmicrohttpd --without-libevent --without-libusb --without-libudev --disable-knc --disable-titan --disable-kncasic'
     - compiler: ": Only CPU"
-      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev yasm' CONFIGURE_ARGS='--disable-other-drivers --enable-cpumining --enable-scrypt'
+      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev yasm' CONFIGURE_ARGS='--disable-other-drivers --enable-cpumining --enable-keccak --enable-scrypt'
     - compiler: ": Only OpenCL"
-      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev libsensors4-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-scrypt'
+      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev libsensors4-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-keccak --enable-scrypt'
     - compiler: ": OpenCL w/o ADL"
-      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev libsensors4-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-scrypt --disable-adl'
+      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev libsensors4-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-keccak --enable-scrypt --disable-adl'
     - compiler: ": OpenCL w/o sensors"
-      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-scrypt'
+      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-keccak --enable-scrypt'
     - compiler: ": OpenCL w/o ADL or sensors"
-      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-scrypt --disable-adl'
+      env: myCC='clang' EXTRA_DEPS='pkg-config libncursesw5-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-opencl --enable-keccak --enable-scrypt --disable-adl'
     - compiler: ": Only bitforce"
       env: myCC='clang' UBUNTU_DEPS='linux-libc-dev' EXTRA_DEPS='pkg-config libncursesw5-dev' CONFIGURE_ARGS='--disable-other-drivers --enable-bitforce'
     - compiler: ": Only icarus"

AUTHORS

@@ -54,6 +54,7 @@ Jean-Luc Cooke <jlcooke@certainkey.com>
 Jonathan Lynch <jonathan.lynch@intel.com>
 Josh Lehan <krellan@krellan.net>
 Lingchao Xu <lingchao.xu@bitmaintech.com>
+Luke Mitchell <Luke.Mitchell.2011@my.bristol.ac.uk>
 Mark Crichton <crichton@gmail.com>
 Martin Danielsen <kalroth@gmail.com>
 Michael Kedzierski <michael@kedzierski.id.au>

Makefile.am

@@ -148,6 +148,14 @@ endif
 endif
 
 
+if USE_KECCAK
+bfgminer_SOURCES += malgo/keccak.c
+
+if USE_OPENCL
+dist_kernels_DATA += $(top_srcdir)/opencl/keccak.cl
+endif
+endif
+
 if USE_SHA256D
 bfgminer_SOURCES += malgo/sha256d.c
 

README

@@ -182,6 +182,7 @@ BFGMiner driver configuration options:
 	--disable-ztex          Compile support for ZTEX (default enabled)
 
 BFGMiner algorithm configuration option:
+	--enable-keccak         Compile support for Keccak (default disabled)
 	--disable-sha256d       Compile support for SHA256d (default enabled)
 	--enable-scrypt         Compile support for scrypt (default disabled)
 

configure.ac

@@ -361,6 +361,7 @@ AM_CONDITIONAL([USE_UDEVRULES_GROUP], [$use_udevrules_group])
 AC_SUBST([UDEVRULES_GROUP], [$udevrules_group])
 
 
+BFG_ALGO(Keccak,no)
 BFG_ALGO(SHA256d,yes)
 BFG_ALGO(scrypt,no)
 
@@ -474,6 +475,10 @@ AC_SUBST(JANSSON_LIBS)
 if test "x$USE_OPENCL" = xyes; then
 	adl="yes"
 	
+	if test "x$USE_KECCAK" = xyes; then
+		AC_DEFINE([USE_OPENCL_FULLHEADER],[1],[Defined to 1 if OpenCL fullheader kernel interface is wanted])
+	fi
+	
 driverlist="$driverlist opencl:sensors/with_sensors"
 AC_ARG_WITH([sensors],
 	[AC_HELP_STRING([--without-sensors],[Build with libsensors monitoring (default enabled)])],

debian/rules

@@ -5,7 +5,7 @@
 
 override_dh_auto_configure:
 	NOSUBMODULES=1 ./autogen.sh
-	dh_auto_configure -- --enable-ztex --enable-bitforce --enable-icarus --enable-cpumining --enable-scrypt --enable-opencl
+	dh_auto_configure -- --enable-ztex --enable-bitforce --enable-icarus --enable-cpumining --enable-keccak --enable-scrypt --enable-opencl
 
 override_dh_auto_install:
 	$(MAKE) DESTDIR=$(CURDIR)/debian/bfgminer install

driver-opencl.c

@@ -1244,6 +1244,34 @@ cl_int queue_scrypt_kernel(const struct opencl_kernel_info * const kinfo, _clSta
 }
 #endif
 
+#ifdef USE_OPENCL_FULLHEADER
+static
+cl_int queue_fullheader_kernel(const struct opencl_kernel_info * const kinfo, _clState * const clState, struct work * const work, __maybe_unused const cl_uint threads)
+{
+	const struct mining_algorithm * const malgo = work_mining_algorithm(work);
+	const cl_kernel * const kernel = &kinfo->kernel;
+	unsigned int num = 0;
+	cl_int status = 0;
+	uint8_t blkheader[80];
+	
+	work->nonce_diff = malgo->opencl_min_nonce_diff;
+	
+	if (!kinfo->goffset)
+	{
+		cl_uint nonce_base = work->blk.nonce;
+		CL_SET_ARG(nonce_base);
+	}
+	
+	swap32yes(blkheader, work->data, 80/4);
+	status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, CL_TRUE, 0, sizeof(blkheader), blkheader, 0, NULL, NULL);
+	
+	CL_SET_ARG(clState->CLbuffer0);
+	CL_SET_ARG(clState->outputBuffer);
+	
+	return status;
+}
+#endif
+
 
 static
 struct opencl_kernel_interface kernel_interfaces[] = {
@@ -1254,6 +1282,9 @@ struct opencl_kernel_interface kernel_interfaces[] = {
 	{"diakgcn", queue_diakgcn_kernel},
 	{"diablo",  queue_diablo_kernel },
 #endif
+#ifdef USE_OPENCL_FULLHEADER
+	{"fullheader", queue_fullheader_kernel },
+#endif
 #ifdef USE_SCRYPT
 	{"scrypt",  queue_scrypt_kernel },
 #endif

findnonce.c

@@ -168,6 +168,10 @@ static void *postcalc_hash(void *userdata)
 
 	for (entry = 0; entry < pcd->res[found]; entry++) {
 		uint32_t nonce = pcd->res[entry];
+#ifdef USE_OPENCL_FULLHEADER
+		if (pcd->kinterface == KL_FULLHEADER)
+			nonce = swab32(nonce);
+#endif
 
 		applog(LOG_DEBUG, "OCL NONCE %u found in slot %d", nonce, entry);
 		submit_nonce(thr, &pcd->work, nonce);

make-release

@@ -97,6 +97,7 @@ for build in "${builds[@]}"; do
 		--enable-icarus \
 		--enable-modminer \
 		--enable-ztex \
+		--enable-keccak \
 		--enable-scrypt \
 		--without-system-libbase58 \
 		--host="$machine"

malgo/keccak.c

@@ -0,0 +1,374 @@
+/*
+ * Copyright 2013-2014 Ronny Van Keer (released as CC0)
+ * Copyright 2014 Luke Mitchell
+ * Copyright 2014 Luke Dashjr
+ *
+ * This program is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License as published by the Free
+ * Software Foundation; either version 3 of the License, or (at your option)
+ * any later version.  See COPYING for more details.
+ */
+
+#include "config.h"
+#include "miner.h"
+
+#include <math.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+
+#include <stdio.h>
+
+#include <uthash.h>
+
+struct uint256 {
+	unsigned char v[32];
+};
+typedef struct uint256 uint256;
+
+typedef unsigned long long UINT64;
+
+#define ROL(a, offset) ((a << offset) | (a >> (64-offset)))
+
+static const UINT64 KeccakF_RoundConstants[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
+};
+
+struct bin32 {
+	UINT64 v0;
+	UINT64 v1;
+	UINT64 v2;
+	UINT64 v3;
+};
+
+static
+void keccak1(unsigned char *out, const unsigned char *inraw, unsigned inrawlen)
+{
+	unsigned char temp[136];
+	unsigned round;
+	
+	UINT64 Aba, Abe, Abi, Abo, Abu;
+	UINT64 Aga, Age, Agi, Ago, Agu;
+	UINT64 Aka, Ake, Aki, Ako, Aku;
+	UINT64 Ama, Ame, Ami, Amo, Amu;
+	UINT64 Asa, Ase, Asi, Aso, Asu;
+	UINT64 BCa, BCe, BCi, BCo, BCu;
+	UINT64 Da, De, Di, Do, Du;
+	UINT64 Eba, Ebe, Ebi, Ebo, Ebu;
+	UINT64 Ega, Ege, Egi, Ego, Egu;
+	UINT64 Eka, Eke, Eki, Eko, Eku;
+	UINT64 Ema, Eme, Emi, Emo, Emu;
+	UINT64 Esa, Ese, Esi, Eso, Esu;
+	
+	memcpy(temp, inraw, inrawlen);
+	temp[inrawlen++] = 1;
+	memset( temp+inrawlen, 0, 136 - inrawlen);
+	temp[136-1] |= 0x80;
+	const UINT64 *in = (const UINT64 *)temp;
+	
+	// copyFromState(A, state)
+	Aba = in[ 0];
+	Abe = in[ 1];
+	Abi = in[ 2];
+	Abo = in[ 3];
+	Abu = in[ 4];
+	Aga = in[ 5];
+	Age = in[ 6];
+	Agi = in[ 7];
+	Ago = in[ 8];
+	Agu = in[ 9];
+	Aka = in[10];
+	Ake = in[11];
+	Aki = in[12];
+	Ako = in[13];
+	Aku = in[14];
+	Ama = in[15];
+	Ame = in[16];
+	Ami = 0;
+	Amo = 0;
+	Amu = 0;
+	Asa = 0;
+	Ase = 0;
+	Asi = 0;
+	Aso = 0;
+	Asu = 0;
+	
+	for (round = 0; round < 24; round += 2)
+	{
+		// prepareTheta
+		BCa = Aba^Aga^Aka^Ama^Asa;
+		BCe = Abe^Age^Ake^Ame^Ase;
+		BCi = Abi^Agi^Aki^Ami^Asi;
+		BCo = Abo^Ago^Ako^Amo^Aso;
+		BCu = Abu^Agu^Aku^Amu^Asu;
+		
+		// thetaRhoPiChiIotaPrepareTheta(round, A, E)
+		Da = BCu^ROL(BCe, 1);
+		De = BCa^ROL(BCi, 1);
+		Di = BCe^ROL(BCo, 1);
+		Do = BCi^ROL(BCu, 1);
+		Du = BCo^ROL(BCa, 1);
+		
+		Aba ^= Da;
+		BCa = Aba;
+		Age ^= De;
+		BCe = ROL(Age, 44);
+		Aki ^= Di;
+		BCi = ROL(Aki, 43);
+		Amo ^= Do;
+		BCo = ROL(Amo, 21);
+		Asu ^= Du;
+		BCu = ROL(Asu, 14);
+		Eba = BCa ^((~BCe) & BCi);
+		Eba ^= KeccakF_RoundConstants[round];
+		Ebe = BCe ^((~BCi) & BCo);
+		Ebi = BCi ^((~BCo) & BCu);
+		Ebo = BCo ^((~BCu) & BCa);
+		Ebu = BCu ^((~BCa) & BCe);
+		
+		Abo ^= Do;
+		BCa = ROL(Abo, 28);
+		Agu ^= Du;
+		BCe = ROL(Agu, 20);
+		Aka ^= Da;
+		BCi = ROL(Aka,  3);
+		Ame ^= De;
+		BCo = ROL(Ame, 45);
+		Asi ^= Di;
+		BCu = ROL(Asi, 61);
+		Ega = BCa ^((~BCe) & BCi);
+		Ege = BCe ^((~BCi) & BCo);
+		Egi = BCi ^((~BCo) & BCu);
+		Ego = BCo ^((~BCu) & BCa);
+		Egu = BCu ^((~BCa) & BCe);
+		
+		Abe ^= De;
+		BCa = ROL(Abe,  1);
+		Agi ^= Di;
+		BCe = ROL(Agi,  6);
+		Ako ^= Do;
+		BCi = ROL(Ako, 25);
+		Amu ^= Du;
+		BCo = ROL(Amu,  8);
+		Asa ^= Da;
+		BCu = ROL(Asa, 18);
+		Eka = BCa ^((~BCe) & BCi);
+		Eke = BCe ^((~BCi) & BCo);
+		Eki = BCi ^((~BCo) & BCu);
+		Eko = BCo ^((~BCu) & BCa);
+		Eku = BCu ^((~BCa) & BCe);
+		
+		Abu ^= Du;
+		BCa = ROL(Abu, 27);
+		Aga ^= Da;
+		BCe = ROL(Aga, 36);
+		Ake ^= De;
+		BCi = ROL(Ake, 10);
+		Ami ^= Di;
+		BCo = ROL(Ami, 15);
+		Aso ^= Do;
+		BCu = ROL(Aso, 56);
+		Ema = BCa ^((~BCe) & BCi);
+		Eme = BCe ^((~BCi) & BCo);
+		Emi = BCi ^((~BCo) & BCu);
+		Emo = BCo ^((~BCu) & BCa);
+		Emu = BCu ^((~BCa) & BCe);
+		
+		Abi ^= Di;
+		BCa = ROL(Abi, 62);
+		Ago ^= Do;
+		BCe = ROL(Ago, 55);
+		Aku ^= Du;
+		BCi = ROL(Aku, 39);
+		Ama ^= Da;
+		BCo = ROL(Ama, 41);
+		Ase ^= De;
+		BCu = ROL(Ase,  2);
+		Esa = BCa ^((~BCe) & BCi);
+		Ese = BCe ^((~BCi) & BCo);
+		Esi = BCi ^((~BCo) & BCu);
+		Eso = BCo ^((~BCu) & BCa);
+		Esu = BCu ^((~BCa) & BCe);
+		
+		// prepareTheta
+		BCa = Eba^Ega^Eka^Ema^Esa;
+		BCe = Ebe^Ege^Eke^Eme^Ese;
+		BCi = Ebi^Egi^Eki^Emi^Esi;
+		BCo = Ebo^Ego^Eko^Emo^Eso;
+		BCu = Ebu^Egu^Eku^Emu^Esu;
+		
+		// thetaRhoPiChiIotaPrepareTheta(round+1, E, A)
+		Da = BCu^ROL(BCe, 1);
+		De = BCa^ROL(BCi, 1);
+		Di = BCe^ROL(BCo, 1);
+		Do = BCi^ROL(BCu, 1);
+		Du = BCo^ROL(BCa, 1);
+		
+		Eba ^= Da;
+		BCa = Eba;
+		Ege ^= De;
+		BCe = ROL(Ege, 44);
+		Eki ^= Di;
+		BCi = ROL(Eki, 43);
+		Emo ^= Do;
+		BCo = ROL(Emo, 21);
+		Esu ^= Du;
+		BCu = ROL(Esu, 14);
+		Aba = BCa ^((~BCe) & BCi);
+		Aba ^= KeccakF_RoundConstants[round+1];
+		Abe = BCe ^((~BCi) & BCo);
+		Abi = BCi ^((~BCo) & BCu);
+		Abo = BCo ^((~BCu) & BCa);
+		Abu = BCu ^((~BCa) & BCe);
+		
+		Ebo ^= Do;
+		BCa = ROL(Ebo, 28);
+		Egu ^= Du;
+		BCe = ROL(Egu, 20);
+		Eka ^= Da;
+		BCi = ROL(Eka, 3);
+		Eme ^= De;
+		BCo = ROL(Eme, 45);
+		Esi ^= Di;
+		BCu = ROL(Esi, 61);
+		Aga = BCa ^((~BCe) & BCi);
+		Age = BCe ^((~BCi) & BCo);
+		Agi = BCi ^((~BCo) & BCu);
+		Ago = BCo ^((~BCu) & BCa);
+		Agu = BCu ^((~BCa) & BCe);
+		
+		Ebe ^= De;
+		BCa = ROL(Ebe, 1);
+		Egi ^= Di;
+		BCe = ROL(Egi, 6);
+		Eko ^= Do;
+		BCi = ROL(Eko, 25);
+		Emu ^= Du;
+		BCo = ROL(Emu, 8);
+		Esa ^= Da;
+		BCu = ROL(Esa, 18);
+		Aka = BCa ^((~BCe) & BCi);
+		Ake = BCe ^((~BCi) & BCo);
+		Aki = BCi ^((~BCo) & BCu);
+		Ako = BCo ^((~BCu) & BCa);
+		Aku = BCu ^((~BCa) & BCe);
+		
+		Ebu ^= Du;
+		BCa = ROL(Ebu, 27);
+		Ega ^= Da;
+		BCe = ROL(Ega, 36);
+		Eke ^= De;
+		BCi = ROL(Eke, 10);
+		Emi ^= Di;
+		BCo = ROL(Emi, 15);
+		Eso ^= Do;
+		BCu = ROL(Eso, 56);
+		Ama = BCa ^((~BCe) & BCi);
+		Ame = BCe ^((~BCi) & BCo);
+		Ami = BCi ^((~BCo) & BCu);
+		Amo = BCo ^((~BCu) & BCa);
+		Amu = BCu ^((~BCa) & BCe);
+		
+		Ebi ^= Di;
+		BCa = ROL(Ebi, 62);
+		Ego ^= Do;
+		BCe = ROL(Ego, 55);
+		Eku ^= Du;
+		BCi = ROL(Eku, 39);
+		Ema ^= Da;
+		BCo = ROL(Ema, 41);
+		Ese ^= De;
+		BCu = ROL(Ese, 2);
+		Asa = BCa ^((~BCe) & BCi);
+		Ase = BCe ^((~BCi) & BCo);
+		Asi = BCi ^((~BCo) & BCu);
+		Aso = BCo ^((~BCu) & BCa);
+		Asu = BCu ^((~BCa) & BCe);
+	}
+	{
+		UINT64 *out64 = (UINT64 *)out;
+		out64[ 0] = Aba;
+		out64[ 1] = Abe;
+		out64[ 2] = Abi;
+		out64[ 3] = Abo;
+	}
+}
+
+static
+void keccak_hash_data(void * const digest, const void * const pdata)
+{
+	uint32_t data[20];
+	swap32yes(data, pdata, 20);
+	keccak1(digest, (unsigned char*)data, 80);
+}
+
+#ifdef USE_OPENCL
+static
+float opencl_oclthreads_to_intensity_keccak(const unsigned long oclthreads)
+{
+	return log2f(oclthreads) - 13.;
+}
+
+static
+unsigned long opencl_intensity_to_oclthreads_keccak(float intensity)
+{
+	return powf(2, intensity + 13);
+}
+
+static
+char *opencl_get_default_kernel_file_keccak(const struct mining_algorithm * const malgo, struct cgpu_info * const cgpu, struct _clState * const clState)
+{
+	return strdup("keccak");
+}
+#endif
+
+static struct mining_algorithm malgo_keccak = {
+	.name = "Keccak",
+	.aliases = "Keccak",
+	
+	.algo = POW_KECCAK,
+	.ui_skip_hash_bytes = 4,
+	.worktime_skip_prevblk_u32 = 1,
+	.reasonable_low_nonce_diff = 1.,
+	
+	.hash_data_f = keccak_hash_data,
+	
+#ifdef USE_OPENCL
+	.opencl_oclthreads_to_intensity = opencl_oclthreads_to_intensity_keccak,
+	.opencl_intensity_to_oclthreads = opencl_intensity_to_oclthreads_keccak,
+	.opencl_min_oclthreads =       0x20,  // intensity -8
+	.opencl_max_oclthreads = 0x20000000,  // intensity 16
+	.opencl_min_nonce_diff = 1./0x10,
+	.opencl_get_default_kernel_file = opencl_get_default_kernel_file_keccak,
+#endif
+};
+
+static
+__attribute__((constructor))
+void init_keccak(void)
+{
+    LL_APPEND(mining_algorithms, (&malgo_keccak));
+}

malgo/sha256d.c

@@ -32,11 +32,13 @@ void hash_data(void *out_hash, const void *data)
 }
 
 #ifdef USE_OPENCL
+static
 float opencl_oclthreads_to_intensity_sha256d(const unsigned long oclthreads)
 {
 	return log2f(oclthreads) - 15.;
 }
 
+static
 unsigned long opencl_intensity_to_oclthreads_sha256d(float intensity)
 {
 	return powf(2, intensity + 15);

miner.h

@@ -277,6 +277,9 @@ struct gpu_adl {
 #endif
 
 enum pow_algorithm {
+#ifdef USE_KECCAK
+	POW_KECCAK,
+#endif
 #ifdef USE_SHA256D
 	POW_SHA256D,
 #endif
@@ -381,6 +384,9 @@ enum cl_kernels {
 	KL_DIAKGCN,
 	KL_DIABLO,
 #endif
+#ifdef USE_OPENCL_FULLHEADER
+	KL_FULLHEADER,
+#endif
 #ifdef USE_SCRYPT
 	KL_SCRYPT,
 #endif

ocl.c

@@ -785,6 +785,11 @@ bool opencl_load_kernel(struct cgpu_info * const cgpu, _clState * const clState,
 			case KL_PHATK:
 				kernel_goffset_support = 0;
 				break;
+#endif
+#ifdef USE_OPENCL_FULLHEADER
+			case KL_FULLHEADER:
+				kernel_goffset_support = 1;
+				break;
 #endif
 			case KL_NONE: case OPENCL_KERNEL_INTERFACE_COUNT:
 #ifdef USE_SCRYPT
@@ -1156,36 +1161,51 @@ built:
 	free((void*)cgpu->kname);
 	cgpu->kname = strdup(kernel_file);
 
-#ifdef USE_SCRYPT
-	if (kernelinfo->interface == KL_SCRYPT && !clState->padbufsize)
+#ifdef MAX_CLBUFFER0_SZ
+	switch (kernelinfo->interface)
 	{
-		size_t ipt = (1024 / data->lookup_gap + (1024 % data->lookup_gap > 0));
-		size_t bufsize = 128 * ipt * data->thread_concurrency;
-
-		/* Use the max alloc value which has been rounded to a power of
-		 * 2 greater >= required amount earlier */
-		if (bufsize > data->max_alloc) {
-			applog(LOG_WARNING, "Maximum buffer memory device %d supports says %lu", gpu, (unsigned long)data->max_alloc);
-			applog(LOG_WARNING, "Your scrypt settings come to %lu", (unsigned long)bufsize);
-		}
-		applog(LOG_DEBUG, "Creating scrypt buffer sized %lu", (unsigned long)bufsize);
-		clState->padbufsize = bufsize;
-
-		/* This buffer is weird and might work to some degree even if
-		 * the create buffer call has apparently failed, so check if we
-		 * get anything back before we call it a failure. */
-		clState->padbuffer8 = NULL;
-		clState->padbuffer8 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, bufsize, NULL, &status);
-		if (status != CL_SUCCESS && !clState->padbuffer8) {
-			applog(LOG_ERR, "Error %d: clCreateBuffer (padbuffer8), decrease TC or increase LG", status);
-			return false;
-		}
-
-		clState->CLbuffer0 = clCreateBuffer(clState->context, CL_MEM_READ_ONLY, 128, NULL, &status);
-		if (status != CL_SUCCESS) {
-			applog(LOG_ERR, "Error %d: clCreateBuffer (CLbuffer0)", status);
-			return false;
-		}
+#ifdef USE_SCRYPT
+		case KL_SCRYPT:
+			if (!clState->padbufsize)
+			{
+				size_t ipt = (1024 / data->lookup_gap + (1024 % data->lookup_gap > 0));
+				size_t bufsize = 128 * ipt * data->thread_concurrency;
+
+				/* Use the max alloc value which has been rounded to a power of
+				 * 2 greater >= required amount earlier */
+				if (bufsize > data->max_alloc) {
+					applog(LOG_WARNING, "Maximum buffer memory device %d supports says %lu", gpu, (unsigned long)data->max_alloc);
+					applog(LOG_WARNING, "Your scrypt settings come to %lu", (unsigned long)bufsize);
+				}
+				applog(LOG_DEBUG, "Creating scrypt buffer sized %lu", (unsigned long)bufsize);
+				clState->padbufsize = bufsize;
+
+				/* This buffer is weird and might work to some degree even if
+				 * the create buffer call has apparently failed, so check if we
+				 * get anything back before we call it a failure. */
+				clState->padbuffer8 = NULL;
+				clState->padbuffer8 = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, bufsize, NULL, &status);
+				if (status != CL_SUCCESS && !clState->padbuffer8) {
+					applog(LOG_ERR, "Error %d: clCreateBuffer (padbuffer8), decrease TC or increase LG", status);
+					return false;
+				}
+			}
+			// NOTE: fallthru
+#endif
+#ifdef USE_OPENCL_FULLHEADER
+		case KL_FULLHEADER:
+#endif
+			if (!clState->CLbuffer0)
+			{
+				clState->CLbuffer0 = clCreateBuffer(clState->context, CL_MEM_READ_ONLY, MAX_CLBUFFER0_SZ, NULL, &status);
+				if (status != CL_SUCCESS) {
+					applog(LOG_ERR, "Error %d: clCreateBuffer (CLbuffer0)", status);
+					return false;
+				}
+			}
+			break;
+		default:
+			break;
 	}
 #endif
 

ocl.h

@@ -8,6 +8,14 @@
 
 #include "miner.h"
 
+#define SCRYPT_CLBUFFER0_SZ      (128)
+#define FULLHEADER_CLBUFFER0_SZ  ( 80)
+#ifdef USE_SCRYPT
+#	define MAX_CLBUFFER0_SZ  SCRYPT_CLBUFFER0_SZ
+#elif USE_OPENCL_FULLHEADER
+#	define MAX_CLBUFFER0_SZ  FULLHEADER_CLBUFFER0_SZ
+#endif
+
 struct mining_algorithm;
 struct opencl_kernel_info;
 typedef struct _clState _clState;
@@ -21,8 +29,10 @@ struct _clState {
 	cl_command_queue commandQueue;
 	
 	cl_mem outputBuffer;
-#ifdef USE_SCRYPT
+#ifdef MAX_CLBUFFER0_SZ
 	cl_mem CLbuffer0;
+#endif
+#ifdef USE_SCRYPT
 	cl_mem padbuffer8;
 	size_t padbufsize;
 	void * cldata;

opencl/keccak.cl

@@ -0,0 +1,133 @@
+/*
+ * Scrypt-jane public domain, OpenCL implementation of scrypt(keccak,chacha,SCRYPTN,1,1) 2013 mtrlt
+ */
+
+// kernel-interface: fullheader Keccak
+
+#define ARGS_25(x) x ## 0, x ## 1, x ## 2, x ## 3, x ## 4, x ## 5, x ## 6, x ## 7, x ## 8, x ## 9, x ## 10, x ## 11, x ## 12, x ## 13, x ## 14, x ## 15, x ## 16, x ## 17, x ## 18, x ## 19, x ## 20, x ## 21, x ## 22, x ## 23, x ## 24
+
+__constant uint2 keccak_round_constants[24] = {
+	(uint2)(0x00000001,0x00000000), (uint2)(0x00008082,0x00000000),
+	(uint2)(0x0000808a,0x80000000), (uint2)(0x80008000,0x80000000),
+	(uint2)(0x0000808b,0x00000000), (uint2)(0x80000001,0x00000000),
+	(uint2)(0x80008081,0x80000000), (uint2)(0x00008009,0x80000000),
+	(uint2)(0x0000008a,0x00000000), (uint2)(0x00000088,0x00000000),
+	(uint2)(0x80008009,0x00000000), (uint2)(0x8000000a,0x00000000),
+	(uint2)(0x8000808b,0x00000000), (uint2)(0x0000008b,0x80000000),
+	(uint2)(0x00008089,0x80000000), (uint2)(0x00008003,0x80000000),
+	(uint2)(0x00008002,0x80000000), (uint2)(0x00000080,0x80000000),
+	(uint2)(0x0000800a,0x00000000), (uint2)(0x8000000a,0x80000000),
+	(uint2)(0x80008081,0x80000000), (uint2)(0x00008080,0x80000000),
+	(uint2)(0x80000001,0x00000000), (uint2)(0x80008008,0x80000000)
+};
+
+uint2 ROTL64_1(const uint2 x, const uint y)
+{
+	return (uint2)((x.x<<y)^(x.y>>(32-y)),(x.y<<y)^(x.x>>(32-y)));
+}
+uint2 ROTL64_2(const uint2 x, const uint y)
+{
+	return (uint2)((x.y<<y)^(x.x>>(32-y)),(x.x<<y)^(x.y>>(32-y)));
+}
+
+#define RND(i) \
+do{  \
+		m0 = *s0 ^ *s5 ^ *s10 ^ *s15 ^ *s20 ^ ROTL64_1(*s2 ^ *s7 ^ *s12 ^ *s17 ^ *s22, 1);\
+		m1 = *s1 ^ *s6 ^ *s11 ^ *s16 ^ *s21 ^ ROTL64_1(*s3 ^ *s8 ^ *s13 ^ *s18 ^ *s23, 1);\
+		m2 = *s2 ^ *s7 ^ *s12 ^ *s17 ^ *s22 ^ ROTL64_1(*s4 ^ *s9 ^ *s14 ^ *s19 ^ *s24, 1);\
+		m3 = *s3 ^ *s8 ^ *s13 ^ *s18 ^ *s23 ^ ROTL64_1(*s0 ^ *s5 ^ *s10 ^ *s15 ^ *s20, 1);\
+		m4 = *s4 ^ *s9 ^ *s14 ^ *s19 ^ *s24 ^ ROTL64_1(*s1 ^ *s6 ^ *s11 ^ *s16 ^ *s21, 1);\
+\
+		m5 = *s1^m0;\
+\
+		*s0 ^= m4;\
+		*s1 = ROTL64_2(*s6^m0, 12);\
+		*s6 = ROTL64_1(*s9^m3, 20);\
+		*s9 = ROTL64_2(*s22^m1, 29);\
+		*s22 = ROTL64_2(*s14^m3, 7);\
+		*s14 = ROTL64_1(*s20^m4, 18);\
+		*s20 = ROTL64_2(*s2^m1, 30);\
+		*s2 = ROTL64_2(*s12^m1, 11);\
+		*s12 = ROTL64_1(*s13^m2, 25);\
+		*s13 = ROTL64_1(*s19^m3,  8);\
+		*s19 = ROTL64_2(*s23^m2, 24);\
+		*s23 = ROTL64_2(*s15^m4, 9);\
+		*s15 = ROTL64_1(*s4^m3, 27);\
+		*s4 = ROTL64_1(*s24^m3, 14);\
+		*s24 = ROTL64_1(*s21^m0,  2);\
+		*s21 = ROTL64_2(*s8^m2, 23);\
+		*s8 = ROTL64_2(*s16^m0, 13);\
+		*s16 = ROTL64_2(*s5^m4, 4);\
+		*s5 = ROTL64_1(*s3^m2, 28);\
+		*s3 = ROTL64_1(*s18^m2, 21);\
+		*s18 = ROTL64_1(*s17^m1, 15);\
+		*s17 = ROTL64_1(*s11^m0, 10);\
+		*s11 = ROTL64_1(*s7^m1,  6);\
+		*s7 = ROTL64_1(*s10^m4,  3);\
+		*s10 = ROTL64_1(      m5,  1);\
+		\
+		m5 = *s0; m6 = *s1; *s0 = bitselect(*s0^*s2,*s0,*s1); *s1 = bitselect(*s1^*s3,*s1,*s2); *s2 = bitselect(*s2^*s4,*s2,*s3); *s3 = bitselect(*s3^m5,*s3,*s4); *s4 = bitselect(*s4^m6,*s4,m5);\
+		m5 = *s5; m6 = *s6; *s5 = bitselect(*s5^*s7,*s5,*s6); *s6 = bitselect(*s6^*s8,*s6,*s7); *s7 = bitselect(*s7^*s9,*s7,*s8); *s8 = bitselect(*s8^m5,*s8,*s9); *s9 = bitselect(*s9^m6,*s9,m5);\
+		m5 = *s10; m6 = *s11; *s10 = bitselect(*s10^*s12,*s10,*s11); *s11 = bitselect(*s11^*s13,*s11,*s12); *s12 = bitselect(*s12^*s14,*s12,*s13); *s13 = bitselect(*s13^m5,*s13,*s14); *s14 = bitselect(*s14^m6,*s14,m5);\
+		m5 = *s15; m6 = *s16; *s15 = bitselect(*s15^*s17,*s15,*s16); *s16 = bitselect(*s16^*s18,*s16,*s17); *s17 = bitselect(*s17^*s19,*s17,*s18); *s18 = bitselect(*s18^m5,*s18,*s19); *s19 = bitselect(*s19^m6,*s19,m5);\
+		m5 = *s20; m6 = *s21; *s20 = bitselect(*s20^*s22,*s20,*s21); *s21 = bitselect(*s21^*s23,*s21,*s22); *s22 = bitselect(*s22^*s24,*s22,*s23); *s23 = bitselect(*s23^m5,*s23,*s24); *s24 = bitselect(*s24^m6,*s24,m5);\
+\
+		*s0 ^= keccak_round_constants[i];  \
+}while(0)
+
+void keccak_block_noabsorb(ARGS_25(uint2* s))
+{
+	uint2 m0,m1,m2,m3,m4,m5,m6;
+#pragma unroll
+	for (uint i = 0; i < 24; ++i)
+		RND(i);
+}
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search(
+#ifndef GOFFSET
+	const uint base,
+#endif
+	__global const uint2*restrict in, __global uint*restrict output)
+{
+#ifdef GOFFSET
+	const uint base = 0;
+#endif
+	uint2 ARGS_25(state);
+	
+	state0 = in[0];
+	state1 = in[1];
+	state2 = in[2];
+	state3 = in[3];
+	state4 = in[4];
+	state5 = in[5];
+	state6 = in[6];
+	state7 = in[7];
+	state8 = in[8];
+	state9 = (uint2)(in[9].x, base + get_global_id(0));
+	state10 = (uint2)(1,0);
+	state11 = 0;
+	state12 = 0;
+	state13 = 0;
+	state14 = 0;
+	state15 = 0;
+	state16 = (uint2)(0,0x80000000U);
+	state17 = 0;
+	state18 = 0;
+	state19 = 0;
+	state20 = 0;
+	state21 = 0;
+	state22 = 0;
+	state23 = 0;
+	state24 = 0;
+	
+	keccak_block_noabsorb(ARGS_25(&state));
+	
+#define FOUND (0x0F)
+#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+	
+	if ((state3.y & 0xFFFFFFF0U) == 0)
+	{
+		SETFOUND(base + get_global_id(0));
+	}
+}

openwrt/bfgminer/Makefile

@@ -62,6 +62,10 @@ config PACKAGE_$(PKG_NAME)_libusb
 	bool "Build with libusb support (X6500 & ZTEX)"
 	depends on PACKAGE_$(PKG_NAME)
 	default y
+config PACKAGE_$(PKG_NAME)_keccak
+	bool "Build with Keccak algorithm support"
+	depends on PACKAGE_$(PKG_NAME)
+	default n
 config PACKAGE_$(PKG_NAME)_scrypt
 	bool "Build with scrypt algorithm support"
 	depends on PACKAGE_$(PKG_NAME)
@@ -90,6 +94,10 @@ ifndef CONFIG_PACKAGE_$(PKG_NAME)_libusb
 CONFIGURE_ARGS += --without-libusb
 endif
 
+ifdef CONFIG_PACKAGE_$(PKG_NAME)_keccak
+CONFIGURE_ARGS += --enable-keccak
+endif
+
 ifdef CONFIG_PACKAGE_$(PKG_NAME)_scrypt
 CONFIGURE_ARGS += --enable-scrypt
 endif