Xn Algorithms Fine Tuning

Added options to fine tune X11-15 algorithms

algorithm.c

@@ -92,14 +92,26 @@ static void append_scrypt_compiler_options(struct _build_kernel_data *data, stru
   strcat(data->binary_filename, buf);
 }
 
-static void append_hamsi_compiler_options(struct _build_kernel_data *data, struct cgpu_info *cgpu, struct _algorithm_t *algorithm)
+static void append_x11_compiler_options(struct _build_kernel_data *data, struct cgpu_info *cgpu, struct _algorithm_t *algorithm)
 {
   char buf[255];
-  sprintf(buf, " -D SPH_HAMSI_EXPAND_BIG=%d -D SPH_HAMSI_SHORT=%d ",
-          opt_hamsi_expand_big, ((opt_hamsi_short)?1:0));
+  sprintf(buf, " -D SPH_COMPACT_BLAKE_64=%d -D SPH_LUFFA_PARALLEL=%d -D SPH_KECCAK_UNROLL=%u ",
+          ((opt_blake_compact)?1:0), ((opt_luffa_parallel)?1:0), (unsigned int)opt_keccak_unroll);
   strcat(data->compiler_options, buf);
 
-  sprintf(buf, "big%u%s", (unsigned int)opt_hamsi_expand_big, ((opt_hamsi_short)?"hs":""));
+  sprintf(buf, "ku%u%s%s", (unsigned int)opt_keccak_unroll, ((opt_blake_compact)?"bc":""), ((opt_luffa_parallel)?"lp":""));
+  strcat(data->binary_filename, buf);
+}
+
+
+static void append_x13_compiler_options(struct _build_kernel_data *data, struct cgpu_info *cgpu, struct _algorithm_t *algorithm)
+{
+  char buf[255];
+  sprintf(buf, " -D SPH_COMPACT_BLAKE_64=%d -D SPH_LUFFA_PARALLEL=%d -D SPH_KECCAK_UNROLL=%u -D SPH_HAMSI_EXPAND_BIG=%d -D SPH_HAMSI_SHORT=%d ",
+          ((opt_blake_compact)?1:0), ((opt_luffa_parallel)?1:0), (unsigned int)opt_keccak_unroll, opt_hamsi_expand_big, ((opt_hamsi_short)?1:0));
+  strcat(data->compiler_options, buf);
+
+  sprintf(buf, "ku%ubig%u%s%s%s", (unsigned int)opt_keccak_unroll, (unsigned int)opt_hamsi_expand_big, ((opt_blake_compact)?"bc":""), ((opt_luffa_parallel)?"lp":""), ((opt_hamsi_short)?"hs":""));
   strcat(data->binary_filename, buf);
 }
 
