Skip to content

Commit

Permalink
Issue #147
Browse files Browse the repository at this point in the history
Linux seems working. Needs to verify windows.
  • Loading branch information
fancyIX committed Jan 9, 2019
1 parent 2c6c071 commit 3054dda
Show file tree
Hide file tree
Showing 16 changed files with 740 additions and 16 deletions.
1 change: 1 addition & 0 deletions Makefile.am
Original file line number Diff line number Diff line change
Expand Up @@ -77,6 +77,7 @@ sgminer_SOURCES += algorithm/lyra2re.c algorithm/lyra2re.h algorithm/lyra2.c alg
sgminer_SOURCES += algorithm/lyra2rev2.c algorithm/lyra2rev2.h
sgminer_SOURCES += algorithm/lyra2rev3.c algorithm/lyra2rev3.h
sgminer_SOURCES += algorithm/lyra2Z.c algorithm/lyra2Z.h
sgminer_SOURCES += algorithm/lyra2Zz.c algorithm/lyra2Zz.h
sgminer_SOURCES += algorithm/lyra2h.c algorithm/lyra2h.h
sgminer_SOURCES += algorithm/pluck.c algorithm/pluck.h
sgminer_SOURCES += algorithm/sia.c algorithm/sia.h
Expand Down
60 changes: 60 additions & 0 deletions algorithm.c
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,7 @@
#include "algorithm/lyra2rev2.h"
#include "algorithm/lyra2rev3.h"
#include "algorithm/lyra2Z.h"
#include "algorithm/lyra2Zz.h"
#include "algorithm/pluck.h"
#include "algorithm/yescrypt.h"
#include "algorithm/credits.h"
Expand Down Expand Up @@ -83,6 +84,7 @@ const char *algorithm_type_str[] = {
"Lyra2REV2",
"Lyra2REV3",
"Lyra2Z",
"Lyra2Zz",
"Lyra2h",
"Pluck",
"Yescrypt",
Expand Down Expand Up @@ -1370,6 +1372,63 @@ static cl_int queue_lyra2z_kernel(struct __clState *clState, struct _dev_blk_ctx
return status;
}

static cl_int queue_lyra2zz_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
cl_kernel *kernel;
unsigned int num;
cl_int status = 0;
cl_ulong le_target;

// le_target = *(cl_uint *)(blk->work->device_target + 28);
le_target = *(cl_ulong *)(blk->work->device_target + 24);
flip112(clState->cldata, blk->work->data);
status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 112, clState->cldata, 0, NULL, NULL);

// blake 112 - search
kernel = &clState->kernel;
num = 0;
// CL_SET_ARG(clState->CLbuffer0);
CL_SET_ARG(clState->padbuffer8);
CL_SET_ARG(blk->work->blk.ctx_a);
CL_SET_ARG(blk->work->blk.ctx_b);
CL_SET_ARG(blk->work->blk.ctx_c);
CL_SET_ARG(blk->work->blk.ctx_d);
CL_SET_ARG(blk->work->blk.ctx_e);
CL_SET_ARG(blk->work->blk.ctx_f);
CL_SET_ARG(blk->work->blk.ctx_g);
CL_SET_ARG(blk->work->blk.ctx_h);
CL_SET_ARG(blk->work->blk.cty_a);
CL_SET_ARG(blk->work->blk.cty_b);
CL_SET_ARG(blk->work->blk.cty_c);
CL_SET_ARG(blk->work->blk.cty_d);
CL_SET_ARG(blk->work->blk.cty_e);
CL_SET_ARG(blk->work->blk.cty_f);
CL_SET_ARG(blk->work->blk.cty_g);
CL_SET_ARG(blk->work->blk.cty_h);
CL_SET_ARG(blk->work->blk.cty_i);
CL_SET_ARG(blk->work->blk.cty_j);
CL_SET_ARG(blk->work->blk.cty_k);

kernel = clState->extra_kernels;
// lyra2_cuda_hash_64 - search1 2 3
num = 0;
CL_SET_ARG(clState->padbuffer8);
CL_SET_ARG(clState->buffer1);
//CL_SET_ARG(clState->buffer2);
num = 0;
CL_NEXTKERNEL_SET_ARG(clState->buffer1);
num = 0;
CL_NEXTKERNEL_SET_ARG(clState->padbuffer8);
CL_SET_ARG(clState->buffer1);
num = 0;
//output
CL_NEXTKERNEL_SET_ARG(clState->padbuffer8);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(le_target);

return status;
}

static cl_int queue_lyra2h_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
cl_kernel *kernel;
Expand Down Expand Up @@ -1622,6 +1681,7 @@ static algorithm_settings_t algos[] = {
{ "lyra2rev2", ALGO_LYRA2REV2, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 8, -1, 0, lyra2rev2_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2rev2_kernel, gen_hash, append_neoscrypt_compiler_options },
{ "lyra2rev3", ALGO_LYRA2REV3, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 8, -1, 0, lyra2rev3_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2rev3_kernel, gen_hash, append_neoscrypt_compiler_options },
{ "lyra2Z" , ALGO_LYRA2Z , "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, -1, 0, lyra2Z_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2z_kernel, gen_hash, NULL },
{ "lyra2Zz" , ALGO_LYRA2ZZ , "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, -1, 0, lyra2Zz_regenhash, blake256_midstate_112, blake256_prepare_work_112, queue_lyra2zz_kernel, gen_hash, NULL },
{ "lyra2h" , ALGO_LYRA2H , "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 1, -1, 0, lyra2h_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2h_kernel, gen_hash, NULL },
{ "allium", ALGO_ALLIUM, "", 1, 128, 128, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 2 * 8 * 4194304, 0, allium_regenhash, blake256_midstate, blake256_prepare_work, queue_allium_kernel, gen_hash, NULL },

Expand Down
1 change: 1 addition & 0 deletions algorithm.h
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@ typedef enum {
ALGO_LYRA2REV2,
ALGO_LYRA2REV3,
ALGO_LYRA2Z,
ALGO_LYRA2ZZ,
ALGO_LYRA2H,
ALGO_PLUCK,
ALGO_YESCRYPT,
Expand Down
43 changes: 43 additions & 0 deletions algorithm/blake256.c
Original file line number Diff line number Diff line change
Expand Up @@ -84,6 +84,49 @@ void blake256_prepare_work(dev_blk_ctx *blk, uint32_t *state, uint32_t *pdata)
blk->cty_c = pdata[18];
}

void blake256_midstate_112(struct work *work)
{
sph_blake256_context ctx_blake;
uint32_t data[16];

be32enc_vect(data, (const uint32_t *)work->data, 16);

sph_blake256_init(&ctx_blake);
sph_blake256(&ctx_blake, (unsigned char *)data, 64);

memcpy(work->midstate, ctx_blake.H, 32);
endian_flip32(work->midstate, work->midstate);

char *strdata, *strmidstate;
strdata = bin2hex(work->data, 112);
strmidstate = bin2hex(work->midstate, 32);
applog(LOG_DEBUG, "data %s midstate %s", strdata, strmidstate);
}

void blake256_prepare_work_112(dev_blk_ctx *blk, uint32_t *state, uint32_t *pdata)
{
blk->ctx_a = state[0];
blk->ctx_b = state[1];
blk->ctx_c = state[2];
blk->ctx_d = state[3];
blk->ctx_e = state[4];
blk->ctx_f = state[5];
blk->ctx_g = state[6];
blk->ctx_h = state[7];

blk->cty_a = pdata[16];
blk->cty_b = pdata[17];
blk->cty_c = pdata[18];
blk->cty_d = pdata[20];
blk->cty_e = pdata[21];
blk->cty_f = pdata[22];
blk->cty_g = pdata[23];
blk->cty_h = pdata[24];
blk->cty_i = pdata[25];
blk->cty_j = pdata[26];
blk->cty_k = pdata[27];
}

static const uint32_t diff1targ = 0x0000ffff;

/* Used externally as confirmation of correct OCL code */
Expand Down
3 changes: 3 additions & 0 deletions algorithm/blake256.h
Original file line number Diff line number Diff line change
Expand Up @@ -8,4 +8,7 @@ extern void blake256_prepare_work(dev_blk_ctx *blk, uint32_t *state, uint32_t *p
extern void blake256_midstate(struct work *work);
extern void blake256_regenhash(struct work *work);

extern void blake256_prepare_work_112(dev_blk_ctx *blk, uint32_t *state, uint32_t *pdata);
extern void blake256_midstate_112(struct work *work);

#endif /* BLAKE256_H */
136 changes: 136 additions & 0 deletions algorithm/lyra2Zz.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,136 @@
/*-
* Copyright 2014 James Lovejoy
* Copyright 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.
*/

#include "config.h"
#include "miner.h"

#include <stdlib.h>
#include <stdint.h>
#include <string.h>

#include "sph/sph_blake.h"
#include "sph/sph_groestl.h"
#include "sph/sph_skein.h"
#include "sph/sph_keccak.h"
#include "sph/sph_bmw.h"
#include "sph/sph_cubehash.h"
#include "lyra2.h"


void lyra2Zzhash(void *state, const void *input)
{
sph_blake256_context ctx_blake;

uint32_t hashA[8], hashB[8];

sph_blake256_init(&ctx_blake);
sph_blake256 (&ctx_blake, input, 112);
sph_blake256_close (&ctx_blake, hashA);

LYRA2(hashB, 32, hashA, 32, hashA, 32, 8, 8, 8);

memcpy(state, hashB, 32);
}

static const uint32_t diff1targ = 0x0000ffff;


/* Used externally as confirmation of correct OCL code */
int lyra2Zz_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[28], ohash[8];

be32enc_vect(data, (const uint32_t *)pdata, 28);
data[19] = htobe32(nonce);
lyra2Zzhash(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 lyra2Zz_regenhash(struct work *work)
{
uint32_t data[28];
uint32_t *nonce = (uint32_t *)(work->data + 76);
uint32_t *ohash = (uint32_t *)(work->hash);

be32enc_vect(data, (const uint32_t *)work->data, 28);
data[19] = htobe32(*nonce);
lyra2Zzhash(ohash, data);
}

bool scanhash_lyra2Zz(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);
uint32_t data[28];
uint32_t tmp_hash7;
uint32_t Htarg = le32toh(((const uint32_t *)ptarget)[7]);
bool ret = false;

be32enc_vect(data, (const uint32_t *)pdata, 28);

while(1) {
uint32_t ostate[8];

*nonce = ++n;
data[19] = (n);
lyra2Zzhash(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;
}



13 changes: 13 additions & 0 deletions algorithm/lyra2Zz.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
#ifndef LYRA2ZZ_H
#define LYRA2ZZ_H

#include "miner.h"
//#define LYRA2Z_SCRATCHBUF_SIZE (24576) // matrix size [12][4][4] uint64_t or equivalent
#define LYRA2ZZ_SCRATCHBUF_SIZE (12*8*8)
// #define LYRA_SCRATCHBUF_SIZE (1536)
#define LYRA_SECBUF_SIZE (4) // (not used)
extern int lyra2Zz_test(unsigned char *pdata, const unsigned char *ptarget,
uint32_t nonce);
extern void lyra2Zz_regenhash(struct work *work);

#endif /* LYRA2ZZ_H */
18 changes: 13 additions & 5 deletions driver-opencl.c
Original file line number Diff line number Diff line change
Expand Up @@ -1424,11 +1424,7 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,

if (clState->goffset)
p_global_work_offset = (size_t *)&work->blk.nonce;

//if (gpu->algorithm.type == ALGO_LYRA2Z) {
// const size_t local_work_size = 256;
// status = clEnqueueNDRangeKernel(clState->commandQueue, clState->kernel, 1, p_global_work_offset, globalThreads, &local_work_size, 0, NULL, NULL); // blake
//} else

status = clEnqueueNDRangeKernel(clState->commandQueue, clState->kernel, 1, p_global_work_offset,
globalThreads, localThreads, 0, NULL, NULL);
if (unlikely(status != CL_SUCCESS)) {
Expand Down Expand Up @@ -1472,6 +1468,18 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
const size_t expand[] = { 4, 5 };
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i], 2, off2, gws, expand, 0, NULL, NULL); // lyra 4w monolithic
}
} else if (gpu->algorithm.type == ALGO_LYRA2ZZ && i == 1) {
if (clState->prebuilt) {
const size_t off2[] = { 0, 0, *p_global_work_offset };
const size_t gws[] = { 4, 4, globalThreads[0] / 2 };
const size_t expand[] = { 4, 4, 16 };
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i], 3, off2, gws, expand, 0, NULL, NULL); // lyra 4w monolithic
} else {
const size_t off2[] = { 0, *p_global_work_offset };
const size_t gws[] = { 4, globalThreads[0] };
const size_t expand[] = { 4, 5 };
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i], 2, off2, gws, expand, 0, NULL, NULL); // lyra 4w monolithic
}
} else if (gpu->algorithm.type == ALGO_ALLIUM && (i == 2 || i == 6)) {
if (clState->prebuilt) {
const size_t off2[] = { 0, 0, *p_global_work_offset };
Expand Down
7 changes: 6 additions & 1 deletion kernel/blake256.cl
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,12 @@ __constant static const sph_u32 c_u256[16] = {
v[b] = SPH_ROTR32(v[b], 7); \
}


__constant static const sph_u32 c_Padding_112[16] = {
0, 0, 0, 0,
0, 0, 0, 0,
0, 0, 0, 0,
0x80000000, 1, 0, 0x380
};



Loading

0 comments on commit 3054dda

Please sign in to comment.