mirror of https://github.com/GOSTSec/sgminer
ystarnaud
11 years ago
10 changed files with 876 additions and 101 deletions
@ -0,0 +1,182 @@
@@ -0,0 +1,182 @@
|
||||
/*-
|
||||
* 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_jh.h" |
||||
#include "sph/sph_keccak.h" |
||||
#include "sph/sph_skein.h" |
||||
|
||||
/* Move init out of loop, so init once externally, and then use one single memcpy with that bigger memory block */ |
||||
typedef struct |
||||
{ |
||||
sph_blake512_context blake1; |
||||
sph_groestl512_context groestl1; |
||||
sph_jh512_context jh1; |
||||
sph_keccak512_context keccak1; |
||||
sph_skein512_context skein1; |
||||
} Xhash_context_holder; |
||||
|
||||
Xhash_context_holder base_contexts; |
||||
|
||||
void init_Nhash_contexts() |
||||
{ |
||||
sph_blake512_init(&base_contexts.blake1); |
||||
sph_groestl512_init(&base_contexts.groestl1); |
||||
sph_jh512_init(&base_contexts.jh1); |
||||
sph_keccak512_init(&base_contexts.keccak1); |
||||
sph_skein512_init(&base_contexts.skein1); |
||||
} |
||||
|
||||
/*
|
||||
* 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 talkhash(void *state, const void *input) |
||||
{ |
||||
init_Nhash_contexts(); |
||||
|
||||
Xhash_context_holder ctx; |
||||
|
||||
uint32_t hashA[16], hashB[16]; |
||||
//blake-bmw-groestl-sken-jh-meccak-luffa-cubehash-shivite-simd-echo
|
||||
memcpy(&ctx, &base_contexts, sizeof(base_contexts)); |
||||
|
||||
sph_blake512 (&ctx.blake1, input, 80); |
||||
sph_blake512_close (&ctx.blake1, hashA); |
||||
|
||||
sph_groestl512 (&ctx.groestl1, hashA, 64); |
||||
sph_groestl512_close(&ctx.groestl1, hashB); |
||||
|
||||
sph_jh512 (&ctx.jh1, hashB, 64); |
||||
sph_jh512_close(&ctx.jh1, hashA); |
||||
|
||||
sph_keccak512 (&ctx.keccak1, hashA, 64); |
||||
sph_keccak512_close(&ctx.keccak1, hashB); |
||||
|
||||
sph_skein512 (&ctx.skein1, hashB, 64); |
||||
sph_skein512_close(&ctx.skein1, hashA); |
||||
|
||||
memcpy(state, hashA, 32); |
||||
} |
||||
|
||||
static const uint32_t diff1targ = 0x0000ffff; |
||||
|
||||
/* Used externally as confirmation of correct OCL code */ |
||||
int talkcoin_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]; |
||||
|
||||
be32enc_vect(data, (const uint32_t *)pdata, 19); |
||||
data[19] = htobe32(nonce); |
||||
talkhash(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 talkcoin_regenhash(struct work *work) |
||||
{ |
||||
uint32_t data[20]; |
||||
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); |
||||
talkhash(ohash, data); |
||||
} |
||||
|
||||
bool scanhash_talkcoin(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[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); |
||||
talkhash(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; |
||||
} |
||||
@ -0,0 +1,9 @@
@@ -0,0 +1,9 @@
|
||||
#ifndef TALKCOIN_H |
||||
#define TALKCOIN_H |
||||
|
||||
#include "miner.h" |
||||
|
||||
extern int talkcoin_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce); |
||||
extern void talkcoin_regenhash(struct work *work); |
||||
|
||||
#endif /* TALKCOIN_H */ |
||||
@ -0,0 +1,396 @@
@@ -0,0 +1,396 @@
|
||||
/* |
||||
* TalkCoin 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 TALKCOIN_MOD_CL |
||||
#define TALKCOIN_MOD_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) (as_uint(x)) |
||||
#define SPH_ROTL32(x, n) rotate(as_uint(x), as_uint(n)) |
||||
#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n))) |
||||
|
||||
#define SPH_C64(x) ((sph_u64)(x ## UL)) |
||||
#define SPH_T64(x) (as_ulong(x)) |
||||
#define SPH_ROTL64(x, n) rotate(as_ulong(x), (n) & 0xFFFFFFFFFFFFFFFFUL) |
||||
#define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n))) |
||||
|
||||
#define SPH_COMPACT_BLAKE_64 0 |
||||
#define SPH_GROESTL_BIG_ENDIAN 0 |
||||
#define SPH_SMALL_FOOTPRINT_GROESTL 0 |
||||
#define SPH_JH_64 1 |
||||
#define SPH_KECCAK_64 1 |
||||
#define SPH_KECCAK_NOCOPY 0 |
||||
#define SPH_KECCAK_UNROLL 0 |
||||
|
||||
#include "blake.cl" |
||||
#include "groestl.cl" |
||||
#include "jh.cl" |
||||
#include "keccak.cl" |
||||
#include "skein.cl" |
||||
|
||||
#define SWAP4(x) as_uint(as_uchar4(x).wzyx) |
||||
#define SWAP8(x) as_ulong(as_uchar8(x).s76543210) |
||||
|
||||
#if SPH_BIG_ENDIAN |
||||
#define DEC64E(x) (x) |
||||
#define DEC64BE(x) (*(const __global sph_u64 *) (x)); |
||||
#else |
||||
#define DEC64E(x) SWAP8(x) |
||||
#define DEC64BE(x) SWAP8(*(const __global sph_u64 *) (x)); |
||||
#endif |
||||
|
||||
|
||||
typedef union { |
||||
unsigned char h1[64]; |
||||
uint h4[16]; |
||||
ulong h8[8]; |
||||
} hash_t; |
||||
|
||||
// blake |
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) |
||||
__kernel void search(__global unsigned char* block, __global hash_t* hashes) |
||||
{ |
||||
uint gid = get_global_id(0); |
||||
__global hash_t *hash = &(hashes[gid-get_global_offset(0)]); |
||||
|
||||
sph_u64 H0 = SPH_C64(0x6A09E667F3BCC908), H1 = SPH_C64(0xBB67AE8584CAA73B); |
||||
sph_u64 H2 = SPH_C64(0x3C6EF372FE94F82B), H3 = SPH_C64(0xA54FF53A5F1D36F1); |
||||
sph_u64 H4 = SPH_C64(0x510E527FADE682D1), H5 = SPH_C64(0x9B05688C2B3E6C1F); |
||||
sph_u64 H6 = SPH_C64(0x1F83D9ABFB41BD6B), H7 = SPH_C64(0x5BE0CD19137E2179); |
||||
sph_u64 S0 = 0, S1 = 0, S2 = 0, S3 = 0; |
||||
sph_u64 T0 = SPH_C64(0xFFFFFFFFFFFFFC00) + (80 << 3), T1 = 0xFFFFFFFFFFFFFFFF;; |
||||
|
||||
if ((T0 = SPH_T64(T0 + 1024)) < 1024) |
||||
{ |
||||
T1 = SPH_T64(T1 + 1); |
||||
} |
||||
sph_u64 M0, M1, M2, M3, M4, M5, M6, M7; |
||||
sph_u64 M8, M9, MA, MB, MC, MD, ME, MF; |
||||
sph_u64 V0, V1, V2, V3, V4, V5, V6, V7; |
||||
sph_u64 V8, V9, VA, VB, VC, VD, VE, VF; |
||||
M0 = DEC64BE(block + 0); |
||||
M1 = DEC64BE(block + 8); |
||||
M2 = DEC64BE(block + 16); |
||||
M3 = DEC64BE(block + 24); |
||||
M4 = DEC64BE(block + 32); |
||||
M5 = DEC64BE(block + 40); |
||||
M6 = DEC64BE(block + 48); |
||||
M7 = DEC64BE(block + 56); |
||||
M8 = DEC64BE(block + 64); |
||||
M9 = DEC64BE(block + 72); |
||||
M9 &= 0xFFFFFFFF00000000; |
||||
M9 ^= SWAP4(gid); |
||||
MA = 0x8000000000000000; |
||||
MB = 0; |
||||
MC = 0; |
||||
MD = 1; |
||||
ME = 0; |
||||
MF = 0x280; |
||||
|
||||
COMPRESS64; |
||||
|
||||
hash->h8[0] = H0; |
||||
hash->h8[1] = H1; |
||||
hash->h8[2] = H2; |
||||
hash->h8[3] = H3; |
||||
hash->h8[4] = H4; |
||||
hash->h8[5] = H5; |
||||
hash->h8[6] = H6; |
||||
hash->h8[7] = H7; |
||||
|
||||
barrier(CLK_GLOBAL_MEM_FENCE); |
||||
} |
||||
|
||||
// groestl |
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) |
||||
__kernel void search1(__global hash_t* hashes) |
||||
{ |
||||
uint gid = get_global_id(0); |
||||
__global hash_t *hash = &(hashes[gid-get_global_offset(0)]); |
||||
|
||||
__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 |
||||
|
||||
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(hash->h8[0]); |
||||
m[1] = DEC64E(hash->h8[1]); |
||||
m[2] = DEC64E(hash->h8[2]); |
||||
m[3] = DEC64E(hash->h8[3]); |
||||
m[4] = DEC64E(hash->h8[4]); |
||||
m[5] = DEC64E(hash->h8[5]); |
||||
m[6] = DEC64E(hash->h8[6]); |
||||
m[7] = DEC64E(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]; |
||||
|
||||
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] = DEC64E(H[u + 8]); |
||||
|
||||
barrier(CLK_GLOBAL_MEM_FENCE); |
||||
} |
||||
|
||||
// jh |
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) |
||||
__kernel void search2(__global hash_t* hashes) |
||||
{ |
||||
uint gid = get_global_id(0); |
||||
__global hash_t *hash = &(hashes[gid-get_global_offset(0)]); |
||||
|
||||
sph_u64 h0h = C64e(0x6fd14b963e00aa17), h0l = C64e(0x636a2e057a15d543), h1h = C64e(0x8a225e8d0c97ef0b), h1l = C64e(0xe9341259f2b3c361), h2h = C64e(0x891da0c1536f801e), h2l = C64e(0x2aa9056bea2b6d80), h3h = C64e(0x588eccdb2075baa6), h3l = C64e(0xa90f3a76baf83bf7); |
||||
sph_u64 h4h = C64e(0x0169e60541e34a69), h4l = C64e(0x46b58a8e2e6fe65a), h5h = C64e(0x1047a7d0c1843c24), h5l = C64e(0x3b6e71b12d5ac199), h6h = C64e(0xcf57f6ec9db1f856), h6l = C64e(0xa706887c5716b156), h7h = C64e(0xe3c2fcdfe68517fb), h7l = C64e(0x545a4678cc8cdd4b); |
||||
sph_u64 tmp; |
||||
|
||||
for(int i = 0; i < 2; i++) |
||||
{ |
||||
if (i == 0) |
||||
{ |
||||
h0h ^= DEC64E(hash->h8[0]); |
||||
h0l ^= DEC64E(hash->h8[1]); |
||||
h1h ^= DEC64E(hash->h8[2]); |
||||
h1l ^= DEC64E(hash->h8[3]); |
||||
h2h ^= DEC64E(hash->h8[4]); |
||||
h2l ^= DEC64E(hash->h8[5]); |
||||
h3h ^= DEC64E(hash->h8[6]); |
||||
h3l ^= DEC64E(hash->h8[7]); |
||||
} |
||||
else if(i == 1) |
||||
{ |
||||
h4h ^= DEC64E(hash->h8[0]); |
||||
h4l ^= DEC64E(hash->h8[1]); |
||||
h5h ^= DEC64E(hash->h8[2]); |
||||
h5l ^= DEC64E(hash->h8[3]); |
||||
h6h ^= DEC64E(hash->h8[4]); |
||||
h6l ^= DEC64E(hash->h8[5]); |
||||
h7h ^= DEC64E(hash->h8[6]); |
||||
h7l ^= DEC64E(hash->h8[7]); |
||||
|
||||
h0h ^= 0x80; |
||||
h3l ^= 0x2000000000000; |
||||
} |
||||
E8; |
||||
} |
||||
|
||||
h4h ^= 0x80; |
||||
h7l ^= 0x2000000000000; |
||||
|
||||
hash->h8[0] = DEC64E(h4h); |
||||
hash->h8[1] = DEC64E(h4l); |
||||
hash->h8[2] = DEC64E(h5h); |
||||
hash->h8[3] = DEC64E(h5l); |
||||
hash->h8[4] = DEC64E(h6h); |
||||
hash->h8[5] = DEC64E(h6l); |
||||
hash->h8[6] = DEC64E(h7h); |
||||
hash->h8[7] = DEC64E(h7l); |
||||
|
||||
barrier(CLK_GLOBAL_MEM_FENCE); |
||||
} |
||||
|
||||
|
||||
// keccak |
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) |
||||
__kernel void search3(__global hash_t* hashes) |
||||
{ |
||||
uint gid = get_global_id(0); |
||||
__global hash_t *hash = &(hashes[gid-get_global_offset(0)]); |
||||
|
||||
// keccak |
||||
|
||||
sph_u64 a00 = 0, a01 = 0, a02 = 0, a03 = 0, a04 = 0; |
||||
sph_u64 a10 = 0, a11 = 0, a12 = 0, a13 = 0, a14 = 0; |
||||
sph_u64 a20 = 0, a21 = 0, a22 = 0, a23 = 0, a24 = 0; |
||||
sph_u64 a30 = 0, a31 = 0, a32 = 0, a33 = 0, a34 = 0; |
||||
sph_u64 a40 = 0, a41 = 0, a42 = 0, a43 = 0, a44 = 0; |
||||
|
||||
a10 = SPH_C64(0xFFFFFFFFFFFFFFFF); |
||||
a20 = SPH_C64(0xFFFFFFFFFFFFFFFF); |
||||
a31 = SPH_C64(0xFFFFFFFFFFFFFFFF); |
||||
a22 = SPH_C64(0xFFFFFFFFFFFFFFFF); |
||||
a23 = SPH_C64(0xFFFFFFFFFFFFFFFF); |
||||
a04 = SPH_C64(0xFFFFFFFFFFFFFFFF); |
||||
|
||||
a00 ^= SWAP8(hash->h8[0]); |
||||
a10 ^= SWAP8(hash->h8[1]); |
||||
a20 ^= SWAP8(hash->h8[2]); |
||||
a30 ^= SWAP8(hash->h8[3]); |
||||
a40 ^= SWAP8(hash->h8[4]); |
||||
a01 ^= SWAP8(hash->h8[5]); |
||||
a11 ^= SWAP8(hash->h8[6]); |
||||
a21 ^= SWAP8(hash->h8[7]); |
||||
a31 ^= 0x8000000000000001; |
||||
KECCAK_F_1600; |
||||
|
||||
// Finalize the "lane complement" |
||||
a10 = ~a10; |
||||
a20 = ~a20; |
||||
|
||||
hash->h8[0] = SWAP8(a00); |
||||
hash->h8[1] = SWAP8(a10); |
||||
hash->h8[2] = SWAP8(a20); |
||||
hash->h8[3] = SWAP8(a30); |
||||
hash->h8[4] = SWAP8(a40); |
||||
hash->h8[5] = SWAP8(a01); |
||||
hash->h8[6] = SWAP8(a11); |
||||
hash->h8[7] = SWAP8(a21); |
||||
|
||||
barrier(CLK_GLOBAL_MEM_FENCE); |
||||
} |
||||
|
||||
|
||||
// skein |
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) |
||||
__kernel void search4(__global hash_t* hashes, __global uint* output, const ulong target) |
||||
{ |
||||
uint gid = get_global_id(0); |
||||
uint offset = get_global_offset(0); |
||||
hash_t hash; |
||||
__global hash_t *hashp = &(hashes[gid-offset]); |
||||
|
||||
for (int i = 0; i < 8; i++) |
||||
hash.h8[i] = hashes[gid-offset].h8[i]; |
||||
|
||||
sph_u64 h0 = SPH_C64(0x4903ADFF749C51CE), h1 = SPH_C64(0x0D95DE399746DF03), h2 = SPH_C64(0x8FD1934127C79BCE), h3 = SPH_C64(0x9A255629FF352CB1), h4 = SPH_C64(0x5DB62599DF6CA7B0), h5 = SPH_C64(0xEABE394CA9D5C3F4), h6 = SPH_C64(0x991112C71A75B523), h7 = SPH_C64(0xAE18A40B660FCC33); |
||||
sph_u64 m0, m1, m2, m3, m4, m5, m6, m7; |
||||
sph_u64 bcount = 0; |
||||
|
||||
m0 = SWAP8(hash.h8[0]); |
||||
m1 = SWAP8(hash.h8[1]); |
||||
m2 = SWAP8(hash.h8[2]); |
||||
m3 = SWAP8(hash.h8[3]); |
||||
m4 = SWAP8(hash.h8[4]); |
||||
m5 = SWAP8(hash.h8[5]); |
||||
m6 = SWAP8(hash.h8[6]); |
||||
m7 = SWAP8(hash.h8[7]); |
||||
|
||||
UBI_BIG(480, 64); |
||||
|
||||
bcount = 0; |
||||
m0 = m1 = m2 = m3 = m4 = m5 = m6 = m7 = 0; |
||||
|
||||
UBI_BIG(510, 8); |
||||
|
||||
hash.h8[0] = SWAP8(h0); |
||||
hash.h8[1] = SWAP8(h1); |
||||
hash.h8[2] = SWAP8(h2); |
||||
hash.h8[3] = SWAP8(h3); |
||||
hash.h8[4] = SWAP8(h4); |
||||
hash.h8[5] = SWAP8(h5); |
||||
hash.h8[6] = SWAP8(h6); |
||||
hash.h8[7] = SWAP8(h7); |
||||
|
||||
bool result = (SWAP8(hash.h8[3]) <= target); |
||||
if (result) |
||||
output[atomic_inc(output+0xFF)] = SWAP4(gid); |
||||
} |
||||
|
||||
#endif // TALKCOIN_MOD_CL |
||||
Loading…
Reference in new issue