@@ -600,17 +612,18 @@ static algorithm_settings_t algos[] = {
 
   { "twecoin", ALGO_TWE, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 0, 0, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, twecoin_regenhash, queue_sph_kernel, sha256, NULL},
   { "maxcoin", ALGO_KECCAK, 1, 256, 1, 4, 15, 0x0F, 0xFFFFULL, 0x000000ffUL, 0, 0, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, maxcoin_regenhash, queue_maxcoin_kernel, sha256, NULL},
-  { "darkcoin-mod", ALGO_X11, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 8 * 16 * 4194304, 0, darkcoin_regenhash, queue_darkcoin_mod_kernel, gen_hash, NULL},
 
-  { "marucoin", ALGO_X13, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 0, 0, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, marucoin_regenhash, queue_sph_kernel, gen_hash, append_hamsi_compiler_options},
-  { "marucoin-mod", ALGO_X13, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 12, 8 * 16 * 4194304, 0, marucoin_regenhash, queue_marucoin_mod_kernel, gen_hash, append_hamsi_compiler_options},
-  { "marucoin-modold", ALGO_X13, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 8 * 16 * 4194304, 0, marucoin_regenhash, queue_marucoin_mod_old_kernel, gen_hash, append_hamsi_compiler_options},
+  { "darkcoin-mod", ALGO_X11, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 8 * 16 * 4194304, 0, darkcoin_regenhash, queue_darkcoin_mod_kernel, gen_hash, append_x11_compiler_options},
+
+  { "marucoin", ALGO_X13, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 0, 0, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, marucoin_regenhash, queue_sph_kernel, gen_hash, append_x13_compiler_options},
+  { "marucoin-mod", ALGO_X13, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 12, 8 * 16 * 4194304, 0, marucoin_regenhash, queue_marucoin_mod_kernel, gen_hash, append_x13_compiler_options},
+  { "marucoin-modold", ALGO_X13, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 8 * 16 * 4194304, 0, marucoin_regenhash, queue_marucoin_mod_old_kernel, gen_hash, append_x13_compiler_options},
 
-  { "x14", ALGO_X14, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 13, 8 * 16 * 4194304, 0, x14_regenhash, queue_x14_kernel, gen_hash, append_hamsi_compiler_options},
-  { "x14old", ALGO_X14, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 8 * 16 * 4194304, 0, x14_regenhash, queue_x14_old_kernel, gen_hash, append_hamsi_compiler_options},
+  { "x14", ALGO_X14, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 13, 8 * 16 * 4194304, 0, x14_regenhash, queue_x14_kernel, gen_hash, append_x13_compiler_options},
+  { "x14old", ALGO_X14, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 8 * 16 * 4194304, 0, x14_regenhash, queue_x14_old_kernel, gen_hash, append_x13_compiler_options},
 
-  { "bitblock", ALGO_X15, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 14, 4 * 16 * 4194304, 0, bitblock_regenhash, queue_bitblock_kernel, gen_hash, append_hamsi_compiler_options},
-  { "bitblockold", ALGO_X15, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 4 * 16 * 4194304, 0, bitblock_regenhash, queue_bitblockold_kernel, gen_hash, append_hamsi_compiler_options},
+  { "bitblock", ALGO_X15, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 14, 4 * 16 * 4194304, 0, bitblock_regenhash, queue_bitblock_kernel, gen_hash, append_x13_compiler_options},
+  { "bitblockold", ALGO_X15, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 4 * 16 * 4194304, 0, bitblock_regenhash, queue_bitblockold_kernel, gen_hash, append_x13_compiler_options},
 
   { "talkcoin-mod", ALGO_NIST, 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4,  8 * 16 * 4194304, 0, talkcoin_regenhash, queue_talkcoin_mod_kernel, gen_hash, NULL},
 
@@ -704,7 +717,7 @@ void set_algorithm(algorithm_t* algo, const char* newname_alias)
   const char* newname;
   //load previous algorithm nfactor in case nfactor was applied before algorithm... or default to 10
   uint8_t old_nfactor = ((algo->nfactor)?algo->nfactor:0);
-  uint8_t nfactor = 0;
+  uint8_t nfactor = 10;
 
   if (!(newname = lookup_algorithm_alias(newname_alias, &nfactor)))
     newname = newname_alias;
@@ -716,9 +729,6 @@ void set_algorithm(algorithm_t* algo, const char* newname_alias)
     nfactor = old_nfactor;
 
   set_algorithm_nfactor(algo, nfactor);
-  // Doesn't matter for non-scrypt algorithms
-  if (nfactor > 0)
-    set_algorithm_nfactor(algo, nfactor);
 }
 
 void set_algorithm_nfactor(algorithm_t* algo, const uint8_t nfactor)

kernel/bitblock.cl

@@ -5,7 +5,7 @@
  *
  * Copyright (c) 2014  phm
  * Copyright (c) 2014 Girino Vey
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -13,10 +13,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -71,12 +71,19 @@ typedef int sph_s32;
 #define SPH_JH_64 1
 #define SPH_SIMD_NOCOPY 0
 #define SPH_KECCAK_NOCOPY 0
