GOSTSec/sgminer
1
0
Fork 0
mirror of https://github.com/GOSTSec/sgminer synced 2025-08-28 14:52:04 +00:00
Code Issues Releases Wiki Activity

Merge branch 'cgminer' of git://github.com/ckolivas/cgminer into cgminer

Browse Source
This commit is contained in:
Ycros 2011-07-23 16:00:49 +10:00
commit a150140eb7
6 changed files with 546 additions and 447 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
Makefile.am
View File

@ -15,7 +15,7 @@ INCLUDES = $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_INCLUDES)
bin_PROGRAMS = cgminer bin_PROGRAMS = cgminer
bin_SCRIPTS = phatk110714.cl poclbm110717.cl bin_SCRIPTS = phatk110722.cl poclbm110717.cl
cgminer_SOURCES = elist.h miner.h compat.h \ cgminer_SOURCES = elist.h miner.h compat.h \
main.c util.c \ main.c util.c \
@ -23,7 +23,7 @@ cgminer_SOURCES = elist.h miner.h compat.h \
sha256_generic.c sha256_4way.c sha256_via.c \ sha256_generic.c sha256_4way.c sha256_via.c \
sha256_cryptopp.c sha256_sse2_amd64.c \ sha256_cryptopp.c sha256_sse2_amd64.c \
sha256_sse4_amd64.c \ sha256_sse4_amd64.c \
phatk110714.cl poclbm110717.cl phatk110722.cl poclbm110717.cl
cgminer_LDFLAGS = $(PTHREAD_FLAGS) cgminer_LDFLAGS = $(PTHREAD_FLAGS)
cgminer_LDADD = @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @OPENCL_LIBS@ @NCURSES_LIBS@ @PDCURSES_LIBS@ lib/libgnu.a ccan/libccan.a cgminer_LDADD = @LIBCURL@ @JANSSON_LIBS@ @PTHREAD_LIBS@ @OPENCL_LIBS@ @NCURSES_LIBS@ @PDCURSES_LIBS@ lib/libgnu.a ccan/libccan.a

13
findnonce.c
View File

@ -72,7 +72,13 @@ void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data) {
blk->cty_a = A; blk->cty_a = A;
blk->cty_b = B; blk->cty_b = B;
blk->cty_c = C; blk->cty_c = C;
blk->C1addK5 = C + 0x59f111f1;
blk->cty_d = D; blk->cty_d = D;
blk->D1A = D + 0xb956c25b;
blk->cty_e = E; blk->cty_e = E;
blk->cty_f = F; blk->cty_f = F;
blk->cty_g = G; blk->cty_g = G;
@ -94,6 +100,10 @@ void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data) {
blk->W16 = blk->fW0 = data[0] + (rotr(data[1], 7) ^ rotr(data[1], 18) ^ (data[1] >> 3)); blk->W16 = blk->fW0 = data[0] + (rotr(data[1], 7) ^ rotr(data[1], 18) ^ (data[1] >> 3));
blk->W17 = blk->fW1 = data[1] + (rotr(data[2], 7) ^ rotr(data[2], 18) ^ (data[2] >> 3)) + 0x01100000; blk->W17 = blk->fW1 = data[1] + (rotr(data[2], 7) ^ rotr(data[2], 18) ^ (data[2] >> 3)) + 0x01100000;
blk->W2 = data[2]; blk->W2 = data[2];
blk->W2A = blk->W2 + (rotr(blk->W16, 19) ^ rotr(blk->W16, 17) ^ (blk->W16 >> 10));
blk->W17_2 = 0x11002000 + (rotr(blk->W17, 19) ^ rotr(blk->W17, 17) ^ (blk->W17 >> 10));
blk->fW2 = data[2] + (rotr(blk->fW0, 17) ^ rotr(blk->fW0, 19) ^ (blk->fW0 >> 10)); blk->fW2 = data[2] + (rotr(blk->fW0, 17) ^ rotr(blk->fW0, 19) ^ (blk->fW0 >> 10));
blk->fW3 = 0x11002000 + (rotr(blk->fW1, 17) ^ rotr(blk->fW1, 19) ^ (blk->fW1 >> 10)); blk->fW3 = 0x11002000 + (rotr(blk->fW1, 17) ^ rotr(blk->fW1, 19) ^ (blk->fW1 >> 10));
blk->fW15 = 0x00000280 + (rotr(blk->fW0, 7) ^ rotr(blk->fW0, 18) ^ (blk->fW0 >> 3)); blk->fW15 = 0x00000280 + (rotr(blk->fW0, 7) ^ rotr(blk->fW0, 18) ^ (blk->fW0 >> 3));
@ -101,6 +111,9 @@ void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data) {
blk->PreVal4 = blk->fcty_e = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + 0xe9b5dba5; blk->PreVal4 = blk->fcty_e = E + (rotr(B, 6) ^ rotr(B, 11) ^ rotr(B, 25)) + (D ^ (B & (C ^ D))) + 0xe9b5dba5;
blk->T1 = blk->fcty_e2 = (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G))); blk->T1 = blk->fcty_e2 = (rotr(F, 2) ^ rotr(F, 13) ^ rotr(F, 22)) + ((F & G) | (H & (F | G)));
blk->PreVal4addT1 = blk->PreVal4 + blk->T1;
blk->T1substate0 = state[0] - blk->T1;
} }
#define P(t) (W[(t)&0xF] = W[(t-16)&0xF] + (rotate(W[(t-15)&0xF], 25) ^ rotate(W[(t-15)&0xF], 14) ^ (W[(t-15)&0xF] >> 3)) + W[(t-7)&0xF] + (rotate(W[(t-2)&0xF], 15) ^ rotate(W[(t-2)&0xF], 13) ^ (W[(t-2)&0xF] >> 10))) #define P(t) (W[(t)&0xF] = W[(t-16)&0xF] + (rotate(W[(t-15)&0xF], 25) ^ rotate(W[(t-15)&0xF], 14) ^ (W[(t-15)&0xF] >> 3)) + W[(t-7)&0xF] + (rotate(W[(t-2)&0xF], 15) ^ rotate(W[(t-2)&0xF], 13) ^ (W[(t-2)&0xF] >> 10)))

96
main.c
View File

@ -225,6 +225,10 @@ static char current_block[37];
static char datestamp[40]; static char datestamp[40];
static char blockdate[40]; static char blockdate[40];
static char *opt_kernel = NULL;
enum cl_kernel chosen_kernel;
struct sigaction termhandler, inthandler; struct sigaction termhandler, inthandler;
struct thread_q *getq; struct thread_q *getq;
@ -503,6 +507,9 @@ static struct opt_table opt_config_table[] = {
OPT_WITH_ARG("--intensity|-I", OPT_WITH_ARG("--intensity|-I",
forced_int_1010, opt_show_intval, &scan_intensity, forced_int_1010, opt_show_intval, &scan_intensity,
"Intensity of GPU scanning (-10 -> 10, default: dynamic to maintain desktop interactivity)"), "Intensity of GPU scanning (-10 -> 10, default: dynamic to maintain desktop interactivity)"),
OPT_WITH_ARG("--kernel|-k",
opt_set_charp, NULL, &opt_kernel,
"Select kernel to use (poclbm or phatk - default: auto)"),
#endif #endif
OPT_WITHOUT_ARG("--load-balance", OPT_WITHOUT_ARG("--load-balance",
set_loadbalance, &pool_strategy, set_loadbalance, &pool_strategy,
@ -2391,7 +2398,7 @@ enum {
#ifdef HAVE_OPENCL #ifdef HAVE_OPENCL
static _clState *clStates[16]; static _clState *clStates[16];
static inline cl_int queue_kernel_parameters(_clState *clState, dev_blk_ctx *blk) static cl_int queue_poclbm_kernel(_clState *clState, dev_blk_ctx *blk)
{ {
cl_kernel *kernel = &clState->kernel; cl_kernel *kernel = &clState->kernel;
cl_int status = 0; cl_int status = 0;
@ -2413,24 +2420,51 @@ static inline cl_int queue_kernel_parameters(_clState *clState, dev_blk_ctx *blk
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->cty_h); status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->cty_h);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->nonce); status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->nonce);
if (clState->hasBitAlign == true) { status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW0);
/* Parameters for phatk kernel */ status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW1);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->W2); status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW2);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->W16); status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW3);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->W17); status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW15);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->PreVal4); status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW01r);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->T1); status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fcty_e);
} else { status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fcty_e2);
/* Parameters for poclbm kernel */
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW0); status |= clSetKernelArg(*kernel, num++, sizeof(clState->outputBuffer),
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW1); (void *)&clState->outputBuffer);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW2);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW3); return status;
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW15); }
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fW01r);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fcty_e); static cl_int queue_phatk_kernel(_clState *clState, dev_blk_ctx *blk)
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->fcty_e2); {
} cl_kernel *kernel = &clState->kernel;
cl_int status = 0;
int num = 0;
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_a);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_b);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_c);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_d);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_e);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_f);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_g);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->ctx_h);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->cty_b);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->cty_c);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->C1addK5);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->D1A);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->cty_f);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->cty_g);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->cty_h);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->nonce);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->W2A);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->W16);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->W17);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->W17_2);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->PreVal4addT1);
status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->T1substate0);
status |= clSetKernelArg(*kernel, num++, sizeof(clState->outputBuffer), status |= clSetKernelArg(*kernel, num++, sizeof(clState->outputBuffer),
(void *)&clState->outputBuffer); (void *)&clState->outputBuffer);
@ -2450,6 +2484,8 @@ static void set_threads_hashes(unsigned int vectors, unsigned int *threads,
static void *gpuminer_thread(void *userdata) static void *gpuminer_thread(void *userdata)
{ {
cl_int (*queue_kernel_parameters)(_clState *, dev_blk_ctx *);
const unsigned long cycle = opt_log_interval / 5 ? : 1; const unsigned long cycle = opt_log_interval / 5 ? : 1;
struct timeval tv_start, tv_end, diff, tv_workstart; struct timeval tv_start, tv_end, diff, tv_workstart;
struct thr_info *mythr = userdata; struct thr_info *mythr = userdata;
@ -2477,6 +2513,16 @@ static void *gpuminer_thread(void *userdata)
bool requested = true; bool requested = true;
uint32_t total_hashes = 0, hash_div = 1; uint32_t total_hashes = 0, hash_div = 1;
switch (chosen_kernel) {
case KL_POCLBM:
queue_kernel_parameters = &queue_poclbm_kernel;
break;
case KL_PHATK:
default:
queue_kernel_parameters = &queue_phatk_kernel;
break;
}
if (opt_dynamic) { if (opt_dynamic) {
/* Minimise impact on desktop if we want dynamic mode */ /* Minimise impact on desktop if we want dynamic mode */
setpriority(PRIO_PROCESS, 0, 19); setpriority(PRIO_PROCESS, 0, 19);
@ -3217,6 +3263,16 @@ int main (int argc, char *argv[])
if (argc != 1) if (argc != 1)
quit(1, "Unexpected extra commandline arguments"); quit(1, "Unexpected extra commandline arguments");
if (opt_kernel) {
if (strcmp(opt_kernel, "poclbm") && strcmp(opt_kernel, "phatk"))
quit(1, "Invalid kernel name specified - must be poclbm or phatk");
if (!strcmp(opt_kernel, "poclbm"))
chosen_kernel = KL_POCLBM;
else
chosen_kernel = KL_PHATK;
} else
chosen_kernel = KL_NONE;
if (total_devices) { if (total_devices) {
if (total_devices > nDevs) if (total_devices > nDevs)
quit(1, "More devices specified than exist"); quit(1, "More devices specified than exist");
@ -3455,7 +3511,7 @@ int main (int argc, char *argv[])
quit(1, "wakeup thread create failed"); quit(1, "wakeup thread create failed");
/* Now that everything's ready put enough work in the queue */ /* Now that everything's ready put enough work in the queue */
for (i = 0; i < opt_queue + mining_threads; i++) { for (i = 0; i < mining_threads; i++) {
if (unlikely(!queue_request())) if (unlikely(!queue_request()))
quit(1, "Failed to queue_request in main"); quit(1, "Failed to queue_request in main");
if (!opt_quiet) if (!opt_quiet)

19
miner.h
View File

@ -255,15 +255,17 @@ extern pthread_mutex_t control_lock;
#ifdef HAVE_OPENCL #ifdef HAVE_OPENCL
typedef struct { typedef struct {
cl_uint ctx_a; cl_uint ctx_b; cl_uint ctx_c; cl_uint ctx_d; cl_uint ctx_a; cl_uint ctx_b; cl_uint ctx_c; cl_uint ctx_d;
cl_uint ctx_e; cl_uint ctx_f; cl_uint ctx_g; cl_uint ctx_h; cl_uint ctx_e; cl_uint ctx_f; cl_uint ctx_g; cl_uint ctx_h;
cl_uint cty_a; cl_uint cty_b; cl_uint cty_c; cl_uint cty_d; cl_uint cty_a; cl_uint cty_b; cl_uint cty_c; cl_uint cty_d;
cl_uint cty_e; cl_uint cty_f; cl_uint cty_g; cl_uint cty_h; cl_uint cty_e; cl_uint cty_f; cl_uint cty_g; cl_uint cty_h;
cl_uint merkle; cl_uint ntime; cl_uint nbits; cl_uint nonce; cl_uint merkle; cl_uint ntime; cl_uint nbits; cl_uint nonce;
cl_uint fW0; cl_uint fW1; cl_uint fW2; cl_uint fW3; cl_uint fW15; cl_uint fW0; cl_uint fW1; cl_uint fW2; cl_uint fW3; cl_uint fW15;
cl_uint fW01r; cl_uint fcty_e; cl_uint fcty_e2; cl_uint fW01r; cl_uint fcty_e; cl_uint fcty_e2;
cl_uint W16; cl_uint W17; cl_uint W2; cl_uint W16; cl_uint W17; cl_uint W2;
cl_uint PreVal4; cl_uint T1; cl_uint PreVal4; cl_uint T1;
cl_uint C1addK5; cl_uint D1A; cl_uint W2A; cl_uint W17_2;
cl_uint PreVal4addT1; cl_uint T1substate0;
} dev_blk_ctx; } dev_blk_ctx;
#else #else
typedef struct { typedef struct {
@ -312,6 +314,12 @@ struct work {
struct pool *pool; struct pool *pool;
}; };
enum cl_kernel {
KL_NONE,
KL_POCLBM,
KL_PHATK,
};
bool submit_nonce(struct thr_info *thr, struct work *work, uint32_t nonce); bool submit_nonce(struct thr_info *thr, struct work *work, uint32_t nonce);
extern void kill_work(void); extern void kill_work(void);
@ -327,5 +335,6 @@ extern void tq_thaw(struct thread_q *tq);
extern bool test_and_set(bool *var); extern bool test_and_set(bool *var);
extern bool test_and_clear(bool *var); extern bool test_and_clear(bool *var);
extern bool successful_connect; extern bool successful_connect;
extern enum cl_kernel chosen_kernel;
#endif /* __MINER_H__ */ #endif /* __MINER_H__ */

31
ocl.c
View File

@ -340,10 +340,25 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
char binaryfilename[255]; char binaryfilename[255];
char numbuf[10]; char numbuf[10];
char filename[16]; char filename[16];
if (clState->hasBitAlign)
strcpy(filename, "phatk110714.cl"); if (chosen_kernel == KL_NONE) {
else if (clState->hasBitAlign)
strcpy(filename, "poclbm110717.cl"); chosen_kernel = KL_PHATK;
else
chosen_kernel = KL_POCLBM;
}
switch (chosen_kernel) {
case KL_POCLBM:
strcpy(filename, "poclbm110717.cl");
strcpy(binaryfilename, "poclbm110717");
break;
case KL_PHATK:
strcpy(filename, "phatk110722.cl");
strcpy(binaryfilename, "phatk110722");
break;
}
FILE *binaryfile; FILE *binaryfile;
size_t *binary_sizes; size_t *binary_sizes;
char **binaries; char **binaries;
@ -368,12 +383,10 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
return NULL; return NULL;
} }
strcpy(binaryfilename, name); strcat(binaryfilename, name);
if (clState->hasBitAlign) { if (clState->hasBitAlign)
strcat(binaryfilename, "phatk110714");
strcat(binaryfilename, "bitalign"); strcat(binaryfilename, "bitalign");
} else
strcat(binaryfilename, "poclbm110717");
strcat(binaryfilename, "v"); strcat(binaryfilename, "v");
sprintf(numbuf, "%d", clState->preferred_vwidth); sprintf(numbuf, "%d", clState->preferred_vwidth);
strcat(binaryfilename, numbuf); strcat(binaryfilename, numbuf);

830
phatk110714.cl → phatk110722.cl
View File

@ -1,411 +1,419 @@
// This file is taken and modified from the public-domain poclbm project, and // This file is taken and modified from the public-domain poclbm project, and
// we have therefore decided to keep it public-domain in Phoenix. // we have therefore decided to keep it public-domain in Phoenix.
// 2011-07-11: further modified by Diapolo and still public-domain // 2011-07-12: further modified by Diapolo and still public-domain
// -ck version to be compatible with cgminer // -ck version to be compatible with cgminer
// 2011-07-14: shorter code // 2011-07-14: shorter code
#define VECTORSX #define VECTORSX
#ifdef VECTORS4 #ifdef VECTORS4
typedef uint4 u; typedef uint4 u;
#elif defined VECTORS2 #elif defined VECTORS2
typedef uint2 u; typedef uint2 u;
#else #else
typedef uint u; typedef uint u;
#endif #endif
__constant uint K[64] = { __constant uint K[64] = {
0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5, 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174, 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da, 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967, 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85, 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070, 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3, 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2
}; };
// H[6] = 0x08909ae5U + 0xb0edbdd0 + K[0] == 0xfc08884d // H[6] = 0x08909ae5U + 0xb0edbdd0 + K[0] == 0xfc08884d
// H[7] = -0x5be0cd19 - (0x90befffa) K[60] == -0xec9fcd13 // H[7] = -0x5be0cd19 - (0x90befffa) K[60] == -0xec9fcd13
__constant uint H[8] = { __constant uint H[8] = {
0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0xfc08884d, 0xec9fcd13 0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0xfc08884d, 0xec9fcd13
}; };
// L = 0xa54ff53a + 0xb0edbdd0 + K[0] == 0x198c7e2a2 // L = 0xa54ff53a + 0xb0edbdd0 + K[0] == 0x198c7e2a2
__constant ulong L = 0x198c7e2a2; __constant ulong L = 0x198c7e2a2;
#define BFI_INTX #define BFI_INTX
#define BITALIGNX #define BITALIGNX
#ifdef BITALIGN #define O 15
#pragma OPENCL EXTENSION cl_amd_media_ops : enable
#define rot(x, y) amd_bitalign(x, x, (u)(32 - y)) #ifdef BITALIGN
#else #pragma OPENCL EXTENSION cl_amd_media_ops : enable
#define rot(x, y) rotate(x, (u)y) #define rot(x, y) amd_bitalign(x, x, (u)(32 - y))
#endif #else
#define rot(x, y) rotate(x, (u)y)
#ifdef BFI_INT #endif
#define Ch(x, y, z) amd_bytealign(x, y, z)
#else #ifdef BFI_INT
#define Ch(x, y, z) bitselect(z, y, x) #define Ch(x, y, z) amd_bytealign(x, y, z)
#endif #else
#define Ch(x, y, z) bitselect(z, y, x)
// Ma now uses the Ch function, if BFI_INT is enabled, the optimized Ch version is used #endif
#define Ma(x, y, z) Ch((z ^ x), y, x)
// Ma now uses the Ch function, if BFI_INT is enabled, the optimized Ch version is used
// Various intermediate calculations for each SHA round #define Ma(x, y, z) Ch((z ^ x), y, x)
#define s0(n) (rot(Vals[(128 - n) % 8], 30) ^ rot(Vals[(128 - n) % 8], 19) ^ rot(Vals[(128 - n) % 8], 10))
#define s1(n) (rot(Vals[(132 - n) % 8], 26) ^ rot(Vals[(132 - n) % 8], 21) ^ rot(Vals[(132 - n) % 8], 7)) // Various intermediate calculations for each SHA round
#define ch(n) (Ch(Vals[(132 - n) % 8], Vals[(133 - n) % 8], Vals[(134 - n) % 8])) #define s0(n) (rot(Vals[(128 - n) % 8], 30) ^ rot(Vals[(128 - n) % 8], 19) ^ rot(Vals[(128 - n) % 8], 10))
#define ma(n) (Ma(Vals[(129 - n) % 8], Vals[(130 - n) % 8], Vals[(128 - n) % 8])) #define s1(n) (rot(Vals[(132 - n) % 8], 26) ^ rot(Vals[(132 - n) % 8], 21) ^ rot(Vals[(132 - n) % 8], 7))
#define t1(n) (K[n % 64] + Vals[(135 - n) % 8] + W[n] + s1(n) + ch(n)) #define ch(n) (Ch(Vals[(132 - n) % 8], Vals[(133 - n) % 8], Vals[(134 - n) % 8]))
#define ma(n) (Ma(Vals[(129 - n) % 8], Vals[(130 - n) % 8], Vals[(128 - n) % 8]))
// intermediate W calculations #define t1(n) (K[n % 64] + Vals[(135 - n) % 8] + W[n - O] + s1(n) + ch(n))
#define P1(x) (rot(W[x - 2], 15) ^ rot(W[x - 2], 13) ^ (W[x - 2] >> 10U)) #define t1_no_W(n) (K[n % 64] + Vals[(135 - n) % 8] + s1(n) + ch(n))
#define P2(x) (rot(W[x - 15], 25) ^ rot(W[x - 15], 14) ^ (W[x - 15] >> 3U))
#define P3(x) W[x - 7] // intermediate W calculations
#define P4(x) W[x - 16] #define P1(x) (rot(W[x - 2 - O], 15) ^ rot(W[x - 2 - O], 13) ^ (W[x - 2 - O] >> 10U))
#define P2(x) (rot(W[x - 15 - O], 25) ^ rot(W[x - 15 - O], 14) ^ (W[x - 15 - O] >> 3U))
// full W calculation #define P3(x) W[x - 7 - O]
#define W(x) (W[x] = P4(x) + P3(x) + P2(x) + P1(x)) #define P4(x) W[x - 16 - O]
// SHA round without W calc // full W calculation
#define sharound(n) { Vals[(131 - n) % 8] += t1(n); Vals[(135 - n) % 8] = t1(n) + s0(n) + ma(n); } #define W(x) (W[x - O] = P4(x) + P3(x) + P2(x) + P1(x))
__kernel void search( const uint state0, const uint state1, const uint state2, const uint state3, // SHA round without W calc
const uint state4, const uint state5, const uint state6, const uint state7, #define sharound(n) { Vals[(131 - n) % 8] += t1(n); Vals[(135 - n) % 8] = t1(n) + s0(n) + ma(n); }
const uint B1, const uint C1, const uint D1, #define sharound_no_W(n) { Vals[(131 - n) % 8] += t1_no_W(n); Vals[(135 - n) % 8] = t1_no_W(n) + s0(n) + ma(n); }
const uint F1, const uint G1, const uint H1,
const uint base, __kernel void search( const uint state0, const uint state1, const uint state2, const uint state3,
const uint W2, const uint state4, const uint state5, const uint state6, const uint state7,
const uint W16, const uint W17, const uint B1, const uint C1, const uint C1addK5, const uint D1,
const uint PreVal4, const uint T1, const uint F1, const uint G1, const uint H1,
__global uint * output) const uint base,
{ const uint W2,
u W[124]; const uint W16, const uint W17, const uint W17_2,
u Vals[8]; const uint PreVal4addT1, const uint T1substate0,
__global uint * output)
Vals[1] = B1; {
Vals[2] = C1; u W[124 - O];
Vals[5] = F1; u Vals[8];
Vals[6] = G1; #ifdef VECTORS4
u W_3 = base + (get_global_id(0) << 2) + (uint4)(0, 1, 2, 3);
W[2] = W2; #elif defined VECTORS2
#ifdef VECTORS4 u W_3 = base + (get_global_id(0) << 1) + (uint2)(0, 1);
Vals[4] = (W[3] = base + (get_global_id(0) << 2) + (uint4)(0, 1, 2, 3)) + PreVal4; #else
#elif defined VECTORS2 u W_3 = base + get_global_id(0);
Vals[4] = (W[3] = base + (get_global_id(0) << 1) + (uint2)(0, 1)) + PreVal4; #endif
#else u Temp;
Vals[4] = (W[3] = base + get_global_id(0)) + PreVal4;
#endif Vals[0] = W_3 + PreVal4addT1 + T1substate0;
// used in: P2(19) == 285220864 (0x11002000), P4(20) Vals[1] = B1;
W[4] = 0x80000000U; Vals[2] = C1;
// P1(x) is 0 for x == 7, 8, 9, 10, 11, 12, 13, 14, 15, 16
// P2(x) is 0 for x == 20, 21, 22, 23, 24, 25, 26, 27, 28, 29 Vals[4] = W_3 + PreVal4addT1;
// P3(x) is 0 for x == 12, 13, 14, 15, 16, 17, 18, 19, 20, 21 Vals[5] = F1;
// P4(x) is 0 for x == 21, 22, 23, 24, 25, 26, 27, 28, 29, 30 Vals[6] = G1;
// W[x] in sharound(x) is 0 for x == 5, 6, 7, 8, 9, 10, 11, 12, 13, 14
W[14] = W[13] = W[12] = W[11] = W[10] = W[9] = W[8] = W[7] = W[6] = W[5] = 0x00000000U; // used in: P2(19) == 285220864 (0x11002000), P4(20)
// used in: P2(30) == 10485845 (0xA00055), P3(22), P4(31) // W[4] = 0x80000000U;
// K[15] + W[15] == 0xc19bf174 + 0x00000280U = 0xc19bf3f4 // P1(x) is 0 for x == 7, 8, 9, 10, 11, 12, 13, 14, 15, 16
W[15] = 0x00000280U; // P2(x) is 0 for x == 20, 21, 22, 23, 24, 25, 26, 27, 28, 29
// P3(x) is 0 for x == 12, 13, 14, 15, 16, 17, 18, 19, 20, 21
W[16] = W16; // P4(x) is 0 for x == 21, 22, 23, 24, 25, 26, 27, 28, 29, 30
W[17] = W17; // W[x] in sharound(x) is 0 for x == 5, 6, 7, 8, 9, 10, 11, 12, 13, 14
// removed P3(18) from add because it is == 0 // W[14] = W[13] = W[12] = W[11] = W[10] = W[9] = W[8] = W[7] = W[6] = W[5] = 0x00000000U;
W[18] = P1(18) + P4(18) + P2(18); // used in: P2(30) == 10485845 (0xA00055), P3(22), P4(31)
// removed P3(19) from add because it is == 0 // K[15] + W[15] == 0xc19bf174 + 0x00000280U = 0xc19bf3f4
W[19] = (u)0x11002000 + P1(19) + P4(19); W[15 - O] = 0x00000280U;
// removed P2(20), P3(20) from add because it is == 0 W[16 - O] = W16;
W[20] = P1(20) + P4(20); W[17 - O] = W17;
W[21] = P1(21); W[18 - O] = W2 + (rot(W_3, 25) ^ rot(W_3, 14) ^ (W_3 >> 3U));
W[22] = P1(22) + P3(22); W[19 - O] = W_3 + W17_2;
W[23] = P1(23) + P3(23); W[20 - O] = (u)0x80000000U + P1(20);
W[24] = P1(24) + P3(24); W[21 - O] = P1(21);
W[25] = P1(25) + P3(25); W[22 - O] = P1(22) + P3(22);
W[26] = P1(26) + P3(26); W[23 - O] = P1(23) + P3(23);
W[27] = P1(27) + P3(27); W[24 - O] = P1(24) + P3(24);
W[28] = P1(28) + P3(28); W[25 - O] = P1(25) + P3(25);
W[29] = P1(29) + P3(29); W[26 - O] = P1(26) + P3(26);
W[30] = (u)0xA00055 + P1(30) + P3(30); W[27 - O] = P1(27) + P3(27);
W[28 - O] = P1(28) + P3(28);
// Round 3 W[29 - O] = P1(29) + P3(29);
Vals[0] = state0 + Vals[4]; W[30 - O] = (u)0xA00055 + P1(30) + P3(30);
Vals[4] += T1;
// Round 4
// Round 4 Temp = D1 + ch(4) + s1(4);
// K[4] + W[4] == 0x3956c25b + 0x80000000U = 0xb956c25b Vals[7] = Temp + H1;
Vals[7] = (Vals[3] = (u)0xb956c25b + D1 + s1(4) + ch(4)) + H1; Vals[3] = Temp + ma(4) + s0(4);
Vals[3] += s0(4) + ma(4);
// Round 5
// Round 5 Temp = C1addK5 + ch(5) + s1(5);
Vals[2] = K[5] + C1 + s1(5) + ch(5) + s0(5) + ma(5); Vals[6] = Temp + G1;
Vals[6] = K[5] + C1 + G1 + s1(5) + ch(5); Vals[2] = Temp + ma(5) + s0(5);
sharound(6); // W[6] to W[14] are 0, so no need to add them!
sharound(7); sharound_no_W(6);
sharound(8); sharound_no_W(7);
sharound(9); sharound_no_W(8);
sharound(10); sharound_no_W(9);
sharound(11); sharound_no_W(10);
sharound(12); sharound_no_W(11);
sharound(13); sharound_no_W(12);
sharound(14); sharound_no_W(13);
sharound(15); sharound_no_W(14);
sharound(16);
sharound(17); // #define sharound(n) { Vals[(131 - n) % 8] += t1(n); Vals[(135 - n) % 8] = t1(n) + s0(n) + ma(n); }
sharound(18); // #define t1(n) (K[n % 64] + Vals[(135 - n) % 8] + W[n] + s1(n) + ch(n))
sharound(19);
sharound(20); // Vals[(131 - 15) % 8] += (Vals[(135 - 15) % 8] = (u)0xc19bf3f4 + Vals[(135 - 15) % 8] + s1(15) + ch(15));
sharound(21); // Vals[(135 - 15) % 8] += s0(15) + ma(15);
sharound(22);
sharound(23); sharound(15);
sharound(24); sharound(16);
sharound(25); sharound(17);
sharound(26); sharound(18);
sharound(27); sharound(19);
sharound(28); sharound(20);
sharound(29); sharound(21);
sharound(30); sharound(22);
sharound(23);
W(31); sharound(24);
sharound(31); sharound(25);
W(32); sharound(26);
sharound(32); sharound(27);
W(33); sharound(28);
sharound(33); sharound(29);
W(34); sharound(30);
sharound(34);
W(35); W(31);
sharound(35); sharound(31);
W(36); W(32);
sharound(36); sharound(32);
W(37); W(33);
sharound(37); sharound(33);
W(38); W(34);
sharound(38); sharound(34);
W(39); W(35);
sharound(39); sharound(35);
W(40); W(36);
sharound(40); sharound(36);
W(41); W(37);
sharound(41); sharound(37);
W(42); W(38);
sharound(42); sharound(38);
W(43); W(39);
sharound(43); sharound(39);
W(44); W(40);
sharound(44); sharound(40);
W(45); W(41);
sharound(45); sharound(41);
W(46); W(42);
sharound(46); sharound(42);
W(47); W(43);
sharound(47); sharound(43);
W(48); W(44);
sharound(48); sharound(44);
W(49); W(45);
sharound(49); sharound(45);
W(50); W(46);
sharound(50); sharound(46);
W(51); W(47);
sharound(51); sharound(47);
W(52); W(48);
sharound(52); sharound(48);
W(53); W(49);
sharound(53); sharound(49);
W(54); W(50);
sharound(54); sharound(50);
W(55); W(51);
sharound(55); sharound(51);
W(56); W(52);
sharound(56); sharound(52);
W(57); W(53);
sharound(57); sharound(53);
W(58); W(54);
sharound(58); sharound(54);
W(59); W(55);
sharound(59); sharound(55);
W(60); W(56);
sharound(60); sharound(56);
W(61); W(57);
sharound(61); sharound(57);
W(62); W(58);
sharound(62); sharound(58);
W(63); W(59);
sharound(63); sharound(59);
W(60);
W[64] = state0 + Vals[0]; sharound(60);
W[65] = state1 + Vals[1]; W(61);
W[66] = state2 + Vals[2]; sharound(61);
W[67] = state3 + Vals[3]; W(62);
W[68] = state4 + Vals[4]; sharound(62);
W[69] = state5 + Vals[5]; W(63);
W[70] = state6 + Vals[6]; sharound(63);
W[71] = state7 + Vals[7];
// used in: P2(87) = 285220864 (0x11002000), P4(88) W[64 - O] = state0 + Vals[0];
// K[72] + W[72] == W[65 - O] = state1 + Vals[1];
W[72] = 0x80000000U; W[66 - O] = state2 + Vals[2];
// P1(x) is 0 for x == 75, 76, 77, 78, 79, 80 W[67 - O] = state3 + Vals[3];
// P2(x) is 0 for x == 88, 89, 90, 91, 92, 93 W[68 - O] = state4 + Vals[4];
// P3(x) is 0 for x == 80, 81, 82, 83, 84, 85 W[69 - O] = state5 + Vals[5];
// P4(x) is 0 for x == 89, 90, 91, 92, 93, 94 W[70 - O] = state6 + Vals[6];
// W[x] in sharound(x) is 0 for x == 73, 74, 75, 76, 77, 78 W[71 - O] = state7 + Vals[7];
W[78] = W[77] = W[76] = W[75] = W[74] = W[73] = 0x00000000U; // used in: P2(87) = 285220864 (0x11002000), P4(88)
// used in: P1(81) = 10485760 (0xA00000), P2(94) = 4194338 (0x400022), P3(86), P4(95) // K[72] + W[72] ==
// K[79] + W[79] == W[72 - O] = 0x80000000U;
W[79] = 0x00000100U; // P1(x) is 0 for x == 75, 76, 77, 78, 79, 80
// P2(x) is 0 for x == 88, 89, 90, 91, 92, 93
Vals[0] = H[0]; // P3(x) is 0 for x == 80, 81, 82, 83, 84, 85
Vals[1] = H[1]; // P4(x) is 0 for x == 89, 90, 91, 92, 93, 94
Vals[2] = H[2]; // W[x] in sharound(x) is 0 for x == 73, 74, 75, 76, 77, 78
Vals[3] = (u)L + W[64]; // W[78] = W[77] = W[76] = W[75] = W[74] = W[73] = 0x00000000U;
Vals[4] = H[3]; // used in: P1(81) = 10485760 (0xA00000), P2(94) = 4194338 (0x400022), P3(86), P4(95)
Vals[5] = H[4]; // K[79] + W[79] ==
Vals[6] = H[5]; W[79 - O] = 0x00000100U;
Vals[7] = H[6] + W[64];
Vals[0] = H[0];
sharound(65); Vals[1] = H[1];
sharound(66); Vals[2] = H[2];
sharound(67); Vals[3] = (u)L + W[64 - O];
sharound(68); Vals[4] = H[3];
sharound(69); Vals[5] = H[4];
sharound(70); Vals[6] = H[5];
sharound(71); Vals[7] = H[6] + W[64 - O];
sharound(72);
sharound(73); sharound(65);
sharound(74); sharound(66);
sharound(75); sharound(67);
sharound(76); sharound(68);
sharound(77); sharound(69);
sharound(78); sharound(70);
sharound(79); sharound(71);
sharound(72);
// removed P1(80), P3(80) from add because it is == 0
W[80] = P2(80) + P4(80); // W is also zero for these rounds
W[81] = (u)0xA00000 + P4(81) + P2(81); sharound_no_W(73);
W[82] = P4(82) + P2(82) + P1(82); sharound_no_W(74);
W[83] = P4(83) + P2(83) + P1(83); sharound_no_W(75);
W[84] = P4(84) + P2(84) + P1(84); sharound_no_W(76);
W[85] = P4(85) + P2(85) + P1(85); sharound_no_W(77);
W(86); sharound_no_W(78);
sharound(80); sharound(79);
sharound(81);
sharound(82); W[80 - O] = P2(80) + P4(80);
sharound(83); W[81 - O] = (u)0xA00000 + P4(81) + P2(81);
sharound(84); W[82 - O] = P4(82) + P2(82) + P1(82);
sharound(85); W[83 - O] = P4(83) + P2(83) + P1(83);
sharound(86); W[84 - O] = P4(84) + P2(84) + P1(84);
W[85 - O] = P4(85) + P2(85) + P1(85);
W[87] = (u)0x11002000 + P4(87) + P3(87) + P1(87); W(86);
sharound(87);
W[88] = P4(88) + P3(88) + P1(88); sharound(80);
sharound(88); sharound(81);
W[89] = P3(89) + P1(89); sharound(82);
sharound(89); sharound(83);
W[90] = P3(90) + P1(90); sharound(84);
sharound(90); sharound(85);
W[91] = P3(91) + P1(91); sharound(86);
sharound(91);
W[92] = P3(92) + P1(92); W[87 - O] = (u)0x11002000 + P4(87) + P3(87) + P1(87);
sharound(92); sharound(87);
// removed P2(93), P4(93) from add because it is == 0 W[88 - O] = (u)0x80000000U + P3(88) + P1(88);
W[93] = P3(93) + P1(93); sharound(88);
sharound(93); W[89 - O] = P3(89) + P1(89);
// removed P4(94) from add because it is == 0 sharound(89);
W[94] = (u)0x400022 + P3(94) + P1(94); W[90 - O] = P3(90) + P1(90);
sharound(94); sharound(90);
W[91 - O] = P3(91) + P1(91);
W(95); sharound(91);
sharound(95); W[92 - O] = P3(92) + P1(92);
W(96); sharound(92);
sharound(96); W[93 - O] = P3(93) + P1(93);
W(97); sharound(93);
sharound(97); W[94 - O] = (u)0x400022 + P3(94) + P1(94);
W(98); sharound(94);
sharound(98); W[95 - O] = (u)0x00000100U + P3(95) + P2(95) + P1(95);
W(99); sharound(95);
sharound(99);
W(100); W(96);
sharound(100); sharound(96);
W(101); W(97);
sharound(101); sharound(97);
W(102); W(98);
sharound(102); sharound(98);
W(103); W(99);
sharound(103); sharound(99);
W(104); W(100);
sharound(104); sharound(100);
W(105); W(101);
sharound(105); sharound(101);
W(106); W(102);
sharound(106); sharound(102);
W(107); W(103);
sharound(107); sharound(103);
W(108); W(104);
sharound(108); sharound(104);
W(109); W(105);
sharound(109); sharound(105);
W(110); W(106);
sharound(110); sharound(106);
W(111); W(107);
sharound(111); sharound(107);
W(112); W(108);
sharound(112); sharound(108);
W(113); W(109);
sharound(113); sharound(109);
W(114); W(110);
sharound(114); sharound(110);
W(115); W(111);
sharound(115); sharound(111);
W(116); W(112);
sharound(116); sharound(112);
W(117); W(113);
sharound(117); sharound(113);
W(118); W(114);
sharound(118); sharound(114);
W(119); W(115);
sharound(119); sharound(115);
W(120); W(116);
sharound(120); sharound(116);
W(121); W(117);
sharound(121); sharound(117);
W(122); W(118);
sharound(122); sharound(118);
W(123); W(119);
sharound(123); sharound(119);
W(120);
// Round 124 sharound(120);
Vals[7] += Vals[3] + P4(124) + P3(124) + P2(124) + P1(124) + s1(124) + ch(124); W(121);
sharound(121);
#define MAXBUFFERS (4095) W(122);
#define NFLAG (0xFFFUL) sharound(122);
W(123);
#if defined(VECTORS4) || defined(VECTORS2) sharound(123);
if (Vals[7].x == -H[7])
{ // Round 124
output[MAXBUFFERS] = output[NFLAG & W[3].x] = W[3].x; Vals[7] += Vals[3] + P4(124) + P3(124) + P2(124) + P1(124) + s1(124) + ch(124);
}
if (Vals[7].y == -H[7]) #define MAXBUFFERS (4095)
{ #define NFLAG (0xFFFUL)
output[MAXBUFFERS] = output[NFLAG & W[3].y] = W[3].y;
} #if defined(VECTORS4) || defined(VECTORS2)
#ifdef VECTORS4 if (Vals[7].x == -H[7])
if (Vals[7].z == -H[7]) {
{ output[MAXBUFFERS] = output[NFLAG & W_3.x] = W_3.x;
output[MAXBUFFERS] = output[NFLAG & W[3].z] = W[3].z; }
} if (Vals[7].y == -H[7])
if (Vals[7].w == -H[7]) {
{ output[MAXBUFFERS] = output[NFLAG & W_3.y] = W_3.y;
output[MAXBUFFERS] = output[NFLAG & W[3].w] = W[3].w; }
} #ifdef VECTORS4
#endif if (Vals[7].z == -H[7])
#else {
if (Vals[7] == -H[7]) output[MAXBUFFERS] = output[NFLAG & W_3.z] = W_3.z;
{ }
output[MAXBUFFERS] = output[NFLAG & W[3]] = W[3]; if (Vals[7].w == -H[7])
} {
#endif output[MAXBUFFERS] = output[NFLAG & W_3.w] = W_3.w;
} }
#endif
#else
if (Vals[7] == -H[7])
{
output[MAXBUFFERS] = output[NFLAG & W_3] = W_3;
}
#endif
}