opencl: Add Keccak support via "fullheader" kernel interface

Makefile.am

@@ -150,6 +150,10 @@ 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

configure.ac

@@ -475,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)])],

malgo/keccak.c

@@ -329,6 +329,17 @@ void keccak_hash_data(void * const digest, const void * const pdata)
 	memcpy(digest, &result, 0x20);
 }
 
+#ifdef USE_OPENCL
+extern float opencl_oclthreads_to_intensity_sha256d(unsigned long oclthreads);
+extern unsigned long opencl_intensity_to_oclthreads_sha256d(float intensity);
+
+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",
@@ -339,6 +350,15 @@ static struct mining_algorithm malgo_keccak = {
 	.reasonable_low_nonce_diff = 1.,
 	
 	.hash_data_f = keccak_hash_data,
+	
+#ifdef USE_OPENCL
+	.opencl_oclthreads_to_intensity = opencl_oclthreads_to_intensity_sha256d,
+	.opencl_intensity_to_oclthreads = opencl_intensity_to_oclthreads_sha256d,
+	.opencl_min_oclthreads =       0x20,  // intensity -10
+	.opencl_max_oclthreads = 0x20000000,  // intensity  14
+	.opencl_min_nonce_diff = 1./0x10,
+	.opencl_get_default_kernel_file = opencl_get_default_kernel_file_keccak,
+#endif
 };
 
 static

opencl/keccak.cl

@@ -1,3 +1,9 @@
+/*
+ * 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] = {
@@ -124,7 +130,3 @@ __kernel void search(__global const uint2*restrict in, __global uint*restrict ou
 		SETFOUND(get_global_id(0));
 	}
 }
-
-/*-
- * Scrypt-jane public domain, OpenCL implementation of scrypt(keccak,chacha,SCRYPTN,1,1) 2013 mtrlt
- */