mirror of
https://github.com/GOSTSec/sgminer
synced 2025-03-09 12:11:02 +00:00
Add Pascal algo for PascalCoin
This commit is contained in:
elbandi
parent
673dc79a40
commit
fb181f5e88
@ -82,6 +82,7 @@ sgminer_SOURCES += algorithm/yescrypt.h algorithm/yescrypt.c algorithm/yescrypt_
|
||||
sgminer_SOURCES += algorithm/blake256.c algorithm/blake256.h
|
||||
sgminer_SOURCES += algorithm/blakecoin.c algorithm/blakecoin.h
|
||||
sgminer_SOURCES += algorithm/decred.c algorithm/decred.h
|
||||
sgminer_SOURCES += algorithm/pascal.c algorithm/pascal.h
|
||||
sgminer_SOURCES += algorithm/lbry.c algorithm/lbry.h
|
||||
|
||||
bin_SCRIPTS = $(top_srcdir)/kernel/*.cl
|
||||
|
||||
23
algorithm.c
23
algorithm.c
@ -41,6 +41,7 @@
|
||||
#include "algorithm/blakecoin.h"
|
||||
#include "algorithm/sia.h"
|
||||
#include "algorithm/decred.h"
|
||||
#include "algorithm/pascal.h"
|
||||
#include "algorithm/lbry.h"
|
||||
#include "algorithm/sibcoin.h"
|
||||
|
||||
@ -54,6 +55,7 @@ const char *algorithm_type_str[] = {
|
||||
"Credits",
|
||||
"Scrypt",
|
||||
"NScrypt",
|
||||
"Pascal",
|
||||
"X11",
|
||||
"X13",
|
||||
"X14",
|
||||
@ -204,6 +206,25 @@ static cl_int queue_scrypt_kernel(struct __clState *clState, struct _dev_blk_ctx
|
||||
return status;
|
||||
}
|
||||
|
||||
static cl_int queue_pascal_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
|
||||
{
|
||||
cl_kernel *kernel = &clState->kernel;
|
||||
unsigned int num = 0;
|
||||
cl_ulong le_target;
|
||||
cl_int status = 0;
|
||||
|
||||
le_target = *(cl_ulong *)(blk->work->device_target + 24);
|
||||
flip196(clState->cldata, blk->work->data);
|
||||
status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 196, clState->cldata, 0, NULL, NULL);
|
||||
|
||||
CL_SET_ARG(clState->CLbuffer0);
|
||||
CL_SET_ARG(clState->outputBuffer);
|
||||
CL_SET_ARG(le_target);
|
||||
CL_SET_ARG(blk->work->midstate);
|
||||
|
||||
return status;
|
||||
}
|
||||
|
||||
static cl_int queue_neoscrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads)
|
||||
{
|
||||
cl_kernel *kernel = &clState->kernel;
|
||||
@ -1206,6 +1227,8 @@ static algorithm_settings_t algos[] = {
|
||||
|
||||
{ "lbry", ALGO_LBRY, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 2, 4 * 8 * 4194304, 0, lbry_regenhash, NULL, NULL, queue_lbry_kernel, gen_hash, NULL },
|
||||
|
||||
{ "pascal", ALGO_PASCAL, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 0, 0, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, pascal_regenhash, pascal_midstate, NULL, queue_pascal_kernel, NULL, NULL },
|
||||
|
||||
// Terminator (do not remove)
|
||||
{ NULL, ALGO_UNK, "", 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, NULL, NULL, NULL, NULL, NULL }
|
||||
};
|
||||
|
||||
@ -16,6 +16,7 @@ typedef enum {
|
||||
ALGO_CRE,
|
||||
ALGO_SCRYPT,
|
||||
ALGO_NSCRYPT,
|
||||
ALGO_PASCAL,
|
||||
ALGO_X11,
|
||||
ALGO_X13,
|
||||
ALGO_X14,
|
||||
|
||||
117
algorithm/pascal.c
Normal file
117
algorithm/pascal.c
Normal file
@ -0,0 +1,117 @@
|
||||
/*
|
||||
* BLAKE implementation.
|
||||
*
|
||||
* ==========================(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
|
||||
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
|
||||
*
|
||||
* Modified for more speed by BlueDragon747 for the Blakecoin project
|
||||
*/
|
||||
|
||||
#include <stddef.h>
|
||||
#include <string.h>
|
||||
#include <limits.h>
|
||||
#include <stdint.h>
|
||||
|
||||
#include "sph/sph_sha2.h"
|
||||
#include "algorithm/pascal.h"
|
||||
|
||||
static const uint32_t diff1targ_pascal = 0x000000ff;
|
||||
|
||||
void pascalhash(void *state, const void *input)
|
||||
{
|
||||
sph_sha256_context ctx_sha;
|
||||
uint32_t hash[16];
|
||||
|
||||
sph_sha256_init(&ctx_sha);
|
||||
sph_sha256(&ctx_sha, input, 200);
|
||||
sph_sha256_close(&ctx_sha, hash);
|
||||
|
||||
sph_sha256_init(&ctx_sha);
|
||||
sph_sha256(&ctx_sha, hash, 32);
|
||||
sph_sha256_close(&ctx_sha, hash);
|
||||
|
||||
memcpy(state, hash, 32);
|
||||
|
||||
}
|
||||
|
||||
void pascal_midstate(struct work *work)
|
||||
{
|
||||
sph_sha256_context ctx_sha;
|
||||
uint32_t data[48];
|
||||
|
||||
memcpy(data, work->data, 192);
|
||||
|
||||
sph_sha256_init(&ctx_sha);
|
||||
sph_sha256 (&ctx_sha, (unsigned char *)data, 192);
|
||||
|
||||
memcpy(work->midstate, ctx_sha.val, 32);
|
||||
endian_flip32(work->midstate, work->midstate);
|
||||
|
||||
/*
|
||||
char *strdata, *strmidstate;
|
||||
strdata = bin2hex(work->data, 192);
|
||||
strmidstate = bin2hex(work->midstate, 32);
|
||||
applog(LOG_DEBUG, "data %s midstate %s", strdata, strmidstate);
|
||||
*/
|
||||
}
|
||||
|
||||
static const uint32_t diff1targ = 0x0000ffff;
|
||||
|
||||
/* Used externally as confirmation of correct OCL code */
|
||||
int pascal_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[50], ohash[8];
|
||||
|
||||
memcpy(data, pdata, 200);
|
||||
data[49] = htobe32(nonce);
|
||||
pascalhash(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 pascal_regenhash(struct work *work)
|
||||
{
|
||||
uint32_t data[64];
|
||||
uint32_t hash[16];
|
||||
uint32_t *nonce = (uint32_t *)(work->data + 196);
|
||||
uint32_t *ohash = (uint32_t *)(work->hash);
|
||||
|
||||
memcpy(data, work->data, 200);
|
||||
data[49] = htole32(*nonce);
|
||||
pascalhash(hash, data);
|
||||
swab256(ohash, hash);
|
||||
}
|
||||
11
algorithm/pascal.h
Normal file
11
algorithm/pascal.h
Normal file
@ -0,0 +1,11 @@
|
||||
#ifndef PASCAL_H
|
||||
#define PASCAL_H
|
||||
|
||||
#include "miner.h"
|
||||
|
||||
extern int pascal_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce);
|
||||
extern void pascal_prepare_work(dev_blk_ctx *blk, uint32_t *state, uint32_t *pdata);
|
||||
extern void pascal_midstate(struct work *work);
|
||||
extern void pascal_regenhash(struct work *work);
|
||||
|
||||
#endif /* PASCAL_H */
|
||||
229
kernel/pascal.cl
Normal file
229
kernel/pascal.cl
Normal file
@ -0,0 +1,229 @@
|
||||
/*
|
||||
* "pascal" kernel implementation.
|
||||
*
|
||||
* ==========================(LICENSE BEGIN)============================
|
||||
*
|
||||
* Copyright (c) 2015 djm34
|
||||
*
|
||||
* 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 djm34
|
||||
*/
|
||||
#if !defined(cl_khr_byte_addressable_store)
|
||||
#error "Device does not support unaligned stores"
|
||||
#endif
|
||||
|
||||
|
||||
#define ROL32(x, n) rotate(x, (uint) n)
|
||||
#define SWAP32(a) (as_uint(as_uchar4(a).wzyx))
|
||||
#define SWAP64(x) as_ulong(as_uchar8(x).s32107654) /// hmm...
|
||||
|
||||
#define SHR(x, n) ((x) >> n)
|
||||
|
||||
#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 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)
|
||||
|
||||
/// generic sha transform
|
||||
inline uint8 sha256_Transform(uint16 data, uint8 state)
|
||||
{
|
||||
uint temp1;
|
||||
uint8 res = state;
|
||||
uint W0 = data.s0;
|
||||
uint W1 = data.s1;
|
||||
uint W2 = data.s2;
|
||||
uint W3 = data.s3;
|
||||
uint W4 = data.s4;
|
||||
uint W5 = data.s5;
|
||||
uint W6 = data.s6;
|
||||
uint W7 = data.s7;
|
||||
uint W8 = data.s8;
|
||||
uint W9 = data.s9;
|
||||
uint W10 = data.sA;
|
||||
uint W11 = data.sB;
|
||||
uint W12 = data.sC;
|
||||
uint W13 = data.sD;
|
||||
uint W14 = data.sE;
|
||||
uint W15 = data.sF;
|
||||
|
||||
#define v0 res.s0
|
||||
#define v1 res.s1
|
||||
#define v2 res.s2
|
||||
#define v3 res.s3
|
||||
#define v4 res.s4
|
||||
#define v5 res.s5
|
||||
#define v6 res.s6
|
||||
#define v7 res.s7
|
||||
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, W0, 0x428A2F98);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, W1, 0x71374491);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, W2, 0xB5C0FBCF);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, W3, 0xE9B5DBA5);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, W4, 0x3956C25B);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, W5, 0x59F111F1);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, W6, 0x923F82A4);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, W7, 0xAB1C5ED5);
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, W8, 0xD807AA98);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, W9, 0x12835B01);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, W10, 0x243185BE);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, W11, 0x550C7DC3);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, W12, 0x72BE5D74);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, W13, 0x80DEB1FE);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, W14, 0x9BDC06A7);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, W15, 0xC19BF174);
|
||||
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, R0, 0xE49B69C1);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, R1, 0xEFBE4786);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, R2, 0x0FC19DC6);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, R3, 0x240CA1CC);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, R4, 0x2DE92C6F);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, R5, 0x4A7484AA);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, R6, 0x5CB0A9DC);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, R7, 0x76F988DA);
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, R8, 0x983E5152);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, R9, 0xA831C66D);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, R10, 0xB00327C8);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, R11, 0xBF597FC7);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, R12, 0xC6E00BF3);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xD5A79147);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, R14, 0x06CA6351);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, R15, 0x14292967);
|
||||
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, R0, 0x27B70A85);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, R1, 0x2E1B2138);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, R2, 0x4D2C6DFC);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, R3, 0x53380D13);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, R4, 0x650A7354);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, R5, 0x766A0ABB);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, R6, 0x81C2C92E);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, R7, 0x92722C85);
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, R8, 0xA2BFE8A1);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, R9, 0xA81A664B);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, R10, 0xC24B8B70);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, R11, 0xC76C51A3);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, R12, 0xD192E819);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xD6990624);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, R14, 0xF40E3585);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, R15, 0x106AA070);
|
||||
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, R0, 0x19A4C116);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, R1, 0x1E376C08);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, R2, 0x2748774C);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, R3, 0x34B0BCB5);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, R4, 0x391C0CB3);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, R5, 0x4ED8AA4A);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, R6, 0x5B9CCA4F);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, R7, 0x682E6FF3);
|
||||
P(v0, v1, v2, v3, v4, v5, v6, v7, R8, 0x748F82EE);
|
||||
P(v7, v0, v1, v2, v3, v4, v5, v6, R9, 0x78A5636F);
|
||||
P(v6, v7, v0, v1, v2, v3, v4, v5, R10, 0x84C87814);
|
||||
P(v5, v6, v7, v0, v1, v2, v3, v4, R11, 0x8CC70208);
|
||||
P(v4, v5, v6, v7, v0, v1, v2, v3, R12, 0x90BEFFFA);
|
||||
P(v3, v4, v5, v6, v7, v0, v1, v2, R13, 0xA4506CEB);
|
||||
P(v2, v3, v4, v5, v6, v7, v0, v1, RD14, 0xBEF9A3F7);
|
||||
P(v1, v2, v3, v4, v5, v6, v7, v0, RD15, 0xC67178F2);
|
||||
#undef v0
|
||||
#undef v1
|
||||
#undef v2
|
||||
#undef v3
|
||||
#undef v4
|
||||
#undef v5
|
||||
#undef v6
|
||||
#undef v7
|
||||
return (res + state);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __constant uint8 H256 = {
|
||||
0x6A09E667, 0xBB67AE85, 0x3C6EF372,
|
||||
0xA54FF53A, 0x510E527F, 0x9B05688C,
|
||||
0x1F83D9AB, 0x5BE0CD19
|
||||
};
|
||||
|
||||
|
||||
static __constant uint16 pad_data =
|
||||
{
|
||||
0x00000000, 0x00000000, 0x80000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000640
|
||||
};
|
||||
|
||||
static __constant uint8 pad_state =
|
||||
{
|
||||
0x80000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000100
|
||||
};
|
||||
|
||||
|
||||
|
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
|
||||
__kernel void search(__global const uchar* restrict input, __global uint* restrict output,const ulong target, uint8 midstate )
|
||||
{
|
||||
uint nonce = get_global_id(0);
|
||||
uint16 in;
|
||||
uint8 state1;
|
||||
in = pad_data;
|
||||
in.s0 = ((__global const uint *)input)[48];
|
||||
in.s1 = nonce;
|
||||
state1 = sha256_Transform(in, midstate);
|
||||
in.lo = state1;
|
||||
in.hi = pad_state;
|
||||
state1 = sha256_Transform(in, H256);
|
||||
|
||||
if (as_ulong(state1.s10) <= target) {
|
||||
output[atomic_inc(output + 0xFF)] = SWAP32(nonce);
|
||||
}
|
||||
}
|
||||
|
||||
18
miner.h
18
miner.h
@ -756,6 +756,16 @@ static inline void flip180(void *dest_p, const void *src_p)
|
||||
dest[i] = swab32(src[i]);
|
||||
}
|
||||
|
||||
static inline void flip196(void *dest_p, const void *src_p)
|
||||
{
|
||||
uint32_t *dest = (uint32_t *)dest_p;
|
||||
const uint32_t *src = (uint32_t *)src_p;
|
||||
int i;
|
||||
|
||||
for (i = 0; i < 49; i++)
|
||||
dest[i] = swab32(src[i]);
|
||||
}
|
||||
|
||||
/*
|
||||
* 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.
|
||||
@ -788,6 +798,10 @@ static inline void endian_flip180(void *dest_p, const void *src_p)
|
||||
{
|
||||
flip180(dest_p, src_p);
|
||||
}
|
||||
static inline void endian_flip196(void *dest_p, const void *src_p)
|
||||
{
|
||||
flip196(dest_p, src_p);
|
||||
}
|
||||
|
||||
#else
|
||||
static inline void
|
||||
@ -807,6 +821,10 @@ static inline void
|
||||
endian_flip180(void __maybe_unused *dest_p, const void __maybe_unused *src_p)
|
||||
{
|
||||
}
|
||||
static inline void
|
||||
endian_flip196(void __maybe_unused *dest_p, const void __maybe_unused *src_p)
|
||||
{
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
1
ocl.c
1
ocl.c
@ -762,6 +762,7 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize, algorithm_t *alg
|
||||
if (algorithm->type == ALGO_CRE) readbufsize = 168;
|
||||
else if (algorithm->type == ALGO_DECRED) readbufsize = 192;
|
||||
else if (algorithm->type == ALGO_LBRY) readbufsize = 112;
|
||||
else if (algorithm->type == ALGO_PASCAL) readbufsize = 196;
|
||||
|
||||
if (algorithm->rw_buffer_size < 0) {
|
||||
// calc buffer size for neoscrypt
|
||||
|
||||
26
sgminer.c
26
sgminer.c
@ -5633,6 +5633,9 @@ static void *stratum_sthread(void *userdata)
|
||||
else if (pool->algorithm.type == ALGO_SIA) {
|
||||
nonce = *((uint32_t *)(work->data + 32));
|
||||
}
|
||||
else if (pool->algorithm.type == ALGO_PASCAL) {
|
||||
nonce = htobe32(*((uint32_t *)(work->data + 196)));
|
||||
}
|
||||
else {
|
||||
nonce = *((uint32_t *)(work->data + 76));
|
||||
}
|
||||
@ -6118,6 +6121,15 @@ static void gen_stratum_work(struct pool *pool, struct work *work)
|
||||
|
||||
cg_wlock(&pool->data_lock);
|
||||
|
||||
if (pool->algorithm.type == ALGO_PASCAL) {
|
||||
/* TODO: refactor this */
|
||||
for (i = 0; i < 56; i += 8) {
|
||||
if (((pool->nonce2 >> i) & 0xff) < 0x2d) pool->nonce2 = (pool->nonce2 & (0xffffffffffffff00 << i)) + (0x002d2d2d2d2d2d2d >> (48 - i));
|
||||
if (((pool->nonce2 >> i) & 0xff) > 0xfe) pool->nonce2 = (pool->nonce2 & (0xffffffffffffff00 << i)) + (0x012d2d2d2d2d2d2d >> (48 - i));
|
||||
}
|
||||
if (((pool->nonce2 >> 56) & 0xff) < 0x2d) pool->nonce2 = 0x2d2d2d2d2d2d2d2d;
|
||||
if (((pool->nonce2 >> 56) & 0xff) > 0xfe) pool->nonce2 = 0x2d2d2d2d2d2d2d2d;
|
||||
}
|
||||
nonce2le = htole64(pool->nonce2);
|
||||
if (pool->algorithm.type != ALGO_DECRED && pool->algorithm.type != ALGO_SIA) {
|
||||
/* Update coinbase. Always use an LE encoded nonce2 to fill in values
|
||||
@ -6130,7 +6142,7 @@ static void gen_stratum_work(struct pool *pool, struct work *work)
|
||||
/* Downgrade to a read lock to read off the pool variables */
|
||||
cg_dwlock(&pool->data_lock);
|
||||
|
||||
if (pool->algorithm.type != ALGO_DECRED && pool->algorithm.type != ALGO_SIA) {
|
||||
if (pool->algorithm.type != ALGO_DECRED && pool->algorithm.type != ALGO_SIA && pool->algorithm.type != ALGO_PASCAL) {
|
||||
/* Generate merkle root */
|
||||
pool->algorithm.gen_hash(pool->coinbase, pool->swork.cb_len, merkle_root);
|
||||
memcpy(merkle_sha, merkle_root, 32);
|
||||
@ -6192,6 +6204,14 @@ static void gen_stratum_work(struct pool *pool, struct work *work)
|
||||
memcpy(work->data + 32 + 8, pool->header_bin + 68, 4); // timestamp
|
||||
flip32(work->data + 32 + 8 + 8, pool->coinbase); // merkleroot
|
||||
}
|
||||
else if (pool->algorithm.type == ALGO_PASCAL) {
|
||||
uint32_t temp;
|
||||
memcpy(work->data, pool->coinbase, pool->swork.cb_len);
|
||||
hex2bin((unsigned char *)&temp, pool->swork.ntime, 4);
|
||||
/* Add the nbits (big endianess). */
|
||||
((uint32_t *)work->data)[48] = be32toh(temp);
|
||||
((uint32_t *)work->data)[49] = 0;
|
||||
}
|
||||
else {
|
||||
data32 = (uint32_t *)merkle_sha;
|
||||
swap32 = (uint32_t *)merkle_root;
|
||||
@ -6216,9 +6236,10 @@ static void gen_stratum_work(struct pool *pool, struct work *work)
|
||||
char *header, *merkle_hash;
|
||||
int datasize = 128;
|
||||
if (pool->algorithm.type == ALGO_DECRED) datasize = 180;
|
||||
else if (pool->algorithm.type == ALGO_PASCAL) datasize = 256;
|
||||
|
||||
header = bin2hex(work->data, datasize);
|
||||
if (pool->algorithm.type != ALGO_DECRED) {
|
||||
if (pool->algorithm.type != ALGO_DECRED && pool->algorithm.type != ALGO_SIA && pool->algorithm.type != ALGO_PASCAL) {
|
||||
merkle_hash = bin2hex((const unsigned char *)merkle_root, 32);
|
||||
applog(LOG_DEBUG, "[THR%d] Generated stratum merkle %s", work->thr_id, merkle_hash);
|
||||
free(merkle_hash);
|
||||
@ -7165,6 +7186,7 @@ static void rebuild_nonce(struct work *work, uint32_t nonce)
|
||||
else if (work->pool->algorithm.type == ALGO_DECRED) nonce_pos = 140;
|
||||
else if (work->pool->algorithm.type == ALGO_LBRY) nonce_pos = 108;
|
||||
else if (work->pool->algorithm.type == ALGO_SIA) nonce_pos = 32;
|
||||
else if (work->pool->algorithm.type == ALGO_PASCAL) nonce_pos = 196;
|
||||
|
||||
uint32_t *work_nonce = (uint32_t *)(work->data + nonce_pos);
|
||||
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user