-#define SPH_COMPACT_BLAKE_64 0
-#define SPH_LUFFA_PARALLEL 0
 #define SPH_SMALL_FOOTPRINT_GROESTL 0
 #define SPH_GROESTL_BIG_ENDIAN 0
 #define SPH_CUBEHASH_UNROLL 0
-#define SPH_KECCAK_UNROLL   1
+
+#ifndef SPH_COMPACT_BLAKE_64
+  #define SPH_COMPACT_BLAKE_64 0
+#endif
+#ifndef SPH_LUFFA_PARALLEL
+  #define SPH_LUFFA_PARALLEL 0
+#endif
+#ifndef SPH_KECCAK_UNROLL
+  #define SPH_KECCAK_UNROLL 0
+#endif
 #ifndef SPH_HAMSI_EXPAND_BIG
   #define SPH_HAMSI_EXPAND_BIG 1
 #endif

kernel/bitblockold.cl

@@ -5,7 +5,7 @@
  *
  * Copyright (c) 2014  phm
  * Copyright (c) 2014 Girino Vey
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -13,10 +13,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -69,12 +69,19 @@ typedef long sph_s64;
 #define SPH_JH_64 1
 #define SPH_SIMD_NOCOPY 0
 #define SPH_KECCAK_NOCOPY 0
-#define SPH_COMPACT_BLAKE_64 0
-#define SPH_LUFFA_PARALLEL 0
 #define SPH_SMALL_FOOTPRINT_GROESTL 0
 #define SPH_GROESTL_BIG_ENDIAN 0
 #define SPH_CUBEHASH_UNROLL 0
-#define SPH_KECCAK_UNROLL   1
+
+#ifndef SPH_COMPACT_BLAKE_64
+  #define SPH_COMPACT_BLAKE_64 0
+#endif
+#ifndef SPH_LUFFA_PARALLEL
+  #define SPH_LUFFA_PARALLEL 0
+#endif
+#ifndef SPH_KECCAK_UNROLL
+  #define SPH_KECCAK_UNROLL 0
+#endif
 #ifndef SPH_HAMSI_EXPAND_BIG
   #define SPH_HAMSI_EXPAND_BIG 4
 #endif
@@ -796,7 +803,7 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
     }
 
     barrier(CLK_LOCAL_MEM_FENCE);
-	
+
     #ifdef INPUT_BIG_LOCAL
       __local sph_u32 T512_L[1024];
       __constant const sph_u32 *T512_C = &T512[0][0];
@@ -819,7 +826,7 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
     	mixtab3[i] = mixtab3_c[i];
     }
     barrier(CLK_LOCAL_MEM_FENCE);
-	
+
 
     for (int i = 0; i < 8; i++) {
         hash.h8[i] = hashes[gid-offset].h8[i];
@@ -1002,7 +1009,7 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
     hash.h4[15] = SWAP4(S30);
 
     }
-	
+
 	//shabal
 	{
     sph_u32 A00 = A_init_512[0], A01 = A_init_512[1], A02 = A_init_512[2], A03 = A_init_512[3], A04 = A_init_512[4], A05 = A_init_512[5], A06 = A_init_512[6], A07 = A_init_512[7],
@@ -1065,9 +1072,9 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
 	hash.h4[12] = BC;
 	hash.h4[13] = BD;
 	hash.h4[14] = BE;
-	hash.h4[15] = BF;	
+	hash.h4[15] = BF;
 	}
-	
+
 	//whirlpool
 	{
     sph_u64 n0, n1, n2, n3, n4, n5, n6, n7;
@@ -1153,7 +1160,7 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
     state[7] ^= n7 ^ 0x2000000000000;
 
     for (unsigned i = 0; i < 8; i ++)
-	hash.h8[i] = state[i];	
+	hash.h8[i] = state[i];
 	}
 
     bool result = (hash.h8[3] <= target);

