Browse Source

Merge branch 'groestlcoin'

Conflicts:
	Makefile.am
	configure.ac
	driver-opencl.c
	miner.h
	ocl.c
	sgminer.c
djm34
phm 11 years ago
parent
commit
983eae1142
  1. 1
      Makefile.am
  2. 1
      configure.ac
  3. 6
      driver-opencl.c
  4. 153
      groestlcoin.c
  5. 10
      groestlcoin.h
  6. 251
      kernel/groestlcoin.cl
  7. 1
      miner.h
  8. 5
      ocl.c
  9. 6
      sgminer.c

1
Makefile.am

@ -49,6 +49,7 @@ sgminer_SOURCES += myriadcoin-groestl.c myriadcoin-groestl.h @@ -49,6 +49,7 @@ sgminer_SOURCES += myriadcoin-groestl.c myriadcoin-groestl.h
sgminer_SOURCES += fuguecoin.c fuguecoin.h
sgminer_SOURCES += inkcoin.c inkcoin.h
sgminer_SOURCES += animecoin.c animecoin.h
sgminer_SOURCES += groestlcoin.c groestlcoin.h
sgminer_SOURCES += kernel/*.cl
bin_SCRIPTS = $(top_srcdir)/kernel/*.cl

1
configure.ac

@ -352,6 +352,7 @@ AC_DEFINE_UNQUOTED([MYRIADCOIN_GROESTL_KERNNAME], ["myriadcoin-groestl"], [Filen @@ -352,6 +352,7 @@ AC_DEFINE_UNQUOTED([MYRIADCOIN_GROESTL_KERNNAME], ["myriadcoin-groestl"], [Filen
AC_DEFINE_UNQUOTED([FUGUECOIN_KERNNAME], ["fuguecoin"], [Filename for FugueCoin optimised kernel])
AC_DEFINE_UNQUOTED([INKCOIN_KERNNAME], ["inkcoin"], [Filename for InkCoin optimised kernel])
AC_DEFINE_UNQUOTED([ANIMECOIN_KERNNAME], ["animecoin"], [Filename for AnimeCoin optimised kernel])
AC_DEFINE_UNQUOTED([GROESTLCOIN_KERNNAME], ["groestlcoin"], [Filename for GroestlCoin optimised kernel])
AC_SUBST(OPENCL_LIBS)
AC_SUBST(OPENCL_FLAGS)

6
driver-opencl.c

@ -221,6 +221,8 @@ static enum cl_kernels select_kernel(char *arg) @@ -221,6 +221,8 @@ static enum cl_kernels select_kernel(char *arg)
return KL_INKCOIN;
if (!strcmp(arg, ANIMECOIN_KERNNAME))
return KL_ANIMECOIN;
if (!strcmp(arg, GROESTLCOIN_KERNNAME))
return KL_GROESTLCOIN;
return KL_NONE;
}
@ -1382,6 +1384,9 @@ static bool opencl_thread_prepare(struct thr_info *thr) @@ -1382,6 +1384,9 @@ static bool opencl_thread_prepare(struct thr_info *thr)
case KL_ANIMECOIN:
cgpu->kname = ANIMECOIN_KERNNAME;
break;
case KL_GROESTLCOIN:
cgpu->kname = GROESTLCOIN_KERNNAME;
break;
default:
break;
}
@ -1424,6 +1429,7 @@ static bool opencl_thread_init(struct thr_info *thr) @@ -1424,6 +1429,7 @@ static bool opencl_thread_init(struct thr_info *thr)
case KL_FUGUECOIN:
case KL_INKCOIN:
case KL_ANIMECOIN:
case KL_GROESTLCOIN:
thrdata->queue_kernel_parameters = &queue_sph_kernel;
break;
default:

153
groestlcoin.c

@ -0,0 +1,153 @@ @@ -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 <stdlib.h>
#include <stdint.h>
#include <string.h>
#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;
}

10
groestlcoin.h

@ -0,0 +1,10 @@ @@ -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 */

251
kernel/groestlcoin.cl

@ -0,0 +1,251 @@ @@ -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 <phm@inbox.com>
*/
#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

1
miner.h

@ -388,6 +388,7 @@ enum cl_kernels { @@ -388,6 +388,7 @@ enum cl_kernels {
KL_DARKCOIN, // kernels starting from this will have difficulty calculated by using bitcoin algorithm
KL_MYRIADCOIN_GROESTL,
KL_FUGUECOIN,
KL_GROESTLCOIN,
};
enum dev_reason {

5
ocl.c

@ -489,6 +489,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) @@ -489,6 +489,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
strcpy(filename, ANIMECOIN_KERNNAME".cl");
strcpy(binaryfilename, ANIMECOIN_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;
}

6
sgminer.c

@ -4248,6 +4248,9 @@ void write_config(FILE *fcfg) @@ -4248,6 +4248,9 @@ void write_config(FILE *fcfg)
case KL_ANIMECOIN:
fprintf(fcfg, ANIMECOIN_KERNNAME);
break;
case KL_GROESTLCOIN:
fprintf(fcfg, GROESTLCOIN_KERNNAME);
break;
}
}
@ -6073,6 +6076,9 @@ static void rebuild_nonce(struct work *work, uint32_t nonce) @@ -6073,6 +6076,9 @@ static void rebuild_nonce(struct work *work, uint32_t nonce)
case KL_ANIMECOIN:
animecoin_regenhash(work);
break;
case KL_GROESTLCOIN:
groestlcoin_regenhash(work);
break;
default:
scrypt_regenhash(work);
break;

Loading…
Cancel
Save