From 764d8706b48cf80166bb5147290e714ea68865dc Mon Sep 17 00:00:00 2001 From: phm Date: Sat, 22 Mar 2014 23:50:21 +0100 Subject: [PATCH] Added support for GroestlCoin. --- Makefile.am | 1 + configure.ac | 1 + driver-opencl.c | 6 + groestlcoin.c | 153 +++++++++++++++++++++++++ groestlcoin.h | 10 ++ kernel/groestlcoin.cl | 251 ++++++++++++++++++++++++++++++++++++++++++ miner.h | 1 + ocl.c | 5 + sgminer.c | 6 + 9 files changed, 434 insertions(+) create mode 100644 groestlcoin.c create mode 100644 groestlcoin.h create mode 100644 kernel/groestlcoin.cl diff --git a/Makefile.am b/Makefile.am index 0c41690f..14ab288b 100644 --- a/Makefile.am +++ b/Makefile.am @@ -46,6 +46,7 @@ sgminer_SOURCES += darkcoin.c darkcoin.h sgminer_SOURCES += qubitcoin.c qubitcoin.h sgminer_SOURCES += quarkcoin.c quarkcoin.h sgminer_SOURCES += myriadcoin-groestl.c myriadcoin-groestl.h +sgminer_SOURCES += groestlcoin.c groestlcoin.h sgminer_SOURCES += kernel/*.cl bin_SCRIPTS = $(top_srcdir)/kernel/*.cl diff --git a/configure.ac b/configure.ac index bd59827e..8c1376f0 100644 --- a/configure.ac +++ b/configure.ac @@ -349,6 +349,7 @@ AC_DEFINE_UNQUOTED([DARKCOIN_KERNNAME], ["darkcoin"], [Filename for DarkCoin opt AC_DEFINE_UNQUOTED([QUBITCOIN_KERNNAME], ["qubitcoin"], [Filename for QubitCoin optimised kernel]) AC_DEFINE_UNQUOTED([QUARKCOIN_KERNNAME], ["quarkcoin"], [Filename for QuarkCoin optimised kernel]) AC_DEFINE_UNQUOTED([MYRIADCOIN_GROESTL_KERNNAME], ["myriadcoin-groestl"], [Filename for MyriadCoin-Groestl optimised kernel]) +AC_DEFINE_UNQUOTED([GROESTLCOIN_KERNNAME], ["groestlcoin"], [Filename for GroestlCoin optimised kernel]) AC_SUBST(OPENCL_LIBS) AC_SUBST(OPENCL_FLAGS) diff --git a/driver-opencl.c b/driver-opencl.c index 664f7e7e..00109797 100644 --- a/driver-opencl.c +++ b/driver-opencl.c @@ -215,6 +215,8 @@ static enum cl_kernels select_kernel(char *arg) return KL_QUARKCOIN; if (!strcmp(arg, MYRIADCOIN_GROESTL_KERNNAME)) return KL_MYRIADCOIN_GROESTL; + if (!strcmp(arg, GROESTLCOIN_KERNNAME)) + return KL_GROESTLCOIN; return KL_NONE; } @@ -1367,6 +1369,9 @@ static bool opencl_thread_prepare(struct thr_info *thr) case KL_MYRIADCOIN_GROESTL: cgpu->kname = MYRIADCOIN_GROESTL_KERNNAME; break; + case KL_GROESTLCOIN: + cgpu->kname = GROESTLCOIN_KERNNAME; + break; default: break; } @@ -1406,6 +1411,7 @@ static bool opencl_thread_init(struct thr_info *thr) case KL_QUBITCOIN: case KL_QUARKCOIN: case KL_MYRIADCOIN_GROESTL: + case KL_GROESTLCOIN: thrdata->queue_kernel_parameters = &queue_sph_kernel; break; default: diff --git a/groestlcoin.c b/groestlcoin.c new file mode 100644 index 00000000..48fe829a --- /dev/null +++ b/groestlcoin.c @@ -0,0 +1,153 @@ +/*- + * Copyright 2009 Colin Percival, 2014 phm + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND + * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE + * ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR OR CONTRIBUTORS BE LIABLE + * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL + * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS + * OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) + * HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY + * OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF + * SUCH DAMAGE. + * + * This file was originally written by Colin Percival as part of the Tarsnap + * online backup system. + */ + +#include "config.h" +#include "miner.h" + +#include +#include +#include + +#include "sph/sph_groestl.h" +#include "sph/sph_sha2.h" + +/* + * Encode a length len/4 vector of (uint32_t) into a length len vector of + * (unsigned char) in big-endian form. Assumes len is a multiple of 4. + */ +static inline void +be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len) +{ + uint32_t i; + + for (i = 0; i < len; i++) + dst[i] = htobe32(src[i]); +} + + +inline void groestlhash(void *state, const void *input) +{ + sph_groestl512_context ctx_groestl; + sph_sha256_context ctx_sha2; + + uint32_t hash[16]; + + sph_groestl512_init(&ctx_groestl); + sph_groestl512(&ctx_groestl, input, 80); + sph_groestl512_close(&ctx_groestl, (void*) hash); + + sph_groestl512_init(&ctx_groestl); + sph_groestl512(&ctx_groestl, hash, 64); + sph_groestl512_close(&ctx_groestl, (void*) hash); + + memcpy(state, hash, 32); +} + +static const uint32_t diff1targ = 0x0000ffff; + + +/* Used externally as confirmation of correct OCL code */ +int groestlcoin_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce) +{ + uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]); + uint32_t data[20], ohash[8]; + //char *scratchbuf; + + be32enc_vect(data, (const uint32_t *)pdata, 19); + data[19] = htobe32(nonce); + //scratchbuf = alloca(SCRATCHBUF_SIZE); + groestlhash(ohash, data); + tmp_hash7 = be32toh(ohash[7]); + + applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx", + (long unsigned int)Htarg, + (long unsigned int)diff1targ, + (long unsigned int)tmp_hash7); + if (tmp_hash7 > diff1targ) + return -1; + if (tmp_hash7 > Htarg) + return 0; + return 1; +} + +void groestlcoin_regenhash(struct work *work) +{ + uint32_t data[20]; + char *scratchbuf; + uint32_t *nonce = (uint32_t *)(work->data + 76); + uint32_t *ohash = (uint32_t *)(work->hash); + + be32enc_vect(data, (const uint32_t *)work->data, 19); + data[19] = htobe32(*nonce); + groestlhash(ohash, data); +} + +bool scanhash_groestlcoin(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate, + unsigned char *pdata, unsigned char __maybe_unused *phash1, + unsigned char __maybe_unused *phash, const unsigned char *ptarget, + uint32_t max_nonce, uint32_t *last_nonce, uint32_t n) +{ + uint32_t *nonce = (uint32_t *)(pdata + 76); + char *scratchbuf; + uint32_t data[20]; + uint32_t tmp_hash7; + uint32_t Htarg = le32toh(((const uint32_t *)ptarget)[7]); + bool ret = false; + + be32enc_vect(data, (const uint32_t *)pdata, 19); + + while(1) { + uint32_t ostate[8]; + + *nonce = ++n; + data[19] = (n); + groestlhash(ostate, data); + tmp_hash7 = (ostate[7]); + + applog(LOG_INFO, "data7 %08lx", + (long unsigned int)data[7]); + + if (unlikely(tmp_hash7 <= Htarg)) { + ((uint32_t *)pdata)[19] = htobe32(n); + *last_nonce = n; + ret = true; + break; + } + + if (unlikely((n >= max_nonce) || thr->work_restart)) { + *last_nonce = n; + break; + } + } + + return ret; +} + + + diff --git a/groestlcoin.h b/groestlcoin.h new file mode 100644 index 00000000..b3ca0477 --- /dev/null +++ b/groestlcoin.h @@ -0,0 +1,10 @@ +#ifndef GROESTLCOIN_H +#define GROESTLCOIN_H + +#include "miner.h" + +extern int groestlcoin_test(unsigned char *pdata, const unsigned char *ptarget, + uint32_t nonce); +extern void groestlcoin_regenhash(struct work *work); + +#endif /* GROESTLCOIN_H */ diff --git a/kernel/groestlcoin.cl b/kernel/groestlcoin.cl new file mode 100644 index 00000000..55b6c0c6 --- /dev/null +++ b/kernel/groestlcoin.cl @@ -0,0 +1,251 @@ +/* + * GroestlCoin kernel implementation. + * + * ==========================(LICENSE BEGIN)============================ + * + * Copyright (c) 2014 phm + * + * 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 + * without limitation the rights to use, copy, modify, merge, publish, + * 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. + * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY + * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, + * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE + * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + * + * ===========================(LICENSE END)============================= + * + * @author phm + */ + +#ifndef GROESTLCOIN_CL +#define GROESTLCOIN_CL + +#if __ENDIAN_LITTLE__ +#define SPH_LITTLE_ENDIAN 1 +#else +#define SPH_BIG_ENDIAN 1 +#endif + +#define SPH_UPTR sph_u64 + +typedef unsigned int sph_u32; +typedef int sph_s32; +#ifndef __OPENCL_VERSION__ +typedef unsigned long long sph_u64; +typedef long long sph_s64; +#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) ((x) & SPH_C32(0xFFFFFFFF)) +#define SPH_ROTL32(x, n) SPH_T32(((x) << (n)) | ((x) >> (32 - (n)))) +#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n))) + +#define SPH_C64(x) ((sph_u64)(x ## UL)) +#define SPH_T64(x) ((x) & SPH_C64(0xFFFFFFFFFFFFFFFF)) +#define SPH_ROTL64(x, n) SPH_T64(((x) << (n)) | ((x) >> (64 - (n)))) +#define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n))) + +#define SPH_ECHO_64 1 +#define SPH_SIMD_NOCOPY 0 +#define SPH_LUFFA_PARALLEL 0 +#define SPH_CUBEHASH_UNROLL 0 + +#include "groestl.cl" + +#define SWAP4(x) as_uint(as_uchar4(x).wzyx) +#define SWAP8(x) as_ulong(as_uchar8(x).s76543210) + +#if SPH_BIG_ENDIAN + #define ENC64E(x) SWAP8(x) + #define DEC64E(x) SWAP8(*(const __global sph_u64 *) (x)); +#else + #define ENC64E(x) (x) + #define DEC64E(x) (*(const __global sph_u64 *) (x)); +#endif + +#define ROL32(x, n) rotate(x, (uint) n) +#define SHR(x, n) ((x) >> n) +#define SWAP32(a) (as_uint(as_uchar4(a).wzyx)) + +#define S0(x) (ROL32(x, 25) ^ ROL32(x, 14) ^ SHR(x, 3)) +#define S1(x) (ROL32(x, 15) ^ ROL32(x, 13) ^ SHR(x, 10)) + +#define S2(x) (ROL32(x, 30) ^ ROL32(x, 19) ^ ROL32(x, 10)) +#define S3(x) (ROL32(x, 26) ^ ROL32(x, 21) ^ ROL32(x, 7)) + +#define P(a,b,c,d,e,f,g,h,x,K) \ +{ \ + temp1 = h + S3(e) + F1(e,f,g) + (K + x); \ + d += temp1; h = temp1 + S2(a) + F0(a,b,c); \ +} + +#define PLAST(a,b,c,d,e,f,g,h,x,K) \ +{ \ + d += h + S3(e) + F1(e,f,g) + (x + K); \ +} + +#define F0(y, x, z) bitselect(z, y, z ^ x) +#define F1(x, y, z) bitselect(z, y, x) + +#define R0 (W0 = S1(W14) + W9 + S0(W1) + W0) +#define R1 (W1 = S1(W15) + W10 + S0(W2) + W1) +#define R2 (W2 = S1(W0) + W11 + S0(W3) + W2) +#define R3 (W3 = S1(W1) + W12 + S0(W4) + W3) +#define R4 (W4 = S1(W2) + W13 + S0(W5) + W4) +#define R5 (W5 = S1(W3) + W14 + S0(W6) + W5) +#define R6 (W6 = S1(W4) + W15 + S0(W7) + W6) +#define R7 (W7 = S1(W5) + W0 + S0(W8) + W7) +#define R8 (W8 = S1(W6) + W1 + S0(W9) + W8) +#define R9 (W9 = S1(W7) + W2 + S0(W10) + W9) +#define R10 (W10 = S1(W8) + W3 + S0(W11) + W10) +#define R11 (W11 = S1(W9) + W4 + S0(W12) + W11) +#define R12 (W12 = S1(W10) + W5 + S0(W13) + W12) +#define R13 (W13 = S1(W11) + W6 + S0(W14) + W13) +#define R14 (W14 = S1(W12) + W7 + S0(W15) + W14) +#define R15 (W15 = S1(W13) + W8 + S0(W0) + W15) + +#define RD14 (S1(W12) + W7 + S0(W15) + W14) +#define RD15 (S1(W13) + W8 + S0(W0) + W15) + +__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) +__kernel void search(__global unsigned char* block, volatile __global uint* output, const ulong target) +{ + uint gid = get_global_id(0); + union { + unsigned char h1[64]; + uint h4[16]; + ulong h8[8]; + } hash; + + __local sph_u64 T0_L[256], T1_L[256], T2_L[256], T3_L[256], T4_L[256], T5_L[256], T6_L[256], T7_L[256]; + int init = get_local_id(0); + int step = get_local_size(0); + for (int i = init; i < 256; i += step) + { + T0_L[i] = T0[i]; + T1_L[i] = T1[i]; + T2_L[i] = T2[i]; + T3_L[i] = T3[i]; + T4_L[i] = T4[i]; + T5_L[i] = T5[i]; + T6_L[i] = T6[i]; + T7_L[i] = T7[i]; + } + barrier(CLK_LOCAL_MEM_FENCE); + +#define T0 T0_L +#define T1 T1_L +#define T2 T2_L +#define T3 T3_L +#define T4 T4_L +#define T5 T5_L +#define T6 T6_L +#define T7 T7_L + + // groestl + + sph_u64 H[16]; + for (unsigned int u = 0; u < 15; u ++) + H[u] = 0; +#if USE_LE + H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40); +#else + H[15] = (sph_u64)512; +#endif + + sph_u64 g[16], m[16]; + m[0] = DEC64E(block + 0 * 8); + m[1] = DEC64E(block + 1 * 8); + m[2] = DEC64E(block + 2 * 8); + m[3] = DEC64E(block + 3 * 8); + m[4] = DEC64E(block + 4 * 8); + m[5] = DEC64E(block + 5 * 8); + m[6] = DEC64E(block + 6 * 8); + m[7] = DEC64E(block + 7 * 8); + m[8] = DEC64E(block + 8 * 8); + m[9] = DEC64E(block + 9 * 8); + m[9] &= 0x00000000FFFFFFFF; + m[9] |= ((sph_u64) gid << 32); + m[10] = 0x80; + m[11] = 0; + m[12] = 0; + m[13] = 0; + m[14] = 0; + m[15] = 0x100000000000000; + for (unsigned int u = 0; u < 16; u ++) + g[u] = m[u] ^ H[u]; + PERM_BIG_P(g); + PERM_BIG_Q(m); + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= g[u] ^ m[u]; + sph_u64 xH[16]; + for (unsigned int u = 0; u < 16; u ++) + xH[u] = H[u]; + PERM_BIG_P(xH); + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= xH[u]; + for (unsigned int u = 0; u < 8; u ++) + hash.h8[u] = ENC64E(H[u + 8]); + + for (unsigned int u = 0; u < 15; u ++) + H[u] = 0; +#if USE_LE + H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40); +#else + H[15] = (sph_u64)512; +#endif + + m[0] = hash.h8[0]; + m[1] = hash.h8[1]; + m[2] = hash.h8[2]; + m[3] = hash.h8[3]; + m[4] = hash.h8[4]; + m[5] = hash.h8[5]; + m[6] = hash.h8[6]; + m[7] = hash.h8[7]; + for (unsigned int u = 0; u < 16; u ++) + g[u] = m[u] ^ H[u]; + m[8] = 0x80; g[8] = m[8] ^ H[8]; + m[9] = 0; g[9] = m[9] ^ H[9]; + m[10] = 0; g[10] = m[10] ^ H[10]; + m[11] = 0; g[11] = m[11] ^ H[11]; + m[12] = 0; g[12] = m[12] ^ H[12]; + m[13] = 0; g[13] = m[13] ^ H[13]; + m[14] = 0; g[14] = m[14] ^ H[14]; + m[15] = 0x100000000000000; g[15] = m[15] ^ H[15]; + PERM_BIG_P(g); + PERM_BIG_Q(m); + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= g[u] ^ m[u]; + for (unsigned int u = 0; u < 16; u ++) + xH[u] = H[u]; + PERM_BIG_P(xH); + for (unsigned int u = 0; u < 16; u ++) + H[u] ^= xH[u]; + for (unsigned int u = 0; u < 8; u ++) + hash.h8[u] = H[u + 8]; + + bool result = (hash.h8[3] <= target); + if (result) + output[output[0xFF]++] = SWAP4(gid); +} + +#endif // GROESTLCOIN_CL diff --git a/miner.h b/miner.h index 8ef6bbc9..c9ffd1b0 100644 --- a/miner.h +++ b/miner.h @@ -385,6 +385,7 @@ enum cl_kernels { KL_QUBITCOIN, KL_DARKCOIN, // kernels starting from this will have difficulty calculated by using bitcoin algorithm KL_MYRIADCOIN_GROESTL, + KL_GROESTLCOIN, }; enum dev_reason { diff --git a/ocl.c b/ocl.c index 3bdd5b21..0b917d0f 100644 --- a/ocl.c +++ b/ocl.c @@ -474,6 +474,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) strcpy(filename, MYRIADCOIN_GROESTL_KERNNAME".cl"); strcpy(binaryfilename, MYRIADCOIN_GROESTL_KERNNAME); break; + case KL_GROESTLCOIN: + applog(LOG_WARNING, "Kernel groestlcoin is experimental."); + strcpy(filename, GROESTLCOIN_KERNNAME".cl"); + strcpy(binaryfilename, GROESTLCOIN_KERNNAME); + break; case KL_NONE: /* Shouldn't happen */ break; } diff --git a/sgminer.c b/sgminer.c index 4723b2ef..c3b2d536 100644 --- a/sgminer.c +++ b/sgminer.c @@ -4239,6 +4239,9 @@ void write_config(FILE *fcfg) case KL_MYRIADCOIN_GROESTL: fprintf(fcfg, MYRIADCOIN_GROESTL_KERNNAME); break; + case KL_GROESTLCOIN: + fprintf(fcfg, GROESTLCOIN_KERNNAME); + break; } } @@ -6055,6 +6058,9 @@ static void rebuild_nonce(struct work *work, uint32_t nonce) case KL_MYRIADCOIN_GROESTL: myriadcoin_groestl_regenhash(work); break; + case KL_GROESTLCOIN: + groestlcoin_regenhash(work); + break; default: scrypt_regenhash(work); break;