kernel/darkcoin-mod.cl

@@ -4,7 +4,7 @@
  *
  * Copyright (c) 2014  phm
  * Copyright (c) 2014 Girino Vey
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -12,10 +12,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -34,54 +34,57 @@
 #endif
 
 #ifndef DARKCOIN_MOD_CL
-  #define DARKCOIN_MOD_CL
+#define DARKCOIN_MOD_CL
 
-  #if __ENDIAN_LITTLE__
+#if __ENDIAN_LITTLE__
   #define SPH_LITTLE_ENDIAN 1
-  #else
+#else
   #define SPH_BIG_ENDIAN 1
-  #endif
+#endif
 
-  #define SPH_UPTR sph_u64
-  typedef unsigned int sph_u32;
-  typedef int sph_s32;
+#define SPH_UPTR sph_u64
+typedef unsigned int sph_u32;
+typedef int sph_s32;
 
-  #ifndef __OPENCL_VERSION__
+#ifndef __OPENCL_VERSION__
   typedef unsigned long long sph_u64;
   typedef long long sph_s64;
-  #else
+#else
   typedef unsigned long sph_u64;
   typedef long sph_s64;
-  #endif
-
-  #define SPH_64 1
-  #define SPH_64_TRUE 1
-
-  #define SPH_C32(x) ((sph_u32)(x ## U))
-  #define SPH_T32(x) (as_uint(x))
-  #define SPH_ROTL32(x, n) rotate(as_uint(x), as_uint(n))
-  #define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n)))
-
-  #define SPH_C64(x) ((sph_u64)(x ## UL))
-  #define SPH_T64(x) (as_ulong(x))
-  #define SPH_ROTL64(x, n) rotate(as_ulong(x), (n) & 0xFFFFFFFFFFFFFFFFUL)
-  #define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n)))
-
-  #define SPH_ECHO_64 1
-  #define SPH_KECCAK_64 1
-  #define SPH_JH_64 1
-  #define SPH_SIMD_NOCOPY 0
-  #define SPH_KECCAK_NOCOPY 0
-  #define SPH_COMPACT_BLAKE_64 0
-  #define SPH_LUFFA_PARALLEL 0
-  #ifndef SPH_SMALL_FOOTPRINT_GROESTL
-  #define SPH_SMALL_FOOTPRINT_GROESTL 0
 #endif
 
+#define SPH_64 1
+#define SPH_64_TRUE 1
+
+#define SPH_C32(x) ((sph_u32)(x ## U))
+#define SPH_T32(x) (as_uint(x))
+#define SPH_ROTL32(x, n) rotate(as_uint(x), as_uint(n))
+#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n)))
+
+#define SPH_C64(x) ((sph_u64)(x ## UL))
+#define SPH_T64(x) (as_ulong(x))
+#define SPH_ROTL64(x, n) rotate(as_ulong(x), (n) & 0xFFFFFFFFFFFFFFFFUL)
+#define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n)))
+
+#define SPH_ECHO_64 1
+#define SPH_KECCAK_64 1
+#define SPH_JH_64 1
+#define SPH_SIMD_NOCOPY 0
+#define SPH_KECCAK_NOCOPY 0
+#define SPH_SMALL_FOOTPRINT_GROESTL 0
 #define SPH_GROESTL_BIG_ENDIAN 0
-
 #define SPH_CUBEHASH_UNROLL 0
-#define SPH_KECCAK_UNROLL   0
+
+#ifndef SPH_COMPACT_BLAKE_64
+  #define SPH_COMPACT_BLAKE_64 0
+#endif
+#ifndef SPH_LUFFA_PARALLEL
+  #define SPH_LUFFA_PARALLEL 0
+#endif
+#ifndef SPH_KECCAK_UNROLL
+  #define SPH_KECCAK_UNROLL   0
+#endif
 
 //#include "aes_helper.cl"
 #include "blake.cl"

kernel/groestl.cl

@@ -5,7 +5,7 @@
  * ==========================(LICENSE BEGIN)============================
  *
  * Copyright (c) 2007-2010  Projet RNRT SAPHIR
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -13,10 +13,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -1400,7 +1400,7 @@ __constant static const sph_u64 T7[] = {
 
 #endif
 
-#define PERM_BIG_P(a)   do { \
+/*#define PERM_BIG_P(a)   do { \
     int r; \
     for (r = 0; r < 14; r += 2) { \
       ROUND_BIG_P(a, r + 0); \
@@ -1415,4 +1415,18 @@ __constant static const sph_u64 T7[] = {
       ROUND_BIG_Q(a, r + 1); \
     } \
   } while (0)
+*/
 
+#define PERM_BIG_P(a)   do { \
+    int r; \
+    for (r = 0; r < 14; ++r) { \
+      ROUND_BIG_P(a, r); \
+    } \
+  } while (0)
+
+#define PERM_BIG_Q(a)   do { \
+    int r; \
+    for (r = 0; r < 14; ++r) { \
+      ROUND_BIG_Q(a, r); \
+    } \
+  } while (0)

kernel/marucoin-mod.cl

@@ -5,7 +5,7 @@
  *
  * Copyright (c) 2014  phm
  * Copyright (c) 2014 Girino Vey
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -13,10 +13,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -71,13 +71,20 @@ typedef int sph_s32;
 #define SPH_JH_64 1
 #define SPH_SIMD_NOCOPY 0
 #define SPH_KECCAK_NOCOPY 0
-#define SPH_COMPACT_BLAKE_64 0
-#define SPH_LUFFA_PARALLEL 0
 #define SPH_SMALL_FOOTPRINT_GROESTL 0
 #define SPH_GROESTL_BIG_ENDIAN 0
 #define SPH_CUBEHASH_UNROLL 0
-#define SPH_KECCAK_UNROLL   1
-#if !defined SPH_HAMSI_EXPAND_BIG
+
+#ifndef SPH_COMPACT_BLAKE_64
+  #define SPH_COMPACT_BLAKE_64 0
+#endif
+#ifndef SPH_LUFFA_PARALLEL
+  #define SPH_LUFFA_PARALLEL 0
+#endif
+#ifndef SPH_KECCAK_UNROLL
+  #define SPH_KECCAK_UNROLL 0
+#endif
+#ifndef SPH_HAMSI_EXPAND_BIG
   #define SPH_HAMSI_EXPAND_BIG 1
 #endif
 

kernel/marucoin-modold.cl

@@ -5,7 +5,7 @@
  *
  * Copyright (c) 2014  phm
  * Copyright (c) 2014 Girino Vey
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -13,10 +13,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -69,12 +69,19 @@ typedef long sph_s64;
 #define SPH_JH_64 1
 #define SPH_SIMD_NOCOPY 0
 #define SPH_KECCAK_NOCOPY 0
-#define SPH_COMPACT_BLAKE_64 0
-#define SPH_LUFFA_PARALLEL 0
 #define SPH_SMALL_FOOTPRINT_GROESTL 0
 #define SPH_GROESTL_BIG_ENDIAN 0
 #define SPH_CUBEHASH_UNROLL 0
-#define SPH_KECCAK_UNROLL   1
+
+#ifndef SPH_COMPACT_BLAKE_64
+  #define SPH_COMPACT_BLAKE_64 0
+#endif
+#ifndef SPH_LUFFA_PARALLEL
+  #define SPH_LUFFA_PARALLEL 0
+#endif
+#ifndef SPH_KECCAK_UNROLL
+  #define SPH_KECCAK_UNROLL 0
+#endif
 #ifndef SPH_HAMSI_EXPAND_BIG
   #define SPH_HAMSI_EXPAND_BIG 4
 #endif

kernel/x14.cl

@@ -5,7 +5,7 @@
  *
  * Copyright (c) 2014  phm
  * Copyright (c) 2014 Girino Vey
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -13,10 +13,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -71,12 +71,19 @@ typedef int sph_s32;
 #define SPH_JH_64 1
 #define SPH_SIMD_NOCOPY 0
 #define SPH_KECCAK_NOCOPY 0
-#define SPH_COMPACT_BLAKE_64 0
-#define SPH_LUFFA_PARALLEL 0
 #define SPH_SMALL_FOOTPRINT_GROESTL 0
 #define SPH_GROESTL_BIG_ENDIAN 0
 #define SPH_CUBEHASH_UNROLL 0
-#define SPH_KECCAK_UNROLL   1
+
+#ifndef SPH_COMPACT_BLAKE_64
+  #define SPH_COMPACT_BLAKE_64 0
+#endif
+#ifndef SPH_LUFFA_PARALLEL
+  #define SPH_LUFFA_PARALLEL 0
+#endif
+#ifndef SPH_KECCAK_UNROLL
+  #define SPH_KECCAK_UNROLL 0
+#endif
 #ifndef SPH_HAMSI_EXPAND_BIG
   #define SPH_HAMSI_EXPAND_BIG 1
 #endif

kernel/x14old.cl

@@ -5,7 +5,7 @@
  *
  * Copyright (c) 2014  phm
  * Copyright (c) 2014 Girino Vey
- * 
+ *
  * Permission is hereby granted, free of charge, to any person obtaining
  * a copy of this software and associated documentation files (the
  * "Software"), to deal in the Software without restriction, including
@@ -13,10 +13,10 @@
  * distribute, sublicense, and/or sell copies of the Software, and to
  * permit persons to whom the Software is furnished to do so, subject to
  * the following conditions:
- * 
+ *
  * The above copyright notice and this permission notice shall be
  * included in all copies or substantial portions of the Software.
- * 
+ *
  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
  * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
  * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
@@ -69,12 +69,19 @@ typedef long sph_s64;
 #define SPH_JH_64 1
 #define SPH_SIMD_NOCOPY 0
 #define SPH_KECCAK_NOCOPY 0
-#define SPH_COMPACT_BLAKE_64 0
-#define SPH_LUFFA_PARALLEL 0
 #define SPH_SMALL_FOOTPRINT_GROESTL 0
 #define SPH_GROESTL_BIG_ENDIAN 0
 #define SPH_CUBEHASH_UNROLL 0
-#define SPH_KECCAK_UNROLL   1
+
+#ifndef SPH_COMPACT_BLAKE_64
+  #define SPH_COMPACT_BLAKE_64 0
+#endif
+#ifndef SPH_LUFFA_PARALLEL
+  #define SPH_LUFFA_PARALLEL 0
+#endif
+#ifndef SPH_KECCAK_UNROLL
+  #define SPH_KECCAK_UNROLL 0
+#endif
 #ifndef SPH_HAMSI_EXPAND_BIG
   #define SPH_HAMSI_EXPAND_BIG 4
 #endif
@@ -819,7 +826,7 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
     	mixtab3[i] = mixtab3_c[i];
     }
     barrier(CLK_LOCAL_MEM_FENCE);
-	
+
 
     for (int i = 0; i < 8; i++) {
         hash.h8[i] = hashes[gid-offset].h8[i];
@@ -1002,7 +1009,7 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
     hash.h4[15] = SWAP4(S30);
 
     }
-	
+
 	//shabal
 	{
     sph_u32 A00 = A_init_512[0], A01 = A_init_512[1], A02 = A_init_512[2], A03 = A_init_512[3], A04 = A_init_512[4], A05 = A_init_512[5], A06 = A_init_512[6], A07 = A_init_512[7],
@@ -1065,9 +1072,9 @@ __kernel void search10(__global hash_t* hashes, __global uint* output, const ulo
 	hash.h4[12] = BC;
 	hash.h4[13] = BD;
 	hash.h4[14] = BE;
-	hash.h4[15] = BF;	
+	hash.h4[15] = BF;
 	}
-	
+
     bool result = (hash.h8[3] <= target);
     if (result)
       output[atomic_inc(output+0xFF)] = SWAP4(gid);

miner.h

@@ -1047,6 +1047,11 @@ extern bool opt_worktime;
 extern int swork_id;
 extern int opt_tcp_keepalive;
 extern bool opt_incognito;
+
+// Xn Algorithm options
+extern int opt_keccak_unroll;
+extern bool opt_blake_compact;
+extern bool opt_luffa_parallel;
 extern int opt_hamsi_expand_big;
 extern bool opt_hamsi_short;
 

sgminer.c

@@ -102,9 +102,16 @@ time_t last_getwork;
 int nDevs;
 int opt_dynamic_interval = 7;
 int opt_g_threads = -1;
+bool opt_restart = true;
+
+/*****************************************
+ * Xn Algorithm options
+ *****************************************/
 int opt_hamsi_expand_big = 4;
+int opt_keccak_unroll = 0;
 bool opt_hamsi_short = false;
-bool opt_restart = true;
+bool opt_blake_compact = false;
+bool opt_luffa_parallel = false;
 
 struct list_head scan_devices;
 bool devices_enabled[MAX_DEVICES];
@@ -1327,6 +1334,9 @@ struct opt_table opt_config_table[] = {
   OPT_WITHOUT_ARG("--balance",
       set_balance, &pool_strategy,
       "Change multipool strategy from failover to even share balance"),
+  OPT_WITHOUT_ARG("--blake-compact",
+      opt_set_bool, &opt_blake_compact,
+      "Set SPH_COMPACT_BLAKE64 for Xn derived algorithms (Can give better hashrate for some GPUs)"),
 #ifdef HAVE_CURSES
   OPT_WITHOUT_ARG("--compact",
       opt_set_bool, &opt_compact,
@@ -1402,15 +1412,21 @@ struct opt_table opt_config_table[] = {
       set_default_gpu_vddc, NULL, NULL,
       "Set the GPU voltage in Volts - one value for all or separate by commas for per card"),
 #endif
-  OPT_WITH_ARG("--lookup-gap",
-      set_default_lookup_gap, NULL, NULL,
-      "Set GPU lookup gap for scrypt mining, comma separated"),
   OPT_WITH_ARG("--hamsi-expand-big",
       set_int_1_to_10, opt_show_intval, &opt_hamsi_expand_big,
       "Set SPH_HAMSI_EXPAND_BIG for X13 derived algorithms (1 or 4 are common)"),
   OPT_WITHOUT_ARG("--hamsi-short",
       opt_set_bool, &opt_hamsi_short,
       "Set SPH_HAMSI_SHORT for X13 derived algorithms (Can give better hashrate for some GPUs)"),
+  OPT_WITH_ARG("--keccak-unroll",
+      set_int_0_to_9999, opt_show_intval, &opt_keccak_unroll,
+      "Set SPH_KECCAK_UNROLL for Xn derived algorithms (Default: 0)"),
+  OPT_WITH_ARG("--lookup-gap",
+      set_default_lookup_gap, NULL, NULL,
+      "Set GPU lookup gap for scrypt mining, comma separated"),
+  OPT_WITHOUT_ARG("--luffa-parallel",
+      opt_set_bool, &opt_luffa_parallel,
+      "Set SPH_LUFFA_PARALLEL for Xn derived algorithms (Can give better hashrate for some GPUs)"),
 #ifdef HAVE_CURSES
   OPT_WITHOUT_ARG("--incognito",
       opt_set_bool, &opt_incognito,