From b3c07acd0846d2c311b1a7e6961078d1efa7e71f Mon Sep 17 00:00:00 2001 From: ystarnaud Date: Mon, 28 Jul 2014 14:43:11 -0400 Subject: [PATCH] Xn Algorithms Fine Tuning Added options to fine tune X11-15 algorithms --- algorithm.c | 42 +++++++++++++-------- kernel/bitblock.cl | 19 +++++++--- kernel/bitblockold.cl | 31 +++++++++------ kernel/darkcoin-mod.cl | 79 ++++++++++++++++++++------------------- kernel/groestl.cl | 22 +++++++++-- kernel/marucoin-mod.cl | 21 +++++++---- kernel/marucoin-modold.cl | 19 +++++++--- kernel/x14.cl | 19 +++++++--- kernel/x14old.cl | 27 ++++++++----- miner.h | 5 +++ sgminer.c | 24 ++++++++++-- 11 files changed, 199 insertions(+), 109 deletions(-) diff --git a/algorithm.c b/algorithm.c index d41322fb..00fb5ebf 100644 --- a/algorithm.c +++ b/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) diff --git a/kernel/bitblock.cl b/kernel/bitblock.cl index 3c8232e8..3603a160 100644 --- a/kernel/bitblock.cl +++ b/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 diff --git a/kernel/bitblockold.cl b/kernel/bitblockold.cl index e5b0af15..87330c9d 100644 --- a/kernel/bitblockold.cl +++ b/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); diff --git a/kernel/darkcoin-mod.cl b/kernel/darkcoin-mod.cl index 0cae9aaf..d4240369 100644 --- a/kernel/darkcoin-mod.cl +++ b/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" diff --git a/kernel/groestl.cl b/kernel/groestl.cl index 4c040c01..80827c65 100644 --- a/kernel/groestl.cl +++ b/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) diff --git a/kernel/marucoin-mod.cl b/kernel/marucoin-mod.cl index f359a5a5..b0bf0dd3 100644 --- a/kernel/marucoin-mod.cl +++ b/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 diff --git a/kernel/marucoin-modold.cl b/kernel/marucoin-modold.cl index 474cc420..4faaad4e 100644 --- a/kernel/marucoin-modold.cl +++ b/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 diff --git a/kernel/x14.cl b/kernel/x14.cl index 2525a714..ff335182 100644 --- a/kernel/x14.cl +++ b/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 diff --git a/kernel/x14old.cl b/kernel/x14old.cl index 55f83235..0982bc39 100644 --- a/kernel/x14old.cl +++ b/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); diff --git a/miner.h b/miner.h index 4c0cd06e..c8190bd5 100644 --- a/miner.h +++ b/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; diff --git a/sgminer.c b/sgminer.c index 4acccceb..84ec2486 100644 --- a/sgminer.c +++ b/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,