Browse Source

add "blake" 256, 14 rounds (for NEOS blake, not BlakeCoin)

also remove "missing" file, its old and not compatible with ubuntu 14.04
master
Tanguy Pruvot 10 years ago
parent
commit
c17d11e377
  1. 3
      Makefile.am
  2. 494
      blake32.cu
  3. 2
      build.sh
  4. 4
      ccminer.vcxproj
  5. 3
      ccminer.vcxproj.filters
  6. 2
      configure.sh
  7. 48
      cpu-miner.c
  8. 5
      miner.h
  9. 367
      missing
  10. 4
      quark/cuda_checkhash.cu
  11. 4
      util.c

3
Makefile.am

@ -32,7 +32,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \
quark/cuda_jh512.cu quark/cuda_quark_blake512.cu quark/cuda_quark_groestl512.cu quark/cuda_skein512.cu \ quark/cuda_jh512.cu quark/cuda_quark_blake512.cu quark/cuda_quark_groestl512.cu quark/cuda_skein512.cu \
quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu quark/quarkcoin.cu quark/animecoin.cu \ quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu quark/quarkcoin.cu quark/animecoin.cu \
quark/cuda_quark_compactionTest.cu \ quark/cuda_quark_compactionTest.cu \
cuda_nist5.cu \ cuda_nist5.cu blake32.cu \
sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \ sph/bmw.c sph/blake.c sph/groestl.c sph/jh.c sph/keccak.c sph/skein.c \
sph/cubehash.c sph/echo.c sph/luffa.c sph/sha2.c sph/shavite.c sph/simd.c \ sph/cubehash.c sph/echo.c sph/luffa.c sph/sha2.c sph/shavite.c sph/simd.c \
sph/hamsi.c sph/hamsi_helper.c sph/sph_hamsi.h \ sph/hamsi.c sph/hamsi_helper.c sph/sph_hamsi.h \
@ -43,6 +43,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \
x15/x14.cu x15/x15.cu x15/cuda_x14_shabal512.cu x15/cuda_x15_whirlpool.cu x15/whirlpool.cu \ x15/x14.cu x15/x15.cu x15/cuda_x14_shabal512.cu x15/cuda_x15_whirlpool.cu x15/whirlpool.cu \
x17/x17.cu x17/cuda_x17_haval512.cu x17/cuda_x17_sha512.cu x17/x17.cu x17/cuda_x17_haval512.cu x17/cuda_x17_sha512.cu
ccminer_LDFLAGS = $(PTHREAD_FLAGS) @CUDA_LDFLAGS@ ccminer_LDFLAGS = $(PTHREAD_FLAGS) @CUDA_LDFLAGS@
ccminer_LDADD = @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @WS2_LIBS@ @CUDA_LIBS@ @OPENMP_CFLAGS@ @LIBS@ ccminer_LDADD = @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @WS2_LIBS@ @CUDA_LIBS@ @OPENMP_CFLAGS@ @LIBS@
ccminer_CPPFLAGS = -msse2 @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES) -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME ccminer_CPPFLAGS = -msse2 @LIBCURL_CPPFLAGS@ @OPENMP_CFLAGS@ $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES) -DSCRYPT_KECCAK512 -DSCRYPT_CHACHA -DSCRYPT_CHOOSE_COMPILETIME

494
blake32.cu

