GOSTSec
/
sgminer
mirror of https://github.com/GOSTSec/sgminer
1
0
Fork 0
Code Issues Releases Wiki Activity
Browse Source

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

nfactor-troky
Paul Sheppard 12 years ago
parent
736c62a91e b768758818
commit
3ef6db81cf
14 changed files with 214 additions and 182 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 11
      API-README
  2. 19
      NEWS
  3. 7
      api.c
  4. 105
      cgminer.c
  5. 2
      configure.ac
  6. 52
      diablo120823.cl
  7. 52
      diakgcn120823.cl
  8. 7
      miner.h
  9. 18
      miner.php
  10. 3
      ocl.c
  11. 52
      phatk120823.cl
  12. 39
      poclbm120823.cl
  13. 12
      scrypt120823.cl
  14. 17
      util.c

11
API-README
Unescape View File

@ -339,6 +339,17 @@ miner.php - an example web page to access the API
Feature Changelog for external applications using the API: Feature Changelog for external applications using the API:
API V1.18
Modified API commands:
'stats' - add 'Work Had Roll Time', 'Work Can Roll', 'Work Had Expire',
'Work Roll Time' to the pool stats
Modified API commands:
'config' - include 'ScanTime'
----------
API V1.17 (cgminer v2.7.1) API V1.17 (cgminer v2.7.1)
Added API commands: Added API commands:

19
NEWS
Unescape View File

@ -1,3 +1,22 @@
Version 2.7.4 - August 23, 2012
- Perform select_pool even when not lagging to allow it to switch back if needed
to the primary.
- Simplify macros in output kernels avoiding apparent loops and local variables.
- Carry the needed bool over the work command queue.
- Move the decision to queue further work upstream before threads are spawned
based on fine grained per-pool stats and increment the queued count immediately.
- Track queued and staged per pool once again for future use.
- OpenCL 1.0 does not have native atomic_add and extremely slow support with
atom_add so detect opencl1.0 and use a non-atomic workaround.
- Pools: add RollTime info to API 'stats' and 'Stats' button in miner.php
Version 2.7.3 - August 22, 2012
- Minimise the number of getwork threads we generate.
Version 2.7.2 - August 22, 2012 Version 2.7.2 - August 22, 2012
- Pick worksize 256 with Cypress if none is specified. - Pick worksize 256 with Cypress if none is specified.

7
api.c
Unescape View File

@ -166,7 +166,7 @@ static const char SEPARATOR = '|';
#define SEPSTR "|" #define SEPSTR "|"
static const char GPUSEP = ','; static const char GPUSEP = ',';
static const char *APIVERSION = "1.17"; static const char *APIVERSION = "1.18";
static const char *DEAD = "Dead"; static const char *DEAD = "Dead";
static const char *SICK = "Sick"; static const char *SICK = "Sick";
static const char *NOSTART = "NoStart"; static const char *NOSTART = "NoStart";
@ -1256,6 +1256,7 @@ static void minerconfig(__maybe_unused SOCKETTYPE c, __maybe_unused char *param,
root = api_add_const(root, "Device Code", DEVICECODE, false); root = api_add_const(root, "Device Code", DEVICECODE, false);
root = api_add_const(root, "OS", OSINFO, false); root = api_add_const(root, "OS", OSINFO, false);
root = api_add_bool(root, "Failover-Only", &opt_fail_only, false); root = api_add_bool(root, "Failover-Only", &opt_fail_only, false);
root = api_add_int(root, "ScanTime", &opt_scantime, false);
root = print_data(root, buf, isjson); root = print_data(root, buf, isjson);
if (isjson) if (isjson)
@ -2676,6 +2677,10 @@ static int itemstats(int i, char *id, struct cgminer_stats *stats, struct cgmine
root = api_add_timeval(root, "Pool Max", &(pool_stats->getwork_wait_max), false); root = api_add_timeval(root, "Pool Max", &(pool_stats->getwork_wait_max), false);
root = api_add_timeval(root, "Pool Min", &(pool_stats->getwork_wait_min), false); root = api_add_timeval(root, "Pool Min", &(pool_stats->getwork_wait_min), false);
root = api_add_double(root, "Pool Av", &(pool_stats->getwork_wait_rolling), false); root = api_add_double(root, "Pool Av", &(pool_stats->getwork_wait_rolling), false);
root = api_add_bool(root, "Work Had Roll Time", &(pool_stats->hadrolltime), false);
root = api_add_bool(root, "Work Can Roll", &(pool_stats->canroll), false);
root = api_add_bool(root, "Work Had Expire", &(pool_stats->hadexpire), false);
root = api_add_uint32(root, "Work Roll Time", &(pool_stats->rolltime), false);
} }
if (extra) if (extra)

105
cgminer.c
Unescape View File

@ -70,7 +70,7 @@ struct workio_cmd {
enum workio_commands cmd; enum workio_commands cmd;
struct thr_info *thr; struct thr_info *thr;
struct work *work; struct work *work;
bool needed; struct pool *pool;
}; };
struct strategies strategies[] = { struct strategies strategies[] = {
@ -1955,11 +1955,9 @@ static inline struct pool *select_pool(bool lagging)
if (pool_strategy == POOL_BALANCE) if (pool_strategy == POOL_BALANCE)
return select_balanced(cp); return select_balanced(cp);
if (pool_strategy != POOL_LOADBALANCE && (!lagging || opt_fail_only)) { if (pool_strategy != POOL_LOADBALANCE && (!lagging || opt_fail_only))
if (cp->prio != 0) pool = cp;
switch_pools(NULL); else
pool = current_pool();
} else
pool = NULL; pool = NULL;
while (!pool) { while (!pool) {
@ -2251,17 +2249,19 @@ static void push_curl_entry(struct curl_ent *ce, struct pool *pool)
/* This is overkill, but at least we'll know accurately how much work is /* This is overkill, but at least we'll know accurately how much work is
* queued to prevent ever being left without work */ * queued to prevent ever being left without work */
static void inc_queued(void) static void inc_queued(struct pool *pool)
{ {
mutex_lock(&qd_lock); mutex_lock(&qd_lock);
total_queued++; total_queued++;
pool->queued++;
mutex_unlock(&qd_lock); mutex_unlock(&qd_lock);
} }
static void dec_queued(void) static void dec_queued(struct pool *pool)
{ {
mutex_lock(&qd_lock); mutex_lock(&qd_lock);
total_queued--; total_queued--;
pool->queued--;
mutex_unlock(&qd_lock); mutex_unlock(&qd_lock);
} }
@ -2380,68 +2380,46 @@ out:
return cloned; return cloned;
} }
static bool queue_request(void);
static void *get_work_thread(void *userdata) static void *get_work_thread(void *userdata)
{ {
struct workio_cmd *wc = (struct workio_cmd *)userdata; struct workio_cmd *wc = (struct workio_cmd *)userdata;
int ts, tq, maxq = opt_queue + mining_threads;
struct pool *pool = current_pool(); struct pool *pool = current_pool();
struct work *ret_work= NULL; struct work *ret_work= NULL;
struct curl_ent *ce = NULL; struct curl_ent *ce = NULL;
bool lagging = false;
pthread_detach(pthread_self()); pthread_detach(pthread_self());
applog(LOG_DEBUG, "Creating extra get work thread"); applog(LOG_DEBUG, "Creating extra get work thread");
retry: pool = wc->pool;
tq = global_queued();
ts = total_staged();
if (ts >= maxq)
goto out;
if (ts >= opt_queue && tq >= maxq) if (clone_available()) {
goto out; dec_queued(pool);
if (clone_available())
goto out; goto out;
}
ret_work = make_work(); ret_work = make_work();
if (wc->thr) ret_work->thr = NULL;
ret_work->thr = wc->thr;
else
ret_work->thr = NULL;
if (opt_benchmark) { if (opt_benchmark) {
get_benchmark_work(ret_work); get_benchmark_work(ret_work);
ret_work->queued = true; ret_work->queued = true;
} else { } else {
ret_work->pool = wc->pool;
if (!ts)
lagging = true;
pool = ret_work->pool = select_pool(lagging);
inc_queued();
if (!ce) if (!ce)
ce = pop_curl_entry(pool); ce = pop_curl_entry(pool);
/* Check that we haven't staged work via other threads while
* waiting for a curl entry */
if (total_staged() >= maxq) {
dec_queued();
free_work(ret_work);
goto out;
}
/* obtain new work from bitcoin via JSON-RPC */ /* obtain new work from bitcoin via JSON-RPC */
if (!get_upstream_work(ret_work, ce->curl)) { if (!get_upstream_work(ret_work, ce->curl)) {
/* pause, then restart work-request loop */ /* pause, then restart work-request loop */
applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying"); applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying");
lagging = true; dec_queued(pool);
dec_queued(); queue_request();
free_work(ret_work); free_work(ret_work);
goto retry; goto out;
} }
ret_work->queued = true; ret_work->queued = true;
@ -2620,8 +2598,6 @@ static struct pool *priority_pool(int choice)
return ret; return ret;
} }
static bool queue_request(struct thr_info *thr, bool needed);
void switch_pools(struct pool *selected) void switch_pools(struct pool *selected)
{ {
struct pool *pool, *last_pool; struct pool *pool, *last_pool;
@ -2696,8 +2672,6 @@ void switch_pools(struct pool *selected)
mutex_lock(&lp_lock); mutex_lock(&lp_lock);
pthread_cond_broadcast(&lp_cond); pthread_cond_broadcast(&lp_cond);
mutex_unlock(&lp_lock); mutex_unlock(&lp_lock);
queue_request(NULL, false);
} }
static void discard_work(struct work *work) static void discard_work(struct work *work)
@ -2721,6 +2695,7 @@ static void discard_stale(void)
HASH_ITER(hh, staged_work, work, tmp) { HASH_ITER(hh, staged_work, work, tmp) {
if (stale_work(work, false)) { if (stale_work(work, false)) {
HASH_DEL(staged_work, work); HASH_DEL(staged_work, work);
work->pool->staged--;
discard_work(work); discard_work(work);
stale++; stale++;
} }
@ -2730,7 +2705,7 @@ static void discard_stale(void)
if (stale) { if (stale) {
applog(LOG_DEBUG, "Discarded %d stales that didn't match current hash", stale); applog(LOG_DEBUG, "Discarded %d stales that didn't match current hash", stale);
while (stale-- > 0) while (stale-- > 0)
queue_request(NULL, false); queue_request();
} }
} }
@ -2927,9 +2902,11 @@ static bool hash_push(struct work *work)
pthread_cond_signal(&getq->cond); pthread_cond_signal(&getq->cond);
mutex_unlock(stgd_lock); mutex_unlock(stgd_lock);
work->pool->staged++;
if (work->queued) { if (work->queued) {
work->queued = false; work->queued = false;
dec_queued(); dec_queued(work->pool);
} }
return rc; return rc;
@ -3919,9 +3896,28 @@ static void pool_resus(struct pool *pool)
switch_pools(NULL); switch_pools(NULL);
} }
static bool queue_request(struct thr_info *thr, bool needed) static bool queue_request(void)
{ {
int ts, tq, maxq = opt_queue + mining_threads;
struct pool *pool, *cp;
struct workio_cmd *wc; struct workio_cmd *wc;
bool lagging;
ts = total_staged();
tq = global_queued();
if (ts && ts + tq >= maxq)
return true;
cp = current_pool();
lagging = !opt_fail_only && cp->lagging && !ts && cp->queued >= maxq;
if (!lagging && cp->staged + cp->queued >= maxq)
return true;
pool = select_pool(lagging);
if (pool->staged + pool->queued >= maxq)
return true;
inc_queued(pool);
/* fill out work request message */ /* fill out work request message */
wc = calloc(1, sizeof(*wc)); wc = calloc(1, sizeof(*wc));
@ -3931,8 +3927,7 @@ static bool queue_request(struct thr_info *thr, bool needed)
} }
wc->cmd = WC_GET_WORK; wc->cmd = WC_GET_WORK;
wc->thr = thr; wc->pool = pool;
wc->needed = needed;
applog(LOG_DEBUG, "Queueing getwork request to work thread"); applog(LOG_DEBUG, "Queueing getwork request to work thread");
@ -3967,12 +3962,13 @@ static struct work *hash_pop(const struct timespec *abstime)
} else } else
work = staged_work; work = staged_work;
HASH_DEL(staged_work, work); HASH_DEL(staged_work, work);
work->pool->staged--;
if (work_rollable(work)) if (work_rollable(work))
staged_rollable--; staged_rollable--;
} }
mutex_unlock(stgd_lock); mutex_unlock(stgd_lock);
queue_request(NULL, false); queue_request();
return work; return work;
} }
@ -4089,7 +4085,10 @@ retry:
pool = work_heap->pool; pool = work_heap->pool;
/* If we make it here we have succeeded in getting fresh work */ /* If we make it here we have succeeded in getting fresh work */
if (!work_heap->mined) { if (!work_heap->mined) {
pool_tclear(pool, &pool->lagging); /* Only clear the lagging flag if we are staging them at a
* rate faster then we're using them */
if (pool->lagging && total_staged())
pool_tclear(pool, &pool->lagging);
if (pool_tclear(pool, &pool->idle)) if (pool_tclear(pool, &pool->idle))
pool_resus(pool); pool_resus(pool);
} }
@ -5722,7 +5721,7 @@ begin_bench:
#endif #endif
for (i = 0; i < mining_threads + opt_queue; i++) for (i = 0; i < mining_threads + opt_queue; i++)
queue_request(NULL, false); queue_request();
/* main loop - simply wait for workio thread to exit. This is not the /* main loop - simply wait for workio thread to exit. This is not the
* normal exit path and only occurs should the workio_thread die * normal exit path and only occurs should the workio_thread die

2
configure.ac
Unescape View File

@ -2,7 +2,7 @@
##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--## ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
m4_define([v_maj], [2]) m4_define([v_maj], [2])
m4_define([v_min], [7]) m4_define([v_min], [7])
m4_define([v_mic], [2]) m4_define([v_mic], [4])
##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--## ##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
m4_define([v_ver], [v_maj.v_min.v_mic]) m4_define([v_ver], [v_maj.v_min.v_mic])
m4_define([lt_rev], m4_eval(v_maj + v_min)) m4_define([lt_rev], m4_eval(v_maj + v_min))

52
diablo120823.cl
Unescape View File

@ -1244,48 +1244,36 @@ void search(
#define FOUND (0x0F) #define FOUND (0x0F)
#if defined(OCL1)
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
#else
#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
#endif
#if defined(VECTORS4) #if defined(VECTORS4)
bool result = any(ZA[924] == 0x136032EDU); bool result = any(ZA[924] == 0x136032EDU);
if (result) { if (result) {
uint found; if (ZA[924].x == 0x136032EDU)
SETFOUND(Znonce.x);
if (ZA[924].x == 0x136032EDU) { if (ZA[924].y == 0x136032EDU)
found = atomic_add(&output[FOUND], 1); SETFOUND(Znonce.y);
output[found] = Znonce.x; if (ZA[924].z == 0x136032EDU)
} SETFOUND(Znonce.z);
if (ZA[924].y == 0x136032EDU) { if (ZA[924].w == 0x136032EDU)
found = atomic_add(&output[FOUND], 1); SETFOUND(Znonce.w);
output[found] = Znonce.y;
}
if (ZA[924].z == 0x136032EDU) {
found = atomic_add(&output[FOUND], 1);
output[found] = Znonce.z;
}
if (ZA[924].w == 0x136032EDU) {
found = atomic_add(&output[FOUND], 1);
output[found] = Znonce.w;
}
} }
#elif defined(VECTORS2) #elif defined(VECTORS2)
bool result = any(ZA[924] == 0x136032EDU); bool result = any(ZA[924] == 0x136032EDU);
if (result) { if (result) {
uint found; if (ZA[924].x == 0x136032EDU)
SETFOUND(Znonce.x);
if (ZA[924].x == 0x136032EDU) { if (ZA[924].y == 0x136032EDU)
found = atomic_add(&output[FOUND], 1); SETFOUND(Znonce.y);
output[found] = Znonce.x;
}
if (ZA[924].y == 0x136032EDU) {
found = atomic_add(&output[FOUND], 1);
output[found] = Znonce.y;
}
} }
#else #else
if (ZA[924] == 0x136032EDU) { if (ZA[924] == 0x136032EDU)
uint found = atomic_add(&output[FOUND], 1); SETFOUND(Znonce);
output[found] = Znonce;
}
#endif #endif
} }

52
diakgcn120823.cl
Unescape View File

@ -573,44 +573,32 @@ __kernel
#define FOUND (0x0F) #define FOUND (0x0F)
#if defined(OCL1)
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
#else
#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
#endif
#ifdef VECTORS4 #ifdef VECTORS4
if ((V[7].x == 0x136032edU) ^ (V[7].y == 0x136032edU) ^ (V[7].z == 0x136032edU) ^ (V[7].w == 0x136032edU)) { if ((V[7].x == 0x136032edU) ^ (V[7].y == 0x136032edU) ^ (V[7].z == 0x136032edU) ^ (V[7].w == 0x136032edU)) {
uint found; if (V[7].x == 0x136032edU)
SETFOUND(nonce.x);
if (V[7].x == 0x136032edU) { if (V[7].y == 0x136032edU)
found = atomic_add(&output[FOUND], 1); SETFOUND(nonce.y);
output[found] = nonce.x; if (V[7].z == 0x136032edU)
} SETFOUND(nonce.z);
if (V[7].y == 0x136032edU) { if (V[7].w == 0x136032edU)
found = atomic_add(&output[FOUND], 1); SETFOUND(nonce.w);
output[found] = nonce.y;
}
if (V[7].z == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.z;
}
if (V[7].w == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.w;
}
} }
#elif defined VECTORS2 #elif defined VECTORS2
if ((V[7].x == 0x136032edU) + (V[7].y == 0x136032edU)) { if ((V[7].x == 0x136032edU) + (V[7].y == 0x136032edU)) {
uint found; if (V[7].x == 0x136032edU)
SETFOUND(nonce.x);
if (V[7].x == 0x136032edU) { if (V[7].y == 0x136032edU)
found = atomic_add(&output[FOUND], 1); SETFOUND(nonce.y);
output[found] = nonce.x;
}
if (V[7].y == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.y;
}
} }
#else #else
if (V[7] == 0x136032edU) { if (V[7] == 0x136032edU)
uint found = atomic_add(&output[FOUND], 1); SETFOUND(nonce);
output[found] = nonce;
}
#endif #endif
} }

7
miner.h
Unescape View File

@ -307,6 +307,10 @@ struct cgminer_pool_stats {
struct timeval getwork_wait_max; struct timeval getwork_wait_max;
struct timeval getwork_wait_min; struct timeval getwork_wait_min;
double getwork_wait_rolling; double getwork_wait_rolling;
bool hadrolltime;
bool canroll;
bool hadexpire;
uint32_t rolltime;
}; };
struct cgpu_info { struct cgpu_info {
@ -732,6 +736,9 @@ struct pool {
int solved; int solved;
int diff1; int diff1;
int queued;
int staged;
bool submit_fail; bool submit_fail;
bool idle; bool idle;
bool lagging; bool lagging;

18
miner.php
Unescape View File

@ -84,8 +84,24 @@ $mobilesum = array(
'DEVS+NOTIFY' => array('DEVS.MHS av', 'DEVS.Accepted', 'DEVS.Rejected', 'DEVS.Utility'), 'DEVS+NOTIFY' => array('DEVS.MHS av', 'DEVS.Accepted', 'DEVS.Rejected', 'DEVS.Utility'),
'POOL' => array('Accepted', 'Rejected')); 'POOL' => array('Accepted', 'Rejected'));
# #
$statspage = array(
'DATE' => null,
'RIGS' => null,
'SUMMARY' => array('Elapsed', 'MHS av', 'Found Blocks=Blks',
'Accepted', 'Rejected=Rej', 'Utility',
'Hardware Errors=HW Errs', 'Network Blocks=Net Blks',
'Work Utility'),
'COIN' => array('*'),
'STATS' => array('*'));
#
$statssum = array(
'SUMMARY' => array('MHS av', 'Found Blocks', 'Accepted',
'Rejected', 'Utility', 'Hardware Errors',
'Work Utility'));
#
# customsummarypages is an array of these Custom Summary Pages # customsummarypages is an array of these Custom Summary Pages
$customsummarypages = array('Mobile' => array($mobilepage, $mobilesum)); $customsummarypages = array('Mobile' => array($mobilepage, $mobilesum),
'Stats' => array($statspage, $statssum));
# #
$here = $_SERVER['PHP_SELF']; $here = $_SERVER['PHP_SELF'];
# #

3
ocl.c
Unescape View File

@ -659,6 +659,9 @@ build:
if (clState->goffset) if (clState->goffset)
strcat(CompilerOptions, " -D GOFFSET"); strcat(CompilerOptions, " -D GOFFSET");
if (!clState->hasOpenCL11plus)
strcat(CompilerOptions, " -D OCL1");
applog(LOG_DEBUG, "CompilerOptions: %s", CompilerOptions); applog(LOG_DEBUG, "CompilerOptions: %s", CompilerOptions);
status = clBuildProgram(clState->program, 1, &devices[gpu], CompilerOptions , NULL, NULL); status = clBuildProgram(clState->program, 1, &devices[gpu], CompilerOptions , NULL, NULL);
free(CompilerOptions); free(CompilerOptions);

52
phatk120823.cl
Unescape View File

@ -389,46 +389,34 @@ void search( const uint state0, const uint state1, const uint state2, const uint
#define FOUND (0x0F) #define FOUND (0x0F)
#if defined(OCL1)
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
#else
#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
#endif
#ifdef VECTORS4 #ifdef VECTORS4
bool result = W[117].x & W[117].y & W[117].z & W[117].w; bool result = W[117].x & W[117].y & W[117].z & W[117].w;
if (!result) { if (!result) {
uint found; if (!W[117].x)
SETFOUND(W[3].x);
if (!W[117].x) { if (!W[117].y)
found = atomic_add(&output[FOUND], 1); SETFOUND(W[3].y);
output[found] = W[3].x; if (!W[117].z)
} SETFOUND(W[3].z);
if (!W[117].y) { if (!W[117].w)
found = atomic_add(&output[FOUND], 1); SETFOUND(W[3].w);
output[found] = W[3].y;
}
if (!W[117].z) {
found = atomic_add(&output[FOUND], 1);
output[found] = W[3].z;
}
if (!W[117].w) {
found = atomic_add(&output[FOUND], 1);
output[found] = W[3].w;
}
} }
#elif defined VECTORS2 #elif defined VECTORS2
bool result = W[117].x & W[117].y; bool result = W[117].x & W[117].y;
if (!result) { if (!result) {
uint found; if (!W[117].x)
SETFOUND(W[3].x);
if (!W[117].x) { if (!W[117].y)
found = atomic_add(&output[FOUND], 1); SETFOUND(W[3].y);
output[found] = W[3].x;
}
if (!W[117].y) {
found = atomic_add(&output[FOUND], 1);
output[found] = W[3].y;
}
} }
#else #else
if (!W[117]) { if (!W[117])
uint found = atomic_add(&output[FOUND], 1); SETFOUND(W[3]);
output[found] = W[3];
}
#endif #endif
} }

39
poclbm120823.cl
Unescape View File

@ -1323,34 +1323,27 @@ Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
#define FOUND (0x0F) #define FOUND (0x0F)
#if defined(VECTORS2) || defined(VECTORS4) #if defined(OCL1)
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
#else
#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
#endif
#if defined(VECTORS2) || defined(VECTORS4)
if (any(Vals[2] == 0x136032edU)) { if (any(Vals[2] == 0x136032edU)) {
uint found; if (Vals[2].x == 0x136032edU)
SETFOUND(nonce.x);
if (Vals[2].x == 0x136032edU) { if (Vals[2].y == 0x136032edU)
found = atomic_add(&output[FOUND], 1); SETFOUND(nonce.y);
output[found] = nonce.x;
}
if (Vals[2].y == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.y;
}
#if defined(VECTORS4) #if defined(VECTORS4)
if (Vals[2].z == 0x136032edU) { if (Vals[2].z == 0x136032edU)
found = atomic_add(&output[FOUND], 1); SETFOUND(nonce.z);
output[found] = nonce.z; if (Vals[2].w == 0x136032edU)
} SETFOUND(nonce.w);
if (Vals[2].w == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.w;
}
#endif #endif
} }
#else #else
if (Vals[2] == 0x136032edU) { if (Vals[2] == 0x136032edU)
uint found = atomic_add(&output[FOUND], 1); SETFOUND(nonce);
output[found] = nonce;
}
#endif #endif
} }

12
scrypt120823.cl
Unescape View File

@ -684,6 +684,12 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
#define FOUND (0x0F) #define FOUND (0x0F)
#if defined(OCL1)
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
#else
#define SETFOUND(Xnonce) output[atomic_add(&output[FOUND], 1)] = Xnonce
#endif
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void search(__global const uint4 * restrict input, __kernel void search(__global const uint4 * restrict input,
volatile __global uint*restrict output, __global uint4*restrict padcache, volatile __global uint*restrict output, __global uint4*restrict padcache,
@ -721,10 +727,8 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(0x80000000U, 0U, 0U, 0U), (uint4)(0U, 0U, 0U, 0x300U)); SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(0x80000000U, 0U, 0U, 0U), (uint4)(0U, 0U, 0U, 0x300U));
bool result = (EndianSwap(ostate1.w) <= target); bool result = (EndianSwap(ostate1.w) <= target);
if (result) { if (result)
uint found = atomic_add(&output[FOUND], 1); SETFOUND(gid);
output[found] = gid;
}
} }
/*- /*-

17
util.c
Unescape View File

@ -58,6 +58,9 @@ struct header_info {
char *lp_path; char *lp_path;
int rolltime; int rolltime;
char *reason; char *reason;
bool hadrolltime;
bool canroll;
bool hadexpire;
}; };
struct tq_ent { struct tq_ent {
@ -157,14 +160,18 @@ static size_t resp_hdr_cb(void *ptr, size_t size, size_t nmemb, void *user_data)
applog(LOG_DEBUG, "HTTP hdr(%s): %s", key, val); applog(LOG_DEBUG, "HTTP hdr(%s): %s", key, val);
if (!strcasecmp("X-Roll-Ntime", key)) { if (!strcasecmp("X-Roll-Ntime", key)) {
hi->hadrolltime = true;
if (!strncasecmp("N", val, 1)) if (!strncasecmp("N", val, 1))
applog(LOG_DEBUG, "X-Roll-Ntime: N found"); applog(LOG_DEBUG, "X-Roll-Ntime: N found");
else { else {
hi->canroll = true;
/* Check to see if expire= is supported and if not, set /* Check to see if expire= is supported and if not, set
* the rolltime to the default scantime */ * the rolltime to the default scantime */
if (strlen(val) > 7 && !strncasecmp("expire=", val, 7)) if (strlen(val) > 7 && !strncasecmp("expire=", val, 7)) {
sscanf(val + 7, "%d", &hi->rolltime); sscanf(val + 7, "%d", &hi->rolltime);
else hi->hadexpire = true;
} else
hi->rolltime = opt_scantime; hi->rolltime = opt_scantime;
applog(LOG_DEBUG, "X-Roll-Ntime expiry set to %d", hi->rolltime); applog(LOG_DEBUG, "X-Roll-Ntime expiry set to %d", hi->rolltime);
} }
@ -258,7 +265,7 @@ json_t *json_rpc_call(CURL *curl, const char *url,
{ {
long timeout = longpoll ? (60 * 60) : 60; long timeout = longpoll ? (60 * 60) : 60;
struct data_buffer all_data = {NULL, 0}; struct data_buffer all_data = {NULL, 0};
struct header_info hi = {NULL, 0, NULL}; struct header_info hi = {NULL, 0, NULL, false, false, false};
char len_hdr[64], user_agent_hdr[128]; char len_hdr[64], user_agent_hdr[128];
char curl_err_str[CURL_ERROR_SIZE]; char curl_err_str[CURL_ERROR_SIZE];
struct curl_slist *headers = NULL; struct curl_slist *headers = NULL;
@ -388,6 +395,10 @@ json_t *json_rpc_call(CURL *curl, const char *url,
} }
*rolltime = hi.rolltime; *rolltime = hi.rolltime;
pool->cgminer_pool_stats.rolltime = hi.rolltime;
pool->cgminer_pool_stats.hadrolltime = hi.hadrolltime;
pool->cgminer_pool_stats.canroll = hi.canroll;
pool->cgminer_pool_stats.hadexpire = hi.hadexpire;
val = JSON_LOADS(all_data.buf, &err); val = JSON_LOADS(all_data.buf, &err);
if (!val) { if (!val) {

Write Preview
Loading…
Cancel
Save