@ -0,0 +1,494 @@
/**
* Blake-256 Cuda Kernel (Tested on SM 5.0)
*
* Tanguy Pruvot - Aug. 2014
*/
#include "miner.h"
extern "C" {
#include "sph/sph_blake.h"
#include <stdint.h>
#include <memory.h>
}
/* hash by cpu with blake 256 */
extern "C" void blake32hash(void *output, const void *input)
{
unsigned char hash[64];
sph_blake256_context ctx;
sph_blake256_init(&ctx);
sph_blake256(&ctx, input, 80);
sph_blake256_close(&ctx, hash);
memcpy(output, hash, 32);
}
#include "cuda_helper.h"
#if __CUDA_ARCH__ < 350
// Kepler (Compute 3.0) + Host
#define ROTR32(x, n) (((x) >> (n)) | ((x) << (32 - (n))))
#else
// Kepler (Compute 3.5 / 5.0)
#define ROTR32(x, n) __funnelshift_r( (x), (x), (n) )
#endif
// in cpu-miner.c
extern bool opt_benchmark;
extern bool opt_debug;
extern int device_map[8];
extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);
// shared for 8 threads of addresses (cudaMalloc)
uint32_t* d_hash[8];
__constant__
static uint32_t pTarget[8];
__constant__
static uint32_t c_PaddedMessage80[32]; // padded message (80 bytes + padding)
static uint32_t *d_resNounce[8];
static uint32_t *h_resNounce[8];
__constant__
static uint8_t c_sigma[16][16];
const uint8_t host_sigma[16][16] =
{
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
{14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
{11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
{ 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
{ 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
{ 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 },
{12, 5, 1, 15, 14, 13, 4, 10, 0, 7, 6, 3, 9, 2, 8, 11 },
{13, 11, 7, 14, 12, 1, 3, 9, 5, 0, 15, 4, 8, 6, 2, 10 },
{ 6, 15, 14, 9, 11, 3, 0, 8, 12, 2, 13, 7, 1, 4, 10, 5 },
{10, 2, 8, 4, 7, 6, 1, 5, 15, 11, 9, 14, 3, 12, 13 , 0 },
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 },
{14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 },
{11, 8, 12, 0, 5, 2, 15, 13, 10, 14, 3, 6, 7, 1, 9, 4 },
{ 7, 9, 3, 1, 13, 12, 11, 14, 2, 6, 5, 10, 4, 0, 15, 8 },
{ 9, 0, 5, 7, 2, 4, 10, 15, 14, 1, 11, 12, 6, 8, 3, 13 },
{ 2, 12, 6, 10, 0, 11, 8, 3, 4, 13, 7, 5, 15, 14, 1, 9 }
};
__device__ __constant__
static const uint32_t c_IV256[8] = {
SPH_C32(0x6A09E667), SPH_C32(0xBB67AE85),
SPH_C32(0x3C6EF372), SPH_C32(0xA54FF53A),
SPH_C32(0x510E527F), SPH_C32(0x9B05688C),
SPH_C32(0x1F83D9AB), SPH_C32(0x5BE0CD19)
};
__device__ __constant__
static const uint32_t c_u256[16] = {
SPH_C32(0x243F6A88), SPH_C32(0x85A308D3),
SPH_C32(0x13198A2E), SPH_C32(0x03707344),
SPH_C32(0xA4093822), SPH_C32(0x299F31D0),
SPH_C32(0x082EFA98), SPH_C32(0xEC4E6C89),
SPH_C32(0x452821E6), SPH_C32(0x38D01377),
SPH_C32(0xBE5466CF), SPH_C32(0x34E90C6C),
SPH_C32(0xC0AC29B7), SPH_C32(0xC97C50DD),
SPH_C32(0x3F84D5B5), SPH_C32(0xB5470917)
};
#if 0
#define GS(m0, m1, c0, c1, a, b, c, d) do { \
a = SPH_T32(a + b + (m0 ^ c1)); \
d = SPH_ROTR32(d ^ a, 16); \
c = SPH_T32(c + d); \
b = SPH_ROTR32(b ^ c, 12); \
a = SPH_T32(a + b + (m1 ^ c0)); \
d = SPH_ROTR32(d ^ a, 8); \
c = SPH_T32(c + d); \
b = SPH_ROTR32(b ^ c, 7); \
} while (0)
#define ROUND_S(r) do { \
GS(Mx(r, 0x0), Mx(r, 0x1), CSx(r, 0x0), CSx(r, 0x1), v[0], v[4], v[0x8], v[0xC]); \
GS(Mx(r, 0x2), Mx(r, 0x3), CSx(r, 0x2), CSx(r, 0x3), v[1], v[5], v[0x9], v[0xD]); \
GS(Mx(r, 0x4), Mx(r, 0x5), CSx(r, 0x4), CSx(r, 0x5), v[2], v[6], v[0xA], v[0xE]); \
GS(Mx(r, 0x6), Mx(r, 0x7), CSx(r, 0x6), CSx(r, 0x7), v[3], v[7], v[0xB], v[0xF]); \
GS(Mx(r, 0x8), Mx(r, 0x9), CSx(r, 0x8), CSx(r, 0x9), v[0], v[5], v[0xA], v[0xF]); \
GS(Mx(r, 0xA), Mx(r, 0xB), CSx(r, 0xA), CSx(r, 0xB), v[1], v[6], v[0xB], v[0xC]); \
GS(Mx(r, 0xC), Mx(r, 0xD), CSx(r, 0xC), CSx(r, 0xD), v[2], v[7], v[0x8], v[0xD]); \
GS(Mx(r, 0xE), Mx(r, 0xF), CSx(r, 0xE), CSx(r, 0xF), v[3], v[4], v[0x9], v[0xE]); \
} while (0)
#endif
#define GS(a,b,c,d,e) { \
v[a] += (m[sigma[i][e]] ^ u256[sigma[i][e+1]]) + v[b]; \
v[d] = ROTR32(v[d] ^ v[a], 16); \
v[c] += v[d]; \
v[b] = ROTR32(v[b] ^ v[c], 12); \
\
v[a] += (m[sigma[i][e+1]] ^ u256[sigma[i][e]]) + v[b]; \
v[d] = ROTR32(v[d] ^ v[a], 8); \
v[c] += v[d]; \
v[b] = ROTR32(v[b] ^ v[c], 7); \
}
__device__ static
void blake256_compress(uint32_t *h, uint32_t *block, uint8_t ((*sigma)[16]), const uint32_t *u256, const uint32_t T0, uint8_t nullt = 1)
{
uint32_t /* __align__(8) */ v[16];
uint32_t /* __align__(8) */ m[16];
//#pragma unroll
for (int i = 0; i < 16; ++i) {
m[i] = cuda_swab32(block[i]);
//m[i] = block[i];
}
#pragma unroll
for(int i = 0; i < 8; i++)
v[i] = h[i];
v[ 8] = u256[0];
v[ 9] = u256[1];
v[10] = u256[2];
v[11] = u256[3];
v[12] = u256[4] ^ T0;
v[13] = u256[5] ^ T0;
v[14] = u256[6];
v[15] = u256[7];
// on a 80-bytes null buffer :
// first : v = {0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, ...}
// second : v = {0xb5bfb2f9, 0x14cfcc63, 0xb85c549c, 0xc9b4184e, ..., 0x299f3350, 0x082efa98, 0xec4e6c89}
//#pragma unroll
for (int i = 0; i < 14; i++) {
/* column step */
GS(0, 4, 0x8, 0xC, 0);
GS(1, 5, 0x9, 0xD, 2);
GS(2, 6, 0xA, 0xE, 4);
GS(3, 7, 0xB, 0xF, 6);
/* diagonal step */
GS(0, 5, 0xA, 0xF, 0x8);
GS(1, 6, 0xB, 0xC, 0xA);
GS(2, 7, 0x8, 0xD, 0xC);
GS(3, 4, 0x9, 0xE, 0xE);
}
//#pragma unroll 16
for(int i = 0; i < 16; i++)
h[i % 8] ^= v[i];
//second H0 = 0x0c7b1594 ... H7 = 0x9051b305
}
#if __CUDA_ARCH__ >= 200
#if (__NV_POINTER_SIZE == 64)
# define SZCT uint64_t
#else
# define SZCT uint32_t
#endif
extern __device__ __device_builtin__ void __nvvm_memset(uint8_t *, unsigned char, SZCT, int);
#endif
__global__
void blake256_gpu_hash_80(int threads, uint32_t startNounce, void *outputHash)
{
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
uint32_t /* __align__(16) */ h[8];
uint32_t /* __align__(16) */ msg[16];
const uint32_t nounce = startNounce + thread;
#pragma unroll
for(int i=0; i<8; i++)
h[i] = c_IV256[i];
blake256_compress(h, c_PaddedMessage80, c_sigma, c_u256, 0x200); /* 512 = 0x200 */
// ------ Close: Bytes 64 to 80 ------
#if 0 /* __CUDA_ARCH__ >= 200 */
__nvvm_memset((uint8_t*)(&msg[4]), 0, sizeof(msg)-16, 16);
#else
msg[5] = 0;
msg[6] = 0;
msg[7] = 0;
msg[8] = 0;
msg[9] = 0;
msg[10] = 0;
msg[11] = 0;
msg[12] = 0;
msg[14] = 0;
#endif
msg[0] = c_PaddedMessage80[16];
msg[1] = c_PaddedMessage80[17];
msg[2] = c_PaddedMessage80[18];
msg[3] = cuda_swab32(nounce); // here or at 80 ?
msg[4] = 0x80; // uchar[16] after buffer
msg[13] = 0x01000000; //((uint8_t*)msg)[55] = 1; // uchar[17 to 55]
msg[15] = 0x80020000; // 60-63 0x280
//h => {0xb5bfb2f9, 0x14cfcc63, 0xb85c549c, 0xc9b4184e, 0x67dfc6ce, 0x29e9904b, 0xd59ee74e, 0xfaa9c653}
//msg {0, 0, 0, 0, 0x80, 0...}
blake256_compress(h, msg, c_sigma, c_u256, 0x280); // or 0x80
//h => {0x0c7b1594, 0x52328517, 0x463db487, 0xdf5e39b7, 0x1322afaf, 0x14ed562c, 0xe9d18d7d, 0x9051b305}
uint32_t *outHash = (uint32_t*) outputHash + 16*thread; // 16 = 4 x sizeof(uint32)
//#pragma unroll
for (int i=0; i < 8; i++) {
outHash[i] = cuda_swab32(h[i]);
}
}
}
__host__
void blake256_cpu_hash_80(int thr_id, int threads, uint32_t startNounce, uint32_t *d_outputHash, int order)
{
const int threadsperblock = 256;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
size_t shared_size = 0;
blake256_gpu_hash_80<<<grid, block, shared_size>>>(threads, startNounce, d_outputHash);
MyStreamSynchronize(NULL, order, thr_id);
}
__global__
void gpu_check_hash_64(int threads, uint32_t startNounce, uint32_t *g_nonceVector, uint32_t *g_hash, uint32_t *resNounce)
{
int thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
uint32_t nounce = (g_nonceVector != NULL) ? g_nonceVector[thread] : (startNounce + thread);
int hashPosition = nounce - startNounce;
uint32_t *inpHash = &g_hash[16 * hashPosition];
uint32_t hash[8];
#pragma unroll 8
for (int i=0; i < 8; i++)
hash[i] = inpHash[i];
int i, position = -1;
bool rc = true;
#pragma unroll 8
for (i = 7; i >= 0; i--) {
if (hash[i] > pTarget[i] && position < i) {
position = i;
rc = false;
}
if (hash[i] < pTarget[i] && position < i) {
position = i;
rc = true;
}
}
if(rc && resNounce[0] > nounce)
resNounce[0] = nounce;
}
}
__host__
uint32_t cpu_check_hash_64(int thr_id, int threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order)
{
uint32_t result = 0xffffffff;
const int threadsperblock = 256;
cudaMemset(d_resNounce[thr_id], 0xff, sizeof(uint32_t));
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
size_t shared_size = 0;
gpu_check_hash_64 <<<grid, block, shared_size>>>(threads, startNounce, d_nonceVector, d_inputHash, d_resNounce[thr_id]);
MyStreamSynchronize(NULL, order, thr_id);
CUDA_SAFE_CALL(cudaMemcpy(h_resNounce[thr_id], d_resNounce[thr_id], sizeof(uint32_t), cudaMemcpyDeviceToHost));
// cudaMemcpy() is asynch!
cudaThreadSynchronize();
result = *h_resNounce[thr_id];
return result;
}
__host__
void blake256_cpu_init(int thr_id)
{
CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_sigma, host_sigma, sizeof(host_sigma), 0, cudaMemcpyHostToDevice));
CUDA_SAFE_CALL(cudaMallocHost(&h_resNounce[thr_id], 1*sizeof(uint32_t)));
CUDA_SAFE_CALL(cudaMalloc(&d_resNounce[thr_id], 1*sizeof(uint32_t)));
}
__host__
void blake256_cpu_setBlock_80(uint32_t *pdata, const void *ptarget)
{
uint32_t PaddedMessage[32];
memcpy(PaddedMessage, pdata, 80);
memset(&PaddedMessage[20], 0, 48);
//for (int i=0; i<20; i++)
// PaddedMessage[i] = cuda_swab32(pdata[i]);
CUDA_SAFE_CALL(cudaMemcpyToSymbol(pTarget, ptarget, 32, 0, cudaMemcpyHostToDevice));
CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_PaddedMessage80, PaddedMessage, sizeof(PaddedMessage), 0, cudaMemcpyHostToDevice));
}
#define NULLTEST 0
extern "C" int scanhash_blake32(int thr_id, uint32_t *pdata, const uint32_t *ptarget,
uint32_t max_nonce, unsigned long *hashes_done)
{
uint32_t endiandata[20];
const uint32_t first_nonce = pdata[19];
const int throughput = 256*256*2;
static bool init[8] = {0,0,0,0,0,0,0,0};
if (opt_benchmark)
((uint32_t*)ptarget)[7] = 0x00000f;
uint32_t Htarg = ptarget[7];
if (!init[thr_id]) {
CUDA_SAFE_CALL(cudaSetDevice(device_map[thr_id]));
CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput));
blake256_cpu_init(thr_id);
init[thr_id] = true;
}
#if NULLTEST
// dev test with a null buffer 0x00000...
for (int k = 0; k < 20; k++)
pdata[k] = 0;
uint32_t vhash[8];
blake32hash(vhash, pdata);
#endif
for (int k=0; k < 20; k++)
be32enc(&endiandata[k], pdata[k]);
blake256_cpu_setBlock_80(endiandata, (void*)ptarget);
do {
int order = 0;
uint32_t foundNonce;
// GPU
blake256_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
#if NULLTEST
uint32_t buf[8]; memset(buf, 0, sizeof buf);
CUDA_SAFE_CALL(cudaMemcpy(buf, d_hash[thr_id], sizeof buf, cudaMemcpyDeviceToHost));
CUDA_SAFE_CALL(cudaThreadSynchronize());
//applog_hash((unsigned char*)buf);
#endif
foundNonce = cpu_check_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
if (foundNonce != 0xffffffff)
{
uint32_t vhashcpu[8];
be32enc(&endiandata[19], foundNonce);
blake32hash(vhashcpu, endiandata);
if (opt_debug)
applog(LOG_DEBUG, "foundNonce = %08x",foundNonce);
if (vhashcpu[7] <= Htarg && fulltest(vhashcpu, ptarget))
{
pdata[19] = foundNonce;
*hashes_done = pdata[19] - first_nonce + 1;
return 1;
} else {
applog(LOG_INFO, "GPU #%d: result for nonce %08x does not validate on CPU!", thr_id, foundNonce);
}
}
pdata[19] += throughput;
} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
*hashes_done = pdata[19] - first_nonce + 1;
return 0;
}
//#define DEBUG_ALGO
__host__
int scanhash_blake256_cpu(int thr_id, uint32_t *pdata, const uint32_t *ptarget,
uint32_t max_nonce, uint64_t *hashes_done)
{
uint32_t n = pdata[19] - 1;
const uint32_t first_nonce = pdata[19];
const uint32_t Htarg = ptarget[7];
uint32_t __align__(32) hash64[8];
uint32_t endiandata[32];
uint64_t htmax[] = {
0,
0xF,
0xFF,
0xFFF,
0xFFFF,
0x10000000
};
uint32_t masks[] = {
0xFFFFFFFF,
0xFFFFFFF0,
0xFFFFFF00,
0xFFFFF000,
0xFFFF0000,
0
};
// we need bigendian data...
for (int kk=0; kk < 32; kk++) {
be32enc(&endiandata[kk], ((uint32_t*)pdata)[kk]);
};
#ifdef DEBUG_ALGO
if (Htarg != 0)
printf("[%d] Htarg=%X\n", thr_id, Htarg);
#endif
for (int m=0; m < 6; m++) {
if (Htarg <= htmax[m]) {
uint32_t mask = masks[m];
do {
pdata[19] = ++n;
be32enc(&endiandata[19], n);
blake32hash(hash64, endiandata);
#ifndef DEBUG_ALGO
if ((!(hash64[7] & mask)) && fulltest(hash64, ptarget)) {
*hashes_done = n - first_nonce + 1;
return true;
}
#else
if (!(n % 0x1000) && !thr_id) printf(".");
if (!(hash64[7] & mask)) {
printf("[%d]",thr_id);
if (fulltest(hash64, ptarget)) {
*hashes_done = n - first_nonce + 1;
return true;
}
}
#endif
} while (n < max_nonce && !work_restart[thr_id].restart);
// see blake.c if else to understand the loop on htmax => mask
break;
}
}
*hashes_done = n - first_nonce + 1;
pdata[19] = n;
return 0;
}

2
build.sh

@ -4,7 +4,7 @@
# export PATH="$PATH:/usr/local/cuda/bin/" # export PATH="$PATH:/usr/local/cuda/bin/"
make distclean || echo clean #make distclean || echo clean
rm -f Makefile.in rm -f Makefile.in
rm -f config.status rm -f config.status

4
ccminer.vcxproj

@ -397,6 +397,10 @@ copy "$(CudaToolkitBinDir)\cudart64*.dll" "$(OutDir)"</Command>
<AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions> <AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions>
<TargetMachinePlatform Condition="'$(Platform)'=='x64'">64</TargetMachinePlatform> <TargetMachinePlatform Condition="'$(Platform)'=='x64'">64</TargetMachinePlatform>
</CudaCompile> </CudaCompile>
<CudaCompile Include="blake32.cu">
<AdditionalOptions Condition="'$(Configuration)'=='Release'">--ptxas-options=-O2 %(AdditionalOptions)</AdditionalOptions>
<AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions>
</CudaCompile>
<CudaCompile Include="quark\animecoin.cu"> <CudaCompile Include="quark\animecoin.cu">
<AdditionalOptions Condition="'$(Configuration)'=='Release'">--ptxas-options=-O2 %(AdditionalOptions)</AdditionalOptions> <AdditionalOptions Condition="'$(Configuration)'=='Release'">--ptxas-options=-O2 %(AdditionalOptions)</AdditionalOptions>
<AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions> <AdditionalOptions Condition="'$(Configuration)'=='Debug'">%(AdditionalOptions)</AdditionalOptions>

3
ccminer.vcxproj.filters

@ -436,5 +436,8 @@
<CudaCompile Include="x17\x17.cu"> <CudaCompile Include="x17\x17.cu">
<Filter>Source Files\CUDA\x17</Filter> <Filter>Source Files\CUDA\x17</Filter>
</CudaCompile> </CudaCompile>
<CudaCompile Include="blake32.cu">
<Filter>Source Files\CUDA</Filter>
</CudaCompile>
</ItemGroup> </ItemGroup>
</Project> </Project>

2
configure.sh

@ -1 +1 @@
./configure "CFLAGS=-O3" "CXXFLAGS=-O3" --with-cuda=/usr/local/cuda ./configure "CFLAGS=-O2" "CXXFLAGS=-O2" --with-cuda=/usr/local/cuda

48
cpu-miner.c

@ -135,6 +135,7 @@ typedef enum {
ALGO_QUARK, ALGO_QUARK,
ALGO_ANIME, ALGO_ANIME,
ALGO_FRESH, ALGO_FRESH,
ALGO_BLAKE,
ALGO_NIST5, ALGO_NIST5,
ALGO_WHC, ALGO_WHC,
ALGO_X11, ALGO_X11,
@ -155,6 +156,7 @@ static const char *algo_names[] = {
"quark", "quark",
"anime", "anime",
"fresh", "fresh",
"blake",
"nist5", "nist5",
"whirl", "whirl",
"x11", "x11",
@ -235,6 +237,7 @@ Options:\n\
jackpot Jackpot hash\n\ jackpot Jackpot hash\n\
quark Quark hash\n\ quark Quark hash\n\
anime Animecoin hash\n\ anime Animecoin hash\n\
blake Blake 256 (like NEOS blake)\n\
fresh Freshcoin hash (shavite 80)\n\ fresh Freshcoin hash (shavite 80)\n\
nist5 NIST5 (TalkCoin) hash\n\ nist5 NIST5 (TalkCoin) hash\n\
whirl Whirlcoin (old whirlpool)\n\ whirl Whirlcoin (old whirlpool)\n\
@ -842,18 +845,23 @@ static void *miner_thread(void *userdata)
int64_t max64; int64_t max64;
int rc; int rc;
// &work.data[19]
int wcmplen = 76;
uint32_t *nonceptr = (uint32_t*) (((char*)work.data) + wcmplen);
if (have_stratum) { if (have_stratum) {
while (time(NULL) >= g_work_time + 120) while (time(NULL) >= g_work_time + 120)
sleep(1); sleep(1);
pthread_mutex_lock(&g_work_lock); pthread_mutex_lock(&g_work_lock);
if (work.data[19] >= end_nonce) if ((*nonceptr) >= end_nonce)
stratum_gen_work(&stratum, &g_work); stratum_gen_work(&stratum, &g_work);
} else { } else {
int min_scantime = have_longpoll ? LP_SCANTIME : opt_scantime;
/* obtain new work from internal workio thread */ /* obtain new work from internal workio thread */
pthread_mutex_lock(&g_work_lock); pthread_mutex_lock(&g_work_lock);
if (!have_stratum && (!have_longpoll || if (!have_stratum &&
time(NULL) >= g_work_time + LP_SCANTIME*3/4 || (time(NULL) - g_work_time >= min_scantime ||
work.data[19] >= end_nonce)) { (*nonceptr) >= end_nonce)) {
if (unlikely(!get_work(mythr, &g_work))) { if (unlikely(!get_work(mythr, &g_work))) {
applog(LOG_ERR, "work retrieval failed, exiting " applog(LOG_ERR, "work retrieval failed, exiting "
"mining thread %d", mythr->id); "mining thread %d", mythr->id);
@ -867,11 +875,11 @@ static void *miner_thread(void *userdata)
continue; continue;
} }
} }
if (memcmp(work.data, g_work.data, 76)) { if (memcmp(work.data, g_work.data, wcmplen)) {
memcpy(&work, &g_work, sizeof(struct work)); memcpy(&work, &g_work, sizeof(struct work));
work.data[19] = 0xffffffffU / opt_n_threads * thr_id; (*nonceptr) = 0xffffffffU / opt_n_threads * thr_id;
} else } else
work.data[19]++; (*nonceptr)++;
pthread_mutex_unlock(&g_work_lock); pthread_mutex_unlock(&g_work_lock);
work_restart[thr_id].restart = 0; work_restart[thr_id].restart = 0;
@ -881,13 +889,26 @@ static void *miner_thread(void *userdata)
else else
max64 = g_work_time + (have_longpoll ? LP_SCANTIME : opt_scantime) max64 = g_work_time + (have_longpoll ? LP_SCANTIME : opt_scantime)
- time(NULL); - time(NULL);
max64 *= (int64_t)thr_hashrates[thr_id]; max64 *= (int64_t)thr_hashrates[thr_id];
if (max64 <= 0)
max64 = (opt_algo == ALGO_JACKPOT) ? 0x1fffLL : 0xfffffLL; if (max64 <= 0) {
if ((int64_t)work.data[19] + max64 > end_nonce) switch (opt_algo) {
case ALGO_JACKPOT:
max64 = 0x1fffLL;
break;
case ALGO_BLAKE:
max64 = 0xffffffLL;
break;
default:
max64 = 0xfffffLL;
break;
}
}
if ((int64_t)(*nonceptr) + max64 > end_nonce)
max_nonce = end_nonce; max_nonce = end_nonce;
else else
max_nonce = (uint32_t)(work.data[19] + max64); max_nonce = (uint32_t)((*nonceptr) + max64);
hashes_done = 0; hashes_done = 0;
gettimeofday(&tv_start, NULL); gettimeofday(&tv_start, NULL);
@ -931,6 +952,11 @@ static void *miner_thread(void *userdata)
max_nonce, &hashes_done); max_nonce, &hashes_done);
break; break;
case ALGO_BLAKE:
rc = scanhash_blake32(thr_id, work.data, work.target,
max_nonce, &hashes_done);
break;
case ALGO_ANIME: case ALGO_ANIME:
rc = scanhash_anime(thr_id, work.data, work.target, rc = scanhash_anime(thr_id, work.data, work.target,
max_nonce, &hashes_done); max_nonce, &hashes_done);

5
miner.h

@ -241,6 +241,10 @@ extern int scanhash_fresh(int thr_id, uint32_t *pdata,
const uint32_t *ptarget, uint32_t max_nonce, const uint32_t *ptarget, uint32_t max_nonce,
unsigned long *hashes_done); unsigned long *hashes_done);
extern int scanhash_blake32(int thr_id, uint32_t *pdata,
const uint32_t *ptarget, uint32_t max_nonce,
unsigned long *hashes_done);
extern int scanhash_nist5(int thr_id, uint32_t *pdata, extern int scanhash_nist5(int thr_id, uint32_t *pdata,
const uint32_t *ptarget, uint32_t max_nonce, const uint32_t *ptarget, uint32_t max_nonce,
unsigned long *hashes_done); unsigned long *hashes_done);
@ -402,6 +406,7 @@ void heavycoin_hash(unsigned char* output, const unsigned char* input, int len);
void groestlhash(void *state, const void *input); void groestlhash(void *state, const void *input);
void myriadhash(void *state, const void *input); void myriadhash(void *state, const void *input);
void fresh_hash(void *state, const void *input); void fresh_hash(void *state, const void *input);
void blake32hash(void *output, const void *input);
void nist5hash(void *state, const void *input); void nist5hash(void *state, const void *input);
void quarkhash(void *state, const void *input); void quarkhash(void *state, const void *input);
void wcoinhash(void *state, const void *input); void wcoinhash(void *state, const void *input);

367
missing

@ -1,367 +0,0 @@
#! /bin/sh
# Common stub for a few missing GNU programs while installing.
scriptversion=2006-05-10.23
# Copyright (C) 1996, 1997, 1999, 2000, 2002, 2003, 2004, 2005, 2006
# Free Software Foundation, Inc.
# Originally by Fran,cois Pinard <pinard@iro.umontreal.ca>, 1996.
# This program is free software; you can redistribute it and/or modify
# it under the terms of the GNU General Public License as published by
# the Free Software Foundation; either version 2, or (at your option)
# any later version.
# This program is distributed in the hope that it will be useful,
# but WITHOUT ANY WARRANTY; without even the implied warranty of
# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
# GNU General Public License for more details.
# You should have received a copy of the GNU General Public License
# along with this program; if not, write to the Free Software
# Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
# 02110-1301, USA.
# As a special exception to the GNU General Public License, if you
# distribute this file as part of a program that contains a
# configuration script generated by Autoconf, you may include it under
# the same distribution terms that you use for the rest of that program.
if test $# -eq 0; then
echo 1>&2 "Try \`$0 --help' for more information"
exit 1
fi
run=:
sed_output='s/.* --output[ =]\([^ ]*\).*/\1/p'
sed_minuso='s/.* -o \([^ ]*\).*/\1/p'
# In the cases where this matters, `missing' is being run in the
# srcdir already.
if test -f configure.ac; then
configure_ac=configure.ac
else
configure_ac=configure.in
fi
msg="missing on your system"
case $1 in
--run)
# Try to run requested program, and just exit if it succeeds.
run=
shift
"$@" && exit 0
# Exit code 63 means version mismatch. This often happens
# when the user try to use an ancient version of a tool on
# a file that requires a minimum version. In this case we
# we should proceed has if the program had been absent, or
# if --run hadn't been passed.
if test $? = 63; then
run=:
msg="probably too old"
fi
;;
-h|--h|--he|--hel|--help)
echo "\
$0 [OPTION]... PROGRAM [ARGUMENT]...
Handle \`PROGRAM [ARGUMENT]...' for when PROGRAM is missing, or return an
error status if there is no known handling for PROGRAM.
Options:
-h, --help display this help and exit
-v, --version output version information and exit
--run try to run the given command, and emulate it if it fails
Supported PROGRAM values:
aclocal touch file \`aclocal.m4'
autoconf touch file \`configure'
autoheader touch file \`config.h.in'
autom4te touch the output file, or create a stub one
automake touch all \`Makefile.in' files
bison create \`y.tab.[ch]', if possible, from existing .[ch]
flex create \`lex.yy.c', if possible, from existing .c
help2man touch the output file
lex create \`lex.yy.c', if possible, from existing .c
makeinfo touch the output file
tar try tar, gnutar, gtar, then tar without non-portable flags
yacc create \`y.tab.[ch]', if possible, from existing .[ch]
Send bug reports to <bug-automake@gnu.org>."
exit $?
;;
-v|--v|--ve|--ver|--vers|--versi|--versio|--version)
echo "missing $scriptversion (GNU Automake)"
exit $?
;;
-*)
echo 1>&2 "$0: Unknown \`$1' option"
echo 1>&2 "Try \`$0 --help' for more information"
exit 1
;;
esac
# Now exit if we have it, but it failed. Also exit now if we
# don't have it and --version was passed (most likely to detect
# the program).
case $1 in
lex|yacc)
# Not GNU programs, they don't have --version.
;;
tar)
if test -n "$run"; then
echo 1>&2 "ERROR: \`tar' requires --run"
exit 1
elif test "x$2" = "x--version" || test "x$2" = "x--help"; then
exit 1
fi
;;
*)
if test -z "$run" && ($1 --version) > /dev/null 2>&1; then
# We have it, but it failed.
exit 1
elif test "x$2" = "x--version" || test "x$2" = "x--help"; then
# Could not run --version or --help. This is probably someone
# running `$TOOL --version' or `$TOOL --help' to check whether
# $TOOL exists and not knowing $TOOL uses missing.
exit 1
fi
;;
esac
# If it does not exist, or fails to run (possibly an outdated version),
# try to emulate it.
case $1 in
aclocal*)
echo 1>&2 "\
WARNING: \`$1' is $msg. You should only need it if
you modified \`acinclude.m4' or \`${configure_ac}'. You might want
to install the \`Automake' and \`Perl' packages. Grab them from
any GNU archive site."
touch aclocal.m4
;;
autoconf)
echo 1>&2 "\
WARNING: \`$1' is $msg. You should only need it if
you modified \`${configure_ac}'. You might want to install the
\`Autoconf' and \`GNU m4' packages. Grab them from any GNU
archive site."
touch configure
;;
autoheader)
echo 1>&2 "\
WARNING: \`$1' is $msg. You should only need it if
you modified \`acconfig.h' or \`${configure_ac}'. You might want
to install the \`Autoconf' and \`GNU m4' packages. Grab them
from any GNU archive site."
files=`sed -n 's/^[ ]*A[CM]_CONFIG_HEADER(\([^)]*\)).*/\1/p' ${configure_ac}`
test -z "$files" && files="config.h"
touch_files=
for f in $files; do
case $f in
*:*) touch_files="$touch_files "`echo "$f" |
sed -e 's/^[^:]*://' -e 's/:.*//'`;;
*) touch_files="$touch_files $f.in";;
esac
done
touch $touch_files
;;
automake*)
echo 1>&2 "\
WARNING: \`$1' is $msg. You should only need it if
you modified \`Makefile.am', \`acinclude.m4' or \`${configure_ac}'.
You might want to install the \`Automake' and \`Perl' packages.
Grab them from any GNU archive site."
find . -type f -name Makefile.am -print |
sed 's/\.am$/.in/' |
while read f; do touch "$f"; done
;;
autom4te)
echo 1>&2 "\
WARNING: \`$1' is needed, but is $msg.
You might have modified some files without having the
proper tools for further handling them.
You can get \`$1' as part of \`Autoconf' from any GNU
archive site."
file=`echo "$*" | sed -n "$sed_output"`
test -z "$file" && file=`echo "$*" | sed -n "$sed_minuso"`
if test -f "$file"; then
touch $file
else
test -z "$file" || exec >$file
echo "#! /bin/sh"
echo "# Created by GNU Automake missing as a replacement of"
echo "# $ $@"
echo "exit 0"
chmod +x $file
exit 1
fi
;;
bison|yacc)
echo 1>&2 "\
WARNING: \`$1' $msg. You should only need it if
you modified a \`.y' file. You may need the \`Bison' package
in order for those modifications to take effect. You can get
\`Bison' from any GNU archive site."
rm -f y.tab.c y.tab.h
if test $# -ne 1; then
eval LASTARG="\${$#}"
case $LASTARG in
*.y)
SRCFILE=`echo "$LASTARG" | sed 's/y$/c/'`
if test -f "$SRCFILE"; then
cp "$SRCFILE" y.tab.c
fi
SRCFILE=`echo "$LASTARG" | sed 's/y$/h/'`
if test -f "$SRCFILE"; then
cp "$SRCFILE" y.tab.h
fi
;;
esac
fi
if test ! -f y.tab.h; then
echo >y.tab.h
fi
if test ! -f y.tab.c; then
echo 'main() { return 0; }' >y.tab.c
fi
;;
lex|flex)
echo 1>&2 "\
WARNING: \`$1' is $msg. You should only need it if
you modified a \`.l' file. You may need the \`Flex' package
in order for those modifications to take effect. You can get
\`Flex' from any GNU archive site."
rm -f lex.yy.c
if test $# -ne 1; then
eval LASTARG="\${$#}"
case $LASTARG in
*.l)
SRCFILE=`echo "$LASTARG" | sed 's/l$/c/'`
if test -f "$SRCFILE"; then
cp "$SRCFILE" lex.yy.c
fi
;;
esac
fi
if test ! -f lex.yy.c; then
echo 'main() { return 0; }' >lex.yy.c
fi
;;
help2man)
echo 1>&2 "\
WARNING: \`$1' is $msg. You should only need it if
you modified a dependency of a manual page. You may need the
\`Help2man' package in order for those modifications to take
effect. You can get \`Help2man' from any GNU archive site."
file=`echo "$*" | sed -n "$sed_output"`
test -z "$file" && file=`echo "$*" | sed -n "$sed_minuso"`
if test -f "$file"; then
touch $file
else
test -z "$file" || exec >$file
echo ".ab help2man is required to generate this page"
exit 1
fi
;;
makeinfo)
echo 1>&2 "\
WARNING: \`$1' is $msg. You should only need it if
you modified a \`.texi' or \`.texinfo' file, or any other file
indirectly affecting the aspect of the manual. The spurious
call might also be the consequence of using a buggy \`make' (AIX,
DU, IRIX). You might want to install the \`Texinfo' package or
the \`GNU make' package. Grab either from any GNU archive site."
# The file to touch is that specified with -o ...
file=`echo "$*" | sed -n "$sed_output"`
test -z "$file" && file=`echo "$*" | sed -n "$sed_minuso"`
if test -z "$file"; then
# ... or it is the one specified with @setfilename ...
infile=`echo "$*" | sed 's/.* \([^ ]*\) *$/\1/'`
file=`sed -n '
/^@setfilename/{
s/.* \([^ ]*\) *$/\1/
p
q
}' $infile`
# ... or it is derived from the source name (dir/f.texi becomes f.info)
test -z "$file" && file=`echo "$infile" | sed 's,.*/,,;s,.[^.]*$,,'`.info
fi
# If the file does not exist, the user really needs makeinfo;
# let's fail without touching anything.
test -f $file || exit 1
touch $file
;;
tar)
shift
# We have already tried tar in the generic part.
# Look for gnutar/gtar before invocation to avoid ugly error
# messages.
if (gnutar --version > /dev/null 2>&1); then
gnutar "$@" && exit 0
fi
if (gtar --version > /dev/null 2>&1); then
gtar "$@" && exit 0
fi
firstarg="$1"
if shift; then
case $firstarg in
*o*)
firstarg=`echo "$firstarg" | sed s/o//`
tar "$firstarg" "$@" && exit 0
;;
esac
case $firstarg in
*h*)
firstarg=`echo "$firstarg" | sed s/h//`
tar "$firstarg" "$@" && exit 0
;;
esac
fi
echo 1>&2 "\
WARNING: I can't seem to be able to run \`tar' with the given arguments.
You may want to install GNU tar or Free paxutils, or check the
command line arguments."
exit 1
;;
*)
echo 1>&2 "\
WARNING: \`$1' is needed, and is $msg.
You might have modified some files without having the
proper tools for further handling them. Check the \`README' file,
it often tells you about the needed prerequisites for installing
this package. You may also peek at any GNU archive site, in case
some other package would contain this missing \`$1' program."
exit 1
;;
esac
exit 0
# Local variables:
# eval: (add-hook 'write-file-hooks 'time-stamp)
# time-stamp-start: "scriptversion="
# time-stamp-format: "%:y-%02m-%02d.%02H"
# time-stamp-end: "$"
# End:

4
quark/cuda_checkhash.cu

@ -6,8 +6,8 @@
// Hash Target gegen das wir testen sollen // Hash Target gegen das wir testen sollen
__constant__ uint32_t pTarget[8]; __constant__ uint32_t pTarget[8];
uint32_t *d_resNounce[8]; static uint32_t *d_resNounce[8];
uint32_t *h_resNounce[8]; static uint32_t *h_resNounce[8];
// aus heavy.cu // aus heavy.cu
extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id); extern cudaError_t MyStreamSynchronize(cudaStream_t stream, int situation, int thr_id);

4
util.c

@ -1393,6 +1393,10 @@ void print_hash_tests(void)
myriadhash(&hash[0], &buf[0]); myriadhash(&hash[0], &buf[0]);
printpfx("myriad", hash); printpfx("myriad", hash);
memset(hash, 0, sizeof hash);
blake32hash(&hash[0], &buf[0]);
printpfx("blake", hash);
memset(hash, 0, sizeof hash); memset(hash, 0, sizeof hash);
nist5hash(&hash[0], &buf[0]); nist5hash(&hash[0], &buf[0]);
printpfx("nist5", hash); printpfx("nist5", hash);

Loading…
Cancel
Save