GOSTSec/sgminer
1
0
Fork 0
mirror of https://github.com/GOSTSec/sgminer synced 2025-02-04 02:55:29 +00:00
Code Issues Releases Wiki Activity

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

Browse Source
This commit is contained in:
Paul Sheppard 2012-08-22 23:31:07 -07:00
commit 736c62a91e
25 changed files with 745 additions and 536 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

49
API-README
View File

@ -42,7 +42,7 @@ To give an IP address/subnet access to a group you use the group letter
in front of the IP address instead of W: e.g. P:192.168.0/32
An IP address/subnet can only be a member of one group
A sample API group would be:
--api-groups P:switchpool:enablepool:addpool:disablepool:removepool.poolpriority:*
--api-groups P:switchpool:enablepool:addpool:disablepool:removepool:poolpriority:*
This would create a group 'P' that can do all current pool commands and all
non-priviliged commands - the '*' means all non-priviledged commands
Without the '*' the group would only have access to the pool commands
@ -281,6 +281,12 @@ The list of requests - a (*) means it requires privileged access - and replies a
none There is no reply section just the STATUS section
stating what failover-only was set to
coin COIN Coin mining information:
Hash Method=sha256/scrypt,
Current Block Time=N.N, <- 0 means none
Current Block Hash=XXXX..., <- blank if none
LP=true/false| <- LP is in use on at least 1 pool
When you enable, disable or restart a GPU or PGA, you will also get Thread messages
in the cgminer status window
@ -333,7 +339,10 @@ miner.php - an example web page to access the API
Feature Changelog for external applications using the API:
API V1.17
API V1.17 (cgminer v2.7.1)
Added API commands:
'coin'
Modified API commands:
'summary' - add 'Work Utility'
@ -745,6 +754,31 @@ Below is each variable that can be changed and an explanation of each
---------
Default:
$dfmt = 'H:i:s j-M-Y \U\T\CP';
Define the date format used to print full length dates
If you get the string 'UTCP' on the end of your dates shown, that
means you are using an older version of PHP and you can instead use:
$dfmt = 'H:i:s j-M-Y \U\T\CO';
The PHP documentation on the date format is here:
http://us.php.net/manual/en/function.date.php
---------
Default:
$title = 'Mine';
Web page title
If you know PHP you can of course use code to define it e.g.
$title = 'My Rig at: '.date($dfmt);
Which would set the web page title to something like:
My Rig at: 10:34:00 22-Aug-2012 UTC+10:00
---------
Default:
$readonly = false;
@ -948,7 +982,7 @@ The section defines what data you want in the summary table and the Fields
define what data you want shown from that section
Standard sections are:
SUMMARY, POOL, PGA, GPU, NOTIFY, CONFIG, DEVDETAILS, DEVS, STATS
SUMMARY, POOL, PGA, GPU, NOTIFY, CONFIG, DEVDETAILS, DEVS, STATS, COIN
Fields are the names as shown on the headers on the normal pages
@ -957,6 +991,7 @@ Fields can be 'name=new name' to display 'name' with a different heading
There are also now joined sections:
SUMMARY+POOL, SUMMARY+DEVS, SUMMARY+CONFIG, DEVS+NOTIFY, DEVS+DEVDETAILS
SUMMARY+COIN
These sections are an SQL join of the two sections and the fields in them
are named section.field where section. is the section the field comes from
@ -967,11 +1002,11 @@ Also note:
- empty columns (e.g. an unknown field) are not shown
- missing field data shows as blank
- the field name '*' matches all fields except in joined sections
(useful for STATS)
(useful for STATS and COIN)
There are 2 hard coded sections:
DATE - displays a date table like 'Summary'
RIGS - displays a rig table like 'Summary'
DATE - displays a date table like at the start of 'Summary'
RIGS - displays a rig table like at the start of 'Summary'
Each custom summary requires a second array, that can be empty, listing fields
to be totaled for each section
@ -1028,7 +1063,7 @@ Each table will have the chosen details for all the rigs specified in $rigs
section.fieldname, not just fieldname
The join code automatically adds 2 fields to each GPU device: 'Name' and 'ID'
They don't exist in the API 'devs' output but you can correctly calculate
They don't exist in the API 'devs' output but I can correctly calculate
them from the GPU device data
These two fields are used to join DEVS to NOTIFY i.e. find the NOTIFY
record that has the same Name and ID as the DEVS record and join them

23
FPGA-README
View File

@ -1,6 +1,7 @@
This README contains extended details about FPGA mining with cgminer
Bitforce
--bfl-range Use nonce range on bitforce devices if supported
@ -13,6 +14,28 @@ a cost of 1% in overall hashrate so this feature is disabled by default. It
is only recommended you enable this if you are mining with a minirig on
p2pool.
C source is included for a bitforce firmware flash utility on Linux only:
bitforce-firmware-flash.c
Using this, you can change the bitstream firmware on bitforce singles.
It is untested with other devices. Use at your own risk!
To compile:
make bitforce-firmware-flash
To flash your BFL, specify the BFL port and the flash file e.g.:
sudo ./bitforce-firmware-flash /dev/ttyUSB0 alphaminer_832.bfl
It takes a bit under 3 minutes to flash a BFL and shows a progress % counter
Once it completes, you may also need to wait about 15 seconds,
then power the BFL off and on again
If you get an error at the end of the BFL flash process stating:
"Error reading response from ZBX"
it may have worked successfully anyway.
Test mining on it to be sure if it worked or not.
You need to give cgminer about 10 minutes mining with the BFL to be sure of
the MH/s value reported with the changed firmware - and the MH/s reported
will be less than the firmware speed since you lose work on every block change.
Icarus

3
Makefile.am
View File

@ -10,7 +10,8 @@ endif
EXTRA_DIST = example.conf m4/gnulib-cache.m4 linux-usb-cgminer \
ADL_SDK/readme.txt api-example.php miner.php \
API.class API.java api-example.c windows-build.txt \
bitstreams/* API-README FPGA-README SCRYPT-README
bitstreams/* API-README FPGA-README SCRYPT-README \
bitforce-firmware-flash.c
SUBDIRS = lib compat ccan

74
NEWS
View File

@ -1,3 +1,77 @@
Version 2.7.2 - August 22, 2012
- Pick worksize 256 with Cypress if none is specified.
- Give warning with sdk2.7 and phatk as well.
- Whitelist sdk2.7 for diablo kernel as well.
- Only keep the last 6 blocks in the uthash database to keep memory usage
constant. Storing more is unhelpful anyway.
- BFL Flash - always distribute source
- Increase kernel versions signifying changed APIs.
- BFL flash - include source in builds and more FPGA-README
- Check we haven't staged work while waiting for a curl entry before proceeding.
- Use atomic ops to never miss a nonce on opencl kernels, including nonce==0,
also allowing us to make the output buffer smaller.
- Remove compile errors/warnings and document compile/usage in FPGA-README
- bitforce-firmware-flash.c by Luke-jr
- Ignore the submit_fail flag when deciding whether to recruit more curls or not
since we have upper bounds on how many curls can be recruited, this test is
redundant and can lead to problems.
- API-README update cgminer version number
- API-README fix groups P: example mistake
- API-README add COIN and other edits
- gpu->hit should be reset on new work as well.
- Do not add time to dynamic opencl calculations over a getwork.
- miner.php allow 'coin' is custom pages
Version 2.7.1 - August 21, 2012
- Update windows build instructions courtesy of sharky.
- Increase max curls to number of mining threads + queue * 2, accounting for up
and downstream comms.
- Queue enough requests to get started.
- There is no point trying to clone_work in get_work() any more since we clone
on every get_work_thread where possible.
- There is no point subtracting 1 from maxq in get_work_thread.
- Only set lagging flag once there are no staged work items.
- select_pool does not switch back to the primary once lagging is disabled.
- miner.php allow page title to be defined in myminer.php
- Free work before retrying in get_work_thread.
- Increment total work counter under mutex lock.
- Increment the queued count after the curl is popped in case there's a delay
waiting on curls and we think we've queued work when in fact we're waiting
- API new command 'coin' with mining information
- Do the dynamic timing in opencl code over a single pass through scanhash to
make sure we're only getting opencl times contributing to the measured inte
- Increase curl reaping time to 5 minutes since comms between curl requests can
be 2 mins apart with lots of rolltime.
- No need for extra variable in hash_push.
- Remove short options -r and -R to allow them to be reused and remove readme
entries for deprecated options.
- Avoid attempting to recursively lock the console mutex by disabling warnings
in gpu_fanpercent when fanspeed monitoring fails on windows. Debugged by l
- Deprecate the opt_fail_pause parameter, leaving a null placeholder for
existing configurations.
- Don't pause after failed getwork, set lagging flag and reassess.
- Add message to share if it's a resubmit.
- We should not be pausing in trying to resubmit shares.
- Get rid of the extending fail pause on failed connects since we discard work
after a period.
- get_work always returns true so turn it into a void function.
- get_work never returns false so get rid of fail pause loop.
- Get rid of pause and retry from get_upstream_work so we only do it from one
place.
- Deprecate the opt_retries feature as no one wants cgminer to automatically
abort. Leave a null placeholder for configurations that still have it.
- Reinstate fix ADL gpu-map not working when there are more ADL devices than
openCL patch by Nite69. Add virtual adl mapping for when none is specified o
- miner.php show summary Diff1 Shares total
- miner.php fix Work Utility totals
- miner.php format new Work Utility and Diff1 Shares
- API V1.17 show Work Utility and Diff1 Shares
Version 2.7.0 - August 18, 2012
- Introduce a new statistic, Work Utility, which is the number of difficulty 1

2
README
View File

@ -160,8 +160,6 @@ Options for both config file and command line:
--quiet|-q Disable logging output, display status and errors
--real-quiet Disable all output
--remove-disabled Remove disabled devices entirely, as if they didn't exist
--retries|-r <arg> Number of times to retry before giving up, if JSON-RPC call fails (-1 means never) (default: -1)
--retry-pause|-R <arg> Number of seconds to pause, between retries (default: 5)
--rotate <arg> Change multipool strategy from failover to regularly rotate at N minutes (default: 0)
--round-robin Change multipool strategy from failover to round robin on failure
--scan-time|-s <arg> Upper bound on time spent scanning current work, in seconds (default: 60)

32
adl.c
View File

@ -306,9 +306,10 @@ void init_adl(int nDevs)
if (gpus[i].mapped) {
vadapters[gpus[i].virtual_adl].virtual_gpu = i;
applog(LOG_INFO, "Mapping OpenCL device %d to ADL device %d", i, gpus[i].virtual_adl);
}
} else
gpus[i].virtual_adl = i;
}
if (!devs_match) {
applog(LOG_ERR, "WARNING: Number of OpenCL and ADL devices did not match!");
applog(LOG_ERR, "Hardware monitoring may NOT match up with devices!");
@ -347,11 +348,12 @@ void init_adl(int nDevs)
int iAdapterIndex;
int lpAdapterID;
ADLODPerformanceLevels *lpOdPerformanceLevels;
int lev;
int lev, adlGpu;
i = vadapters[gpu].id;
adlGpu = gpus[gpu].virtual_adl;
i = vadapters[adlGpu].id;
iAdapterIndex = lpInfo[i].iAdapterIndex;
gpus[gpu].virtual_gpu = vadapters[gpu].virtual_gpu;
gpus[gpu].virtual_gpu = vadapters[adlGpu].virtual_gpu;
/* Get unique identifier of the adapter, 0 means not AMD */
result = ADL_Adapter_ID_Get(iAdapterIndex, &lpAdapterID);
@ -361,11 +363,11 @@ void init_adl(int nDevs)
}
if (gpus[gpu].deven == DEV_DISABLED) {
gpus[i].gpu_engine =
gpus[i].gpu_memclock =
gpus[i].gpu_vddc =
gpus[i].gpu_fan =
gpus[i].gpu_powertune = 0;
gpus[gpu].gpu_engine =
gpus[gpu].gpu_memclock =
gpus[gpu].gpu_vddc =
gpus[gpu].gpu_fan =
gpus[gpu].gpu_powertune = 0;
continue;
}
@ -701,6 +703,8 @@ int gpu_fanpercent(int gpu)
ret = __gpu_fanpercent(ga);
unlock_adl();
if (unlikely(ga->has_fanspeed && ret == -1)) {
#if 0
/* Recursive calling applog causes a hang, so disable messages */
applog(LOG_WARNING, "GPU %d stopped reporting fanspeed due to driver corruption", gpu);
if (opt_restart) {
applog(LOG_WARNING, "Restart enabled, will attempt to restart cgminer");
@ -714,6 +718,14 @@ int gpu_fanpercent(int gpu)
ga->twin->twin = NULL;;
ga->twin = NULL;
}
#endif
if (opt_restart)
app_restart();
ga->has_fanspeed = false;
if (ga->twin) {
ga->twin->twin = NULL;;
ga->twin = NULL;
}
}
return ret;
}

60
api.c
View File

@ -187,6 +187,11 @@ static const char *NULLSTR = "(null)";
static const char *TRUESTR = "true";
static const char *FALSESTR = "false";
#ifdef USE_SCRYPT
static const char *SCRYPTSTR = "scrypt";
#endif
static const char *SHA256STR = "sha256";
static const char *DEVICECODE = ""
#ifdef HAVE_OPENCL
"GPU "
@ -232,7 +237,7 @@ static const char *OSINFO =
#define _SUMMARY "SUMMARY"
#define _STATUS "STATUS"
#define _VERSION "VERSION"
#define _MINECON "CONFIG"
#define _MINECONFIG "CONFIG"
#define _GPU "GPU"
#ifdef HAVE_AN_FPGA
@ -252,6 +257,7 @@ static const char *OSINFO =
#define _RESTART "RESTART"
#define _MINESTATS "STATS"
#define _CHECK "CHECK"
#define _MINECOIN "COIN"
static const char ISJSON = '{';
#define JSON0 "{"
@ -267,7 +273,7 @@ static const char ISJSON = '{';
#define JSON_SUMMARY JSON1 _SUMMARY JSON2
#define JSON_STATUS JSON1 _STATUS JSON2
#define JSON_VERSION JSON1 _VERSION JSON2
#define JSON_MINECON JSON1 _MINECON JSON2
#define JSON_MINECONFIG JSON1 _MINECONFIG JSON2
#define JSON_GPU JSON1 _GPU JSON2
#ifdef HAVE_AN_FPGA
@ -288,6 +294,7 @@ static const char ISJSON = '{';
#define JSON_CLOSE JSON3
#define JSON_MINESTATS JSON1 _MINESTATS JSON2
#define JSON_CHECK JSON1 _CHECK JSON2
#define JSON_MINECOIN JSON1 _MINECOIN JSON2
#define JSON_END JSON4 JSON5
static const char *JSON_COMMAND = "command";
@ -329,7 +336,7 @@ static const char *JSON_PARAMETER = "parameter";
#define MSG_NOGPUADL 30
#define MSG_INVINT 31
#define MSG_GPUINT 32
#define MSG_MINECON 33
#define MSG_MINECONFIG 33
#define MSG_GPUMERR 34
#define MSG_GPUMEM 35
#define MSG_GPUEERR 36
@ -382,6 +389,7 @@ static const char *JSON_PARAMETER = "parameter";
#define MSG_MISBOOL 75
#define MSG_INVBOOL 76
#define MSG_FOO 77
#define MSG_MINECOIN 78
enum code_severity {
SEVERITY_ERR,
@ -496,7 +504,7 @@ struct CODES {
{ SEVERITY_ERR, MSG_NOGPUADL,PARAM_GPU, "GPU %d does not have ADL" },
{ SEVERITY_ERR, MSG_INVINT, PARAM_STR, "Invalid intensity (%s) - must be '" _DYNAMIC "' or range " _MIN_INTENSITY_STR " - " _MAX_INTENSITY_STR },
{ SEVERITY_INFO, MSG_GPUINT, PARAM_BOTH, "GPU %d set new intensity to %s" },
{ SEVERITY_SUCC, MSG_MINECON, PARAM_NONE, "CGMiner config" },
{ SEVERITY_SUCC, MSG_MINECONFIG,PARAM_NONE, "CGMiner config" },
#ifdef HAVE_OPENCL
{ SEVERITY_ERR, MSG_GPUMERR, PARAM_BOTH, "Setting GPU %d memoryclock to (%s) reported failure" },
{ SEVERITY_SUCC, MSG_GPUMEM, PARAM_BOTH, "Setting GPU %d memoryclock to (%s) reported success" },
@ -534,6 +542,7 @@ struct CODES {
{ SEVERITY_ERR, MSG_MISBOOL, PARAM_NONE, "Missing parameter: true/false" },
{ SEVERITY_ERR, MSG_INVBOOL, PARAM_NONE, "Invalid parameter should be true or false" },
{ SEVERITY_SUCC, MSG_FOO, PARAM_BOOL, "Failover-Only set to %s" },
{ SEVERITY_SUCC, MSG_MINECOIN,PARAM_NONE, "CGMiner coin" },
{ SEVERITY_FAIL, 0, 0, NULL }
};
@ -1232,9 +1241,9 @@ static void minerconfig(__maybe_unused SOCKETTYPE c, __maybe_unused char *param,
#endif
sprintf(io_buffer, isjson
? "%s," JSON_MINECON
: "%s" _MINECON ",",
message(MSG_MINECON, 0, NULL, isjson));
? "%s," JSON_MINECONFIG
: "%s" _MINECONFIG ",",
message(MSG_MINECONFIG, 0, NULL, isjson));
root = api_add_int(root, "GPU Count", &gpucount, false);
root = api_add_int(root, "PGA Count", &pgacount, false);
@ -2741,6 +2750,42 @@ static void failoveronly(__maybe_unused SOCKETTYPE c, char *param, bool isjson,
strcpy(io_buffer, message(MSG_FOO, tf, NULL, isjson));
}
static void minecoin(__maybe_unused SOCKETTYPE c, __maybe_unused char *param, bool isjson, __maybe_unused char group)
{
struct api_data *root = NULL;
char buf[TMPBUFSIZ];
sprintf(io_buffer, isjson
? "%s," JSON_MINECOIN
: "%s" _MINECOIN ",",
message(MSG_MINECOIN, 0, NULL, isjson));
#ifdef USE_SCRYPT
if (opt_scrypt)
root = api_add_const(root, "Hash Method", SCRYPTSTR, false);
else
#endif
root = api_add_const(root, "Hash Method", SHA256STR, false);
mutex_lock(&ch_lock);
if (current_fullhash && *current_fullhash) {
root = api_add_timeval(root, "Current Block Time", &block_timeval, true);
root = api_add_string(root, "Current Block Hash", current_fullhash, true);
} else {
struct timeval t = {0,0};
root = api_add_timeval(root, "Current Block Time", &t, true);
root = api_add_const(root, "Current Block Hash", BLANK, false);
}
mutex_unlock(&ch_lock);
root = api_add_bool(root, "LP", &have_longpoll, false);
root = print_data(root, buf, isjson);
if (isjson)
strcat(buf, JSON_CLOSE);
strcat(io_buffer, buf);
}
static void checkcommand(__maybe_unused SOCKETTYPE c, char *param, bool isjson, char group);
struct CMDS {
@ -2792,6 +2837,7 @@ struct CMDS {
{ "stats", minerstats, false },
{ "check", checkcommand, false },
{ "failover-only", failoveronly, true },
{ "coin", minecoin, false },
{ NULL, NULL, false }
};

108
bitforce-firmware-flash.c Normal file
View File

@ -0,0 +1,108 @@
/*
* Copyright 2012 Luke Dashjr
*
* This program is free software; you can redistribute it and/or modify it
* under the terms of the GNU General Public License as published by the Free
* Software Foundation; either version 3 of the License, or (at your option)
* any later version. See COPYING for more details.
*/
#define _BSD_SOURCE
#include <stdint.h>
#include <stdio.h>
#include <string.h>
#include <libgen.h>
#include <arpa/inet.h>
#define BFL_FILE_MAGIC "BFLDATA"
#define BFL_UPLOAD_MAGIC "NGH-STREAM"
#define myassert(expr, n, ...) \
do { \
if (!(expr)) { \
fprintf(stderr, __VA_ARGS__); \
return n; \
} \
} while(0)
#define ERRRESP(buf) buf, (buf[strlen(buf)-1] == '\n' ? "" : "\n")
#define WAITFOROK(n, msg) \
do { \
myassert(fgets(buf, sizeof(buf), BFL), n, "Error reading response from " msg "\n"); \
myassert(!strcmp(buf, "OK\n"), n, "Invalid response from " msg ": %s%s", ERRRESP(buf)); \
} while(0)
int main(int argc, char**argv)
{
myassert(argc == 3, 1, "Usage: %s <serialdev> <firmware.bfl>\n", argv[0]);
setbuf(stdout, NULL);
// Check filename
char *FWname = basename(strdup(argv[2]));
size_t FWnameLen = strlen(FWname);
myassert(FWnameLen <= 255, 0x0f, "Firmware filename '%s' is too long\n", FWname);
uint8_t n8 = FWnameLen;
// Open and check firmware file
FILE *FW = fopen(argv[2], "r");
myassert(FW, 0x10, "Failed to open '%s' for reading\n", argv[2]);
char buf[0x20];
myassert(1 == fread(buf, 7, 1, FW), 0x10, "Failed to read from '%s'\n", argv[2]);
myassert(!memcmp(buf, BFL_FILE_MAGIC, sizeof(BFL_FILE_MAGIC)-1), 0x11, "'%s' doesn't look like a BFL firmware\n", argv[2]);
myassert(!fseek(FW, 0, SEEK_END), 0x12, "Failed to find end of '%s'\n", argv[2]);
long FWlen = ftell(FW);
myassert(FWlen > 0, 0x12, "Couldn't get size of '%s'\n", argv[2]);
myassert(!fseek(FW, 7, SEEK_SET), 0x12, "Failed to rewind firmware file after getting size\n");
FWlen -= 7;
printf("Firmware file looks OK :)\n");
// Open device
FILE *BFL = fopen(argv[1], "r+");
myassert(BFL, 0x20, "Failed to open '%s' for read/write\n", argv[1]);
myassert(!setvbuf(BFL, NULL, _IOFBF, 1032), 0x21, "Failed to setup buffer for device");
// ZAX: Start firmware upload
printf("Starting firmware upload... ");
myassert(1 == fwrite("ZAX", 3, 1, BFL), 0x22, "Failed to issue ZAX command\n");
WAITFOROK(0x22, "ZAX");
// Firmware upload header
myassert(1 == fwrite(BFL_UPLOAD_MAGIC, sizeof(BFL_UPLOAD_MAGIC)-1, 1, BFL), 0x23, "Failed to send firmware upload header (magic)\n");
uint32_t n32 = htonl(FWlen - FWlen / 6);
myassert(1 == fwrite(&n32, sizeof(n32), 1, BFL), 0x23, "Failed to send firmware upload header (size)\n");
myassert(1 == fwrite("\0\0", 2 , 1, BFL), 0x23, "Failed to send firmware upload header (padding 1)\n");
myassert(1 == fwrite(&n8, sizeof(n8) , 1, BFL), 0x23, "Failed to send firmware upload header (filename length)\n");
myassert(1 == fwrite(FWname, n8 , 1, BFL), 0x23, "Failed to send firmware upload header (filename)\n");
myassert(1 == fwrite("\0>>>>>>>>", 9 , 1, BFL), 0x23, "Failed to send firmware upload header (padding 2)\n");
WAITFOROK(0x23, "firmware upload header");
printf("OK, sending...\n");
// Actual firmware upload
long i, j;
for (i = 0, j = 0; i < FWlen; ++i) {
myassert(1 == fread(&n8, sizeof(n8), 1, FW), 0x30, "Error reading data from firmware file\n");
if (5 == i % 6)
continue;
n8 ^= 0x2f;
myassert(1 == fwrite(&n8, sizeof(n8), 1, BFL), 0x31, "Error sending data to device\n");
if (!(++j % 0x400)) {
myassert(1 == fwrite(">>>>>>>>", 8, 1, BFL), 0x32, "Error sending block-finish to device\n");
printf("\r%5.2f%% complete", (double)i * 100. / (double)FWlen);
WAITFOROK(0x32, "block-finish");
}
}
printf("\r100%% complete :)\n");
myassert(1 == fwrite(">>>>>>>>", 8, 1, BFL), 0x3f, "Error sending upload-finished to device\n");
myassert(fgets(buf, sizeof(buf), BFL), 0x3f, "Error reading response from upload-finished\n");
myassert(!strcmp(buf, "DONE\n"), 0x3f, "Invalid response from upload-finished: %s%s", ERRRESP(buf));
// ZBX: Finish programming
printf("Waiting for device... ");
myassert(1 == fwrite("ZBX", 3, 1, BFL), 0x40, "Failed to issue ZBX command\n");
WAITFOROK(0x40, "ZBX");
printf("All done! Try mining to test the flash succeeded.\n");
return 0;
}

264
cgminer.c
View File

@ -85,16 +85,13 @@ static char packagename[255];
bool opt_protocol;
static bool opt_benchmark;
static bool have_longpoll;
bool have_longpoll;
static bool want_per_device_stats;
bool use_syslog;
bool opt_quiet;
static bool opt_realquiet;
bool opt_loginput;
const int opt_cutofftemp = 95;
static int opt_retries = -1;
static int opt_fail_pause = 5;
static int fail_pause = 5;
int opt_log_interval = 5;
static int opt_queue = 1;
int opt_scantime = 60;
@ -170,7 +167,7 @@ static pthread_mutex_t hash_lock;
static pthread_mutex_t qd_lock;
static pthread_mutex_t *stgd_lock;
pthread_mutex_t console_lock;
static pthread_mutex_t ch_lock;
pthread_mutex_t ch_lock;
static pthread_rwlock_t blk_lock;
pthread_rwlock_t netacc_lock;
@ -213,8 +210,10 @@ bool curses_active;
static char current_block[37];
static char *current_hash;
char *current_fullhash;
static char datestamp[40];
static char blocktime[30];
struct timeval block_timeval;
struct block {
char hash[37];
@ -732,6 +731,11 @@ static char *set_icarus_timing(const char *arg)
}
#endif
static char *set_null(const char __maybe_unused *arg)
{
return NULL;
}
/* These options are available from config file or commandline */
static struct opt_table opt_config_table[] = {
#ifdef WANT_CPUMINE
@ -954,12 +958,12 @@ static struct opt_table opt_config_table[] = {
OPT_WITHOUT_ARG("--remove-disabled",
opt_set_bool, &opt_removedisabled,
"Remove disabled devices entirely, as if they didn't exist"),
OPT_WITH_ARG("--retries|-r",
opt_set_intval, opt_show_intval, &opt_retries,
"Number of times to retry before giving up, if JSON-RPC call fails (-1 means never)"),
OPT_WITH_ARG("--retry-pause|-R",
set_int_0_to_9999, opt_show_intval, &opt_fail_pause,
"Number of seconds to pause, between retries"),
OPT_WITH_ARG("--retries",
set_null, NULL, NULL,
opt_hidden),
OPT_WITH_ARG("--retry-pause",
set_null, NULL, NULL,
opt_hidden),
OPT_WITH_ARG("--rotate",
set_rotate, opt_show_intval, &opt_rotate_period,
"Change multipool strategy from failover to regularly rotate at N minutes"),
@ -1739,7 +1743,7 @@ static void reject_pool(struct pool *pool)
pool->enabled = POOL_REJECTING;
}
static bool submit_upstream_work(const struct work *work, CURL *curl)
static bool submit_upstream_work(const struct work *work, CURL *curl, bool resubmit)
{
char *hexstr = NULL;
json_t *val, *res;
@ -1814,11 +1818,11 @@ static bool submit_upstream_work(const struct work *work, CURL *curl)
applog(LOG_DEBUG, "PROOF OF WORK RESULT: true (yay!!!)");
if (!QUIET) {
if (total_pools > 1)
applog(LOG_NOTICE, "Accepted %s %s %d pool %d",
hashshow, cgpu->api->name, cgpu->device_id, work->pool->pool_no);
applog(LOG_NOTICE, "Accepted %s %s %d pool %d %s",
hashshow, cgpu->api->name, cgpu->device_id, work->pool->pool_no, resubmit ? "(resubmit)" : "");
else
applog(LOG_NOTICE, "Accepted %s %s %d",
hashshow, cgpu->api->name, cgpu->device_id);
applog(LOG_NOTICE, "Accepted %s %s %d %s",
hashshow, cgpu->api->name, cgpu->device_id, resubmit ? "(resubmit)" : "");
}
sharelog("accept", work);
if (opt_shares && total_accepted >= opt_shares) {
@ -1867,8 +1871,8 @@ static bool submit_upstream_work(const struct work *work, CURL *curl)
} else
strcpy(reason, "");
applog(LOG_NOTICE, "Rejected %s %s %d %s%s",
hashshow, cgpu->api->name, cgpu->device_id, where, reason);
applog(LOG_NOTICE, "Rejected %s %s %d %s%s %s",
hashshow, cgpu->api->name, cgpu->device_id, where, reason, resubmit ? "(resubmit)" : "");
sharelog(disposition, work);
}
@ -1951,9 +1955,11 @@ static inline struct pool *select_pool(bool lagging)
if (pool_strategy == POOL_BALANCE)
return select_balanced(cp);
if (pool_strategy != POOL_LOADBALANCE && (!lagging || opt_fail_only))
pool = cp;
else
if (pool_strategy != POOL_LOADBALANCE && (!lagging || opt_fail_only)) {
if (cp->prio != 0)
switch_pools(NULL);
pool = current_pool();
} else
pool = NULL;
while (!pool) {
@ -1989,7 +1995,6 @@ static bool get_upstream_work(struct work *work, CURL *curl)
struct timeval tv_start, tv_end, tv_elapsed;
json_t *val = NULL;
bool rc = false;
int retries = 0;
char *url;
applog(LOG_DEBUG, "DBG: sending %s get RPC call: %s", pool->rpc_url, rpc_req);
@ -1997,23 +2002,17 @@ static bool get_upstream_work(struct work *work, CURL *curl)
url = pool->rpc_url;
gettimeofday(&tv_start, NULL);
retry:
/* A single failure response here might be reported as a dead pool and
* there may be temporary denied messages etc. falsely reporting
* failure so retry a few times before giving up */
while (!val && retries++ < 3) {
pool_stats->getwork_attempts++;
val = json_rpc_call(curl, url, pool->rpc_userpass, rpc_req,
false, false, &work->rolltime, pool, false);
}
if (unlikely(!val)) {
applog(LOG_DEBUG, "Failed json_rpc_call in get_upstream_work");
goto out;
}
rc = work_decode(json_object_get(val, "result"), work);
if (!rc && retries < 3)
goto retry;
val = json_rpc_call(curl, url, pool->rpc_userpass, rpc_req, false,
false, &work->rolltime, pool, false);
pool_stats->getwork_attempts++;
if (likely(val)) {
rc = work_decode(json_object_get(val, "result"), work);
if (unlikely(!rc))
applog(LOG_DEBUG, "Failed to decode work in get_upstream_work");
} else
applog(LOG_DEBUG, "Failed json_rpc_call in get_upstream_work");
gettimeofday(&tv_end, NULL);
timersub(&tv_end, &tv_start, &tv_elapsed);
@ -2036,8 +2035,8 @@ retry:
total_getworks++;
pool->getwork_requested++;
json_decref(val);
out:
if (likely(val))
json_decref(val);
return rc;
}
@ -2048,7 +2047,9 @@ static struct work *make_work(void)
if (unlikely(!work))
quit(1, "Failed to calloc work in make_work");
mutex_lock(&control_lock);
work->id = total_work++;
mutex_unlock(&control_lock);
return work;
}
@ -2218,7 +2219,7 @@ static void recruit_curl(struct pool *pool)
* network delays/outages. */
static struct curl_ent *pop_curl_entry(struct pool *pool)
{
int curl_limit = opt_delaynet ? 5 : mining_threads * 4 / 3;
int curl_limit = opt_delaynet ? 5 : (mining_threads + opt_queue) * 2;
struct curl_ent *ce;
mutex_lock(&pool->pool_lock);
@ -2226,7 +2227,7 @@ retry:
if (!pool->curls)
recruit_curl(pool);
else if (list_empty(&pool->curlring)) {
if (pool->submit_fail || pool->curls >= curl_limit) {
if (pool->curls >= curl_limit) {
pthread_cond_wait(&pool->cr_cond, &pool->pool_lock);
goto retry;
} else
@ -2384,21 +2385,22 @@ static void *get_work_thread(void *userdata)
struct workio_cmd *wc = (struct workio_cmd *)userdata;
int ts, tq, maxq = opt_queue + mining_threads;
struct pool *pool = current_pool();
struct work *ret_work= NULL;
struct curl_ent *ce = NULL;
struct work *ret_work;
int failures = 0;
bool lagging = false;
pthread_detach(pthread_self());
applog(LOG_DEBUG, "Creating extra get work thread");
retry:
tq = global_queued();
ts = total_staged();
if (ts >= maxq)
goto out;
if (ts >= opt_queue && tq >= maxq - 1)
if (ts >= opt_queue && tq >= maxq)
goto out;
if (clone_available())
@ -2410,34 +2412,37 @@ static void *get_work_thread(void *userdata)
else
ret_work->thr = NULL;
if (opt_benchmark)
if (opt_benchmark) {
get_benchmark_work(ret_work);
else {
bool lagging = false;
ret_work->queued = true;
} else {
if (ts <= opt_queue)
if (!ts)
lagging = true;
pool = ret_work->pool = select_pool(lagging);
inc_queued();
ce = pop_curl_entry(pool);
if (!ce)
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 */
while (!get_upstream_work(ret_work, ce->curl)) {
if (unlikely((opt_retries >= 0) && (++failures > opt_retries))) {
applog(LOG_ERR, "json_rpc_call failed, terminating workio thread");
free_work(ret_work);
kill_work();
goto out;
}
if (!get_upstream_work(ret_work, ce->curl)) {
/* pause, then restart work-request loop */
applog(LOG_DEBUG, "json_rpc_call failed on get work, retry after %d seconds",
fail_pause);
sleep(fail_pause);
fail_pause += opt_fail_pause;
applog(LOG_DEBUG, "json_rpc_call failed on get work, retrying");
lagging = true;
dec_queued();
free_work(ret_work);
goto retry;
}
fail_pause = opt_fail_pause;
ret_work->queued = true;
}
@ -2534,8 +2539,8 @@ static void *submit_work_thread(void *userdata)
struct workio_cmd *wc = (struct workio_cmd *)userdata;
struct work *work = wc->work;
struct pool *pool = work->pool;
bool resubmit = false;
struct curl_ent *ce;
int failures = 0;
pthread_detach(pthread_self());
@ -2560,26 +2565,18 @@ static void *submit_work_thread(void *userdata)
ce = pop_curl_entry(pool);
/* submit solution to bitcoin via JSON-RPC */
while (!submit_upstream_work(work, ce->curl)) {
while (!submit_upstream_work(work, ce->curl, resubmit)) {
resubmit = true;
if (stale_work(work, true)) {
applog(LOG_NOTICE, "Share became stale while retrying submit, discarding");
total_stale++;
pool->stale_shares++;
break;
}
if (unlikely((opt_retries >= 0) && (++failures > opt_retries))) {
applog(LOG_ERR, "Failed %d retries ...terminating workio thread", opt_retries);
kill_work();
break;
}
/* pause, then restart work-request loop */
applog(LOG_INFO, "json_rpc_call failed on submit_work, retry after %d seconds",
fail_pause);
sleep(fail_pause);
fail_pause += opt_fail_pause;
applog(LOG_INFO, "json_rpc_call failed on submit_work, retrying");
}
fail_pause = opt_fail_pause;
push_curl_entry(ce, pool);
out:
workio_cmd_free(wc);
@ -2785,22 +2782,27 @@ static void restart_threads(void)
static void set_curblock(char *hexstr, unsigned char *hash)
{
unsigned char hash_swap[32];
struct timeval tv_now;
unsigned char block_hash_swap[32];
char *old_hash;
strcpy(current_block, hexstr);
gettimeofday(&tv_now, NULL);
get_timestamp(blocktime, &tv_now);
swap256(hash_swap, hash);
swap256(block_hash_swap, hash+4);
/* Don't free current_hash directly to avoid dereferencing when read
* elsewhere */
* elsewhere - and update block_timeval inside the same lock */
mutex_lock(&ch_lock);
gettimeofday(&block_timeval, NULL);
old_hash = current_hash;
current_hash = bin2hex(hash_swap, 16);
free(old_hash);
old_hash = current_fullhash;
current_fullhash = bin2hex(block_hash_swap, 32);
free(old_hash);
mutex_unlock(&ch_lock);
get_timestamp(blocktime, &block_timeval);
if (unlikely(!current_hash))
quit (1, "set_curblock OOM");
applog(LOG_INFO, "New block: %s...", current_hash);
@ -2856,6 +2858,20 @@ static void test_work_current(struct work *work)
quit (1, "test_work_current OOM");
strcpy(s->hash, hexstr);
wr_lock(&blk_lock);
/* Only keep the last 6 blocks in memory since work from blocks
* before this is virtually impossible and we want to prevent
* memory usage from continually rising */
if (HASH_COUNT(blocks) > 5) {
struct block *blocka, *blockb;
int count = 0;
HASH_ITER(hh, blocks, blocka, blockb) {
if (count++ < 6)
continue;
HASH_DEL(blocks, blocka);
free(blocka);
}
}
HASH_ADD_STR(blocks, hash, s);
wr_unlock(&blk_lock);
set_curblock(hexstr, work->data);
@ -2898,12 +2914,7 @@ static bool work_rollable(struct work *work)
static bool hash_push(struct work *work)
{
bool rc = true, dec = false;
if (work->queued) {
work->queued = false;
dec = true;
}
bool rc = true;
mutex_lock(stgd_lock);
if (work_rollable(work))
@ -2916,8 +2927,10 @@ static bool hash_push(struct work *work)
pthread_cond_signal(&getq->cond);
mutex_unlock(stgd_lock);
if (dec)
if (work->queued) {
work->queued = false;
dec_queued();
}
return rc;
}
@ -3507,9 +3520,9 @@ static void set_options(void)
immedok(logwin, true);
clear_logwin();
retry:
wlogprint("[Q]ueue: %d\n[S]cantime: %d\n[E]xpiry: %d\n[R]etries: %d\n"
"[P]ause: %d\n[W]rite config file\n[C]gminer restart\n",
opt_queue, opt_scantime, opt_expiry, opt_retries, opt_fail_pause);
wlogprint("[Q]ueue: %d\n[S]cantime: %d\n[E]xpiry: %d\n"
"[W]rite config file\n[C]gminer restart\n",
opt_queue, opt_scantime, opt_expiry);
wlogprint("Select an option or any other key to return\n");
input = getch();
@ -3537,22 +3550,6 @@ retry:
}
opt_expiry = selected;
goto retry;
} else if (!strncasecmp(&input, "r", 1)) {
selected = curses_int("Retries before failing (-1 infinite)");
if (selected < -1 || selected > 9999) {
wlogprint("Invalid selection\n");
goto retry;
}
opt_retries = selected;
goto retry;
} else if (!strncasecmp(&input, "p", 1)) {
selected = curses_int("Seconds to pause before network retries");
if (selected < 1 || selected > 9999) {
wlogprint("Invalid selection\n");
goto retry;
}
opt_fail_pause = selected;
goto retry;
} else if (!strncasecmp(&input, "w", 1)) {
FILE *fcfg;
char *str, filename[PATH_MAX], prompt[PATH_MAX + 50];
@ -4027,14 +4024,12 @@ static struct work *clone_work(struct work *work)
return work;
}
static bool get_work(struct work *work, struct thr_info *thr, const int thr_id)
static void get_work(struct work *work, struct thr_info *thr, const int thr_id)
{
struct timespec abstime = {0, 0};
struct work *work_heap;
struct timeval now;
struct pool *pool;
int failures = 0;
bool ret = false;
/* Tell the watchdog thread this thread is waiting on getwork and
* should not be restarted */
@ -4042,17 +4037,14 @@ static bool get_work(struct work *work, struct thr_info *thr, const int thr_id)
if (opt_benchmark) {
get_benchmark_work(work);
thread_reportin(thr);
return true;
goto out;
}
retry:
pool = current_pool();
if (reuse_work(work)) {
ret = true;
if (reuse_work(work))
goto out;
}
if (!pool->lagging && !total_staged() && global_queued() >= mining_threads + opt_queue) {
struct cgpu_info *cgpu = thr->cgpu;
@ -4102,29 +4094,13 @@ retry:
pool_resus(pool);
}
work_heap = clone_work(work_heap);
memcpy(work, work_heap, sizeof(struct work));
free_work(work_heap);
ret = true;
out:
if (unlikely(ret == false)) {
if ((opt_retries >= 0) && (++failures > opt_retries)) {
applog(LOG_ERR, "Failed %d times to get_work");
return ret;
}
applog(LOG_DEBUG, "Retrying after %d seconds", fail_pause);
sleep(fail_pause);
fail_pause += opt_fail_pause;
goto retry;
}
fail_pause = opt_fail_pause;
work->thr_id = thr_id;
thread_reportin(thr);
if (ret)
work->mined = true;
return ret;
work->mined = true;
}
bool submit_work_sync(struct thr_info *thr, const struct work *work_in)
@ -4281,11 +4257,9 @@ void *miner_thread(void *userdata)
mythr->work_restart = false;
if (api->free_work && likely(work->pool))
api->free_work(mythr, work);
if (unlikely(!get_work(work, mythr, thr_id))) {
applog(LOG_ERR, "work retrieval failed, exiting "
"mining thread %d", thr_id);
break;
}
get_work(work, mythr, thr_id);
cgpu->new_work = true;
gettimeofday(&tv_workstart, NULL);
work->blk.nonce = 0;
cgpu->max_hashes = 0;
@ -4570,16 +4544,9 @@ retry_pool:
gettimeofday(&end, NULL);
if (end.tv_sec - start.tv_sec > 30)
continue;
if (opt_retries == -1 || failures++ < opt_retries) {
if (failures == 1)
applog(LOG_WARNING,
"longpoll failed for %s, retrying every 30s", pool->lp_url);
sleep(30);
} else {
applog(LOG_ERR,
"longpoll failed for %s, ending thread", pool->lp_url);
goto out;
}
if (failures == 1)
applog(LOG_WARNING, "longpoll failed for %s, retrying every 30s", pool->lp_url);
sleep(30);
}
if (pool != cp) {
pool = select_longpoll_pool(cp);
@ -4618,7 +4585,7 @@ static void reap_curl(struct pool *pool)
list_for_each_entry_safe(ent, iter, &pool->curlring, node) {
if (pool->curls < 2)
break;
if (now.tv_sec - ent->tv.tv_sec > 60) {
if (now.tv_sec - ent->tv.tv_sec > 300) {
reaped++;
pool->curls--;
list_del(&ent->node);
@ -5754,6 +5721,9 @@ begin_bench:
pthread_detach(thr->pth);
#endif
for (i = 0; i < mining_threads + opt_queue; i++)
queue_request(NULL, false);
/* main loop - simply wait for workio thread to exit. This is not the
* normal exit path and only occurs should the workio_thread die
* unexpectedly */

12
configure.ac
View File

@ -2,7 +2,7 @@
##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
m4_define([v_maj], [2])
m4_define([v_min], [7])
m4_define([v_mic], [0])
m4_define([v_mic], [2])
##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
m4_define([v_ver], [v_maj.v_min.v_mic])
m4_define([lt_rev], m4_eval(v_maj + v_min))
@ -389,11 +389,11 @@ fi
AC_DEFINE_UNQUOTED([CGMINER_PREFIX], ["$prefix/bin"], [Path to cgminer install])
AC_DEFINE_UNQUOTED([PHATK_KERNNAME], ["phatk120724"], [Filename for phatk kernel])
AC_DEFINE_UNQUOTED([POCLBM_KERNNAME], ["poclbm120724"], [Filename for poclbm kernel])
AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn120724"], [Filename for diakgcn kernel])
AC_DEFINE_UNQUOTED([DIABLO_KERNNAME], ["diablo120724"], [Filename for diablo kernel])
AC_DEFINE_UNQUOTED([SCRYPT_KERNNAME], ["scrypt120724"], [Filename for scrypt kernel])
AC_DEFINE_UNQUOTED([PHATK_KERNNAME], ["phatk120823"], [Filename for phatk kernel])
AC_DEFINE_UNQUOTED([POCLBM_KERNNAME], ["poclbm120823"], [Filename for poclbm kernel])
AC_DEFINE_UNQUOTED([DIAKGCN_KERNNAME], ["diakgcn120823"], [Filename for diakgcn kernel])
AC_DEFINE_UNQUOTED([DIABLO_KERNNAME], ["diablo120823"], [Filename for diablo kernel])
AC_DEFINE_UNQUOTED([SCRYPT_KERNNAME], ["scrypt120823"], [Filename for scrypt kernel])
AC_SUBST(OPENCL_LIBS)

51
diablo120724.cl → diablo120823.cl
View File

@ -62,7 +62,7 @@ void search(
const uint c1_plus_k5, const uint b1_plus_k6,
const uint state0, const uint state1, const uint state2, const uint state3,
const uint state4, const uint state5, const uint state6, const uint state7,
__global uint * output)
volatile __global uint * output)
{
z ZA[930];
@ -1242,33 +1242,50 @@ void search(
ZA[924] = (ZCh(ZA[922], ZA[920], ZA[918]) + ZA[923]) + ZR26(ZA[922]);
#define FOUND (0x800)
#define NFLAG (0x7FF)
#define FOUND (0x0F)
#if defined(VECTORS4)
bool result = any(ZA[924] == 0x136032EDU);
if (result) {
if (ZA[924].x == 0x136032EDU)
output[FOUND] = output[NFLAG & Znonce.x] = Znonce.x;
if (ZA[924].y == 0x136032EDU)
output[FOUND] = output[NFLAG & Znonce.y] = Znonce.y;
if (ZA[924].z == 0x136032EDU)
output[FOUND] = output[NFLAG & Znonce.z] = Znonce.z;
if (ZA[924].w == 0x136032EDU)
output[FOUND] = output[NFLAG & Znonce.w] = Znonce.w;
uint found;
if (ZA[924].x == 0x136032EDU) {
found = atomic_add(&output[FOUND], 1);
output[found] = Znonce.x;
}
if (ZA[924].y == 0x136032EDU) {
found = atomic_add(&output[FOUND], 1);
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)
bool result = any(ZA[924] == 0x136032EDU);
if (result) {
if (ZA[924].x == 0x136032EDU)
output[FOUND] = output[NFLAG & Znonce.x] = Znonce.x;
if (ZA[924].y == 0x136032EDU)
output[FOUND] = output[NFLAG & Znonce.y] = Znonce.y;
uint found;
if (ZA[924].x == 0x136032EDU) {
found = atomic_add(&output[FOUND], 1);
output[found] = Znonce.x;
}
if (ZA[924].y == 0x136032EDU) {
found = atomic_add(&output[FOUND], 1);
output[found] = Znonce.y;
}
}
#else
if (ZA[924] == 0x136032EDU)
output[FOUND] = output[NFLAG & Znonce] = Znonce;
if (ZA[924] == 0x136032EDU) {
uint found = atomic_add(&output[FOUND], 1);
output[found] = Znonce;
}
#endif
}

47
diakgcn120724.cl → diakgcn120823.cl
View File

@ -48,7 +48,7 @@ __kernel
const uint state0A, const uint state0B,
const uint state1A, const uint state2A, const uint state3A, const uint state4A,
const uint state5A, const uint state6A, const uint state7A,
__global uint * output)
volatile __global uint * output)
{
u V[8];
u W[16];
@ -571,17 +571,46 @@ __kernel
V[7] += V[3] + W[12] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
#define FOUND (0x800)
#define NFLAG (0x7FF)
#define FOUND (0x0F)
#ifdef VECTORS4
if ((V[7].x == 0x136032edU) ^ (V[7].y == 0x136032edU) ^ (V[7].z == 0x136032edU) ^ (V[7].w == 0x136032edU))
output[FOUND] = output[NFLAG & nonce.x] = (V[7].x == 0x136032edU) ? nonce.x : ((V[7].y == 0x136032edU) ? nonce.y : ((V[7].z == 0x136032edU) ? nonce.z : nonce.w));
if ((V[7].x == 0x136032edU) ^ (V[7].y == 0x136032edU) ^ (V[7].z == 0x136032edU) ^ (V[7].w == 0x136032edU)) {
uint found;
if (V[7].x == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.x;
}
if (V[7].y == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
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
if ((V[7].x == 0x136032edU) + (V[7].y == 0x136032edU))
output[FOUND] = output[NFLAG & nonce.x] = (V[7].x == 0x136032edU) ? nonce.x : nonce.y;
if ((V[7].x == 0x136032edU) + (V[7].y == 0x136032edU)) {
uint found;
if (V[7].x == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.x;
}
if (V[7].y == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.y;
}
}
#else
if (V[7] == 0x136032edU)
output[FOUND] = output[NFLAG & nonce] = nonce;
if (V[7] == 0x136032edU) {
uint found = atomic_add(&output[FOUND], 1);
output[found] = nonce;
}
#endif
}

2
driver-bitforce.c
View File

@ -175,7 +175,7 @@ static int bitforce_autodetect_ftdi(void)
applog(LOG_DEBUG, "FTDI reports %u devices", (unsigned)numDevs);
buf = alloca(65 * numDevs);
bufptrs = alloca(numDevs + 1);
bufptrs = alloca(sizeof(*bufptrs) * (numDevs + 1));
for (i = 0; i < numDevs; ++i)
bufptrs[i] = &buf[i * 65];

103
driver-opencl.c
View File

@ -1460,6 +1460,10 @@ static void opencl_free_work(struct thr_info *thr, struct work *work)
const int thr_id = thr->id;
struct opencl_thread_data *thrdata = thr->cgpu_data;
_clState *clState = clStates[thr_id];
struct cgpu_info *gpu = thr->cgpu;
if (gpu->dynamic)
return;
clFinish(clState->commandQueue);
if (thrdata->res[FOUND]) {
@ -1491,7 +1495,6 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
const cl_kernel *kernel = &clState->kernel;
const int dynamic_us = opt_dynamic_interval * 1000;
struct timeval tv_gpuend;
cl_bool blocking;
cl_int status;
size_t globalThreads[1];
@ -1499,57 +1502,19 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
unsigned int threads;
int64_t hashes;
if (gpu->dynamic)
blocking = CL_TRUE;
else
blocking = CL_FALSE;
/* This finish flushes the readbuffer set with CL_FALSE later */
if (!blocking)
if (!gpu->dynamic)
clFinish(clState->commandQueue);
if (gpu->dynamic) {
double gpu_us;
/* Windows returns the same time for gettimeofday due to its
* 15ms timer resolution, so we must average the result over
* at least 5 values that are actually different to get an
* accurate result */
gpu->intervals++;
gettimeofday(&tv_gpuend, NULL);
gpu_us = us_tdiff(&tv_gpuend, &gpu->tv_gpumid);
if (gpu_us > 0 && ++gpu->hit > 4) {
gpu_us = us_tdiff(&tv_gpuend, &gpu->tv_gpustart) / gpu->intervals;
gpu->gpu_us_average = (gpu->gpu_us_average + gpu_us * 0.63) / 1.63;
/* Try to not let the GPU be out for longer than
* opt_dynamic_interval in ms, but increase
* intensity when the system is idle in dynamic mode */
if (gpu->gpu_us_average > dynamic_us) {
if (gpu->intensity > MIN_INTENSITY)
--gpu->intensity;
} else if (gpu->gpu_us_average < dynamic_us / 2) {
if (gpu->intensity < MAX_INTENSITY)
++gpu->intensity;
}
gpu->intervals = gpu->hit = 0;
}
}
set_threads_hashes(clState->vwidth, &threads, &hashes, globalThreads,
localThreads[0], gpu->intensity);
if (hashes > gpu->max_hashes)
gpu->max_hashes = hashes;
status = thrdata->queue_kernel_parameters(clState, &work->blk, globalThreads[0]);
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error: clSetKernelArg of all params failed.");
return -1;
}
/* MAXBUFFERS entry is used as a flag to say nonces exist */
/* FOUND entry is used as a counter to say how many nonces exist */
if (thrdata->res[FOUND]) {
/* Clear the buffer again */
status = clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, blocking, 0,
status = clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_FALSE, 0,
BUFFERSIZE, blank_res, 0, NULL, NULL);
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error: clEnqueueWriteBuffer failed.");
@ -1564,14 +1529,25 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
postcalc_hash_async(thr, work, thrdata->res);
}
memset(thrdata->res, 0, BUFFERSIZE);
if (!blocking)
clFinish(clState->commandQueue);
clFinish(clState->commandQueue);
}
gettimeofday(&gpu->tv_gpumid, NULL);
if (!gpu->intervals) {
gpu->tv_gpustart.tv_sec = gpu->tv_gpumid.tv_sec;
gpu->tv_gpustart.tv_usec = gpu->tv_gpumid.tv_usec;
if (gpu->dynamic) {
gettimeofday(&gpu->tv_gpumid, NULL);
if (gpu->new_work) {
gpu->new_work = false;
gpu->intervals = gpu->hit = 0;
}
if (!gpu->intervals) {
gpu->tv_gpustart.tv_sec = gpu->tv_gpumid.tv_sec;
gpu->tv_gpustart.tv_usec = gpu->tv_gpumid.tv_usec;
}
}
status = thrdata->queue_kernel_parameters(clState, &work->blk, globalThreads[0]);
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error: clSetKernelArg of all params failed.");
return -1;
}
if (clState->goffset) {
@ -1588,13 +1564,42 @@ static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
return -1;
}
status = clEnqueueReadBuffer(clState->commandQueue, clState->outputBuffer, blocking, 0,
status = clEnqueueReadBuffer(clState->commandQueue, clState->outputBuffer, CL_FALSE, 0,
BUFFERSIZE, thrdata->res, 0, NULL, NULL);
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error: clEnqueueReadBuffer failed error %d. (clEnqueueReadBuffer)", status);
return -1;
}
if (gpu->dynamic) {
double gpu_us;
clFinish(clState->commandQueue);
/* Windows returns the same time for gettimeofday due to its
* 15ms timer resolution, so we must average the result over
* at least 5 values that are actually different to get an
* accurate result */
gpu->intervals++;
gettimeofday(&tv_gpuend, NULL);
gpu_us = us_tdiff(&tv_gpuend, &gpu->tv_gpumid);
if (gpu_us > 0 && ++gpu->hit > 4) {
gpu_us = us_tdiff(&tv_gpuend, &gpu->tv_gpustart) / gpu->intervals;
gpu->gpu_us_average = (gpu->gpu_us_average + gpu_us * 0.63) / 1.63;
/* Try to not let the GPU be out for longer than
* opt_dynamic_interval in ms, but increase
* intensity when the system is idle in dynamic mode */
if (gpu->gpu_us_average > dynamic_us) {
if (gpu->intensity > MIN_INTENSITY)
--gpu->intensity;
} else if (gpu->gpu_us_average < dynamic_us / 2) {
if (gpu->intensity < MAX_INTENSITY)
++gpu->intensity;
}
gpu->intervals = gpu->hit = 0;
}
}
/* The amount of work scanned can fluctuate when intensity changes
* and since we do this one cycle behind, we increment the work more
* than enough to prevent repeating work */

1
example.conf
View File

@ -33,7 +33,6 @@
"gpu-threads" : "2",
"log" : "5",
"queue" : "1",
"retry-pause" : "5",
"scan-time" : "60",
"temp-hysteresis" : "3",

25
findnonce.c
View File

@ -172,6 +172,7 @@ struct pc_data {
struct work *work;
uint32_t res[MAXBUFFERS];
pthread_t pth;
int found;
};
static void send_sha_nonce(struct pc_data *pcd, cl_uint nonce)
@ -237,32 +238,22 @@ static void send_scrypt_nonce(struct pc_data *pcd, uint32_t nonce)
static void *postcalc_hash(void *userdata)
{
struct pc_data *pcd = (struct pc_data *)userdata;
struct thr_info *thr = pcd->thr;
int entry = 0, nonces = 0;
unsigned int entry = 0;
pthread_detach(pthread_self());
for (entry = 0; entry < FOUND; entry++) {
for (entry = 0; entry < pcd->res[FOUND]; entry++) {
uint32_t nonce = pcd->res[entry];
if (nonce) {
applog(LOG_DEBUG, "OCL NONCE %u", nonce);
if (opt_scrypt)
send_scrypt_nonce(pcd, nonce);
else
send_sha_nonce(pcd, nonce);
nonces++;
}
applog(LOG_DEBUG, "OCL NONCE %u found in slot %d", nonce, entry);
if (opt_scrypt)
send_scrypt_nonce(pcd, nonce);
else
send_sha_nonce(pcd, nonce);
}
free(pcd);
if (unlikely(!nonces)) {
applog(LOG_DEBUG, "No nonces found! Error in OpenCL code?");
hw_errors++;
thr->cgpu->hw_errors++;
}
return NULL;
}

5
findnonce.h
View File

@ -4,10 +4,9 @@
#include "config.h"
#define MAXTHREADS (0xFFFFFFFEULL)
#define MAXBUFFERS (0xFFF)
#define MAXBUFFERS (0x10)
#define BUFFERSIZE (sizeof(uint32_t) * MAXBUFFERS)
#define FOUND (0x800)
/* #define NFLAG (0x7FF) Just for reference */
#define FOUND (0x0F)
#ifdef HAVE_OPENCL
extern void precalc_hash(dev_blk_ctx *blk, uint32_t *state, uint32_t *data);

6
miner.h
View File

@ -374,6 +374,8 @@ struct cgpu_info {
int intervals, hit;
#endif
bool new_work;
float temp;
int cutofftemp;
@ -552,6 +554,7 @@ static inline void rwlock_init(pthread_rwlock_t *lock)
struct pool;
extern bool opt_protocol;
extern bool have_longpoll;
extern char *opt_kernel_path;
extern char *opt_socks_proxy;
extern char *cgminer_path;
@ -595,6 +598,7 @@ extern bool fulltest(const unsigned char *hash, const unsigned char *target);
extern int opt_scantime;
extern pthread_mutex_t console_lock;
extern pthread_mutex_t ch_lock;
extern pthread_mutex_t restart_lock;
extern pthread_cond_t restart_cond;
@ -669,6 +673,8 @@ extern unsigned int total_go, total_ro;
extern const int opt_cutofftemp;
extern int opt_log_interval;
extern unsigned long long global_hashrate;
extern char *current_fullhash;
extern struct timeval block_timeval;
#ifdef HAVE_OPENCL
typedef struct {

25
miner.php
View File

@ -1,7 +1,7 @@
<?php
session_start();
#
global $miner, $port, $readonly, $notify, $rigs;
global $title, $miner, $port, $readonly, $notify, $rigs;
global $rigtotals, $forcerigtotals;
global $socksndtimeoutsec, $sockrcvtimeoutsec;
global $checklastshare, $poolinputs, $hidefields;
@ -13,6 +13,9 @@ global $colouroverride, $placebuttons;
# See API-README for more details of these variables and how
# to configure miner.php
#
# Web page title
$title = 'Mine';
#
# Set $readonly to true to force miner.php to be readonly
# Set $readonly to false then it will check cgminer 'privileged'
$readonly = false;
@ -185,7 +188,7 @@ function getdom($domname)
#
function htmlhead($checkapi, $rig, $pg = null)
{
global $miner_font_family, $miner_font_size;
global $title, $miner_font_family, $miner_font_size;
global $error, $readonly, $poolinputs, $here;
global $ignorerefresh, $autorefresh;
@ -216,7 +219,7 @@ function htmlhead($checkapi, $rig, $pg = null)
$miner_font = "font-family:$miner_font_family; font-size:$miner_font_size;";
echo "<html><head>$refreshmeta
<title>Mine</title>
<title>$title</title>
<style type='text/css'>
td { $miner_font ".getcss('td')."}
td.two { $miner_font ".getcss('td.two')."}
@ -635,6 +638,10 @@ function fmt($section, $name, $value, $when, $alldata)
$class = $hiclass;
}
break;
case 'SUMMARY.Work Utility':
case 'total.Work Utility':
$ret = $value.'/m';
break;
case 'PGA.Temperature':
case 'GPU.Temperature':
case 'DEVS.Temperature':
@ -732,6 +739,8 @@ function fmt($section, $name, $value, $when, $alldata)
case 'SUMMARY.Discarded':
case 'POOL.Discarded':
case 'total.Discarded':
case 'POOL.Diff1 Shares':
case 'total.Diff1 Shares':
$parts = explode('.', $value, 2);
if (count($parts) == 1)
$dec = '';
@ -753,6 +762,7 @@ function fmt($section, $name, $value, $when, $alldata)
$class = $warnclass;
break;
case 'STATUS.When':
case 'COIN.Current Block Time':
$ret = date($dfmt, $value);
break;
case 'BUTTON.Rig':
@ -818,7 +828,8 @@ $singlerigsum = array(
'devs' => array('MHS av' => 1, 'MHS 5s' => 1, 'Accepted' => 1, 'Rejected' => 1,
'Hardware Errors' => 1, 'Utility' => 1, 'Total MH' => 1),
'pools' => array('Getworks' => 1, 'Accepted' => 1, 'Rejected' => 1, 'Discarded' => 1,
'Stale' => 1, 'Get Failures' => 1, 'Remote Failures' => 1),
'Stale' => 1, 'Get Failures' => 1, 'Remote Failures' => 1,
'Diff1 Shares' => 1),
'notify' => array('*' => 1));
#
function showtotal($total, $when, $oldvalues)
@ -1447,7 +1458,8 @@ $sectionmap = array(
'NOTIFY' => 'notify',
'DEVDETAILS' => 'devdetails',
'STATS' => 'stats',
'CONFIG' => 'config');
'CONFIG' => 'config',
'COIN' => 'coin');
#
function joinfields($section1, $section2, $join, $results)
{
@ -1582,6 +1594,7 @@ function joinsections($sections, $results, $errors)
case 'POOL':
case 'DEVS':
case 'CONFIG':
case 'COIN':
$sectionmap[$section] = $section;
$results[$section] = joinall($both[0], $both[1], $results);
break;
@ -2023,7 +2036,7 @@ function display()
newtable();
doforeach('version', 'rig summary', array(), array(), true);
$sum = array('MHS av', 'Getworks', 'Found Blocks', 'Accepted', 'Rejected', 'Discarded', 'Stale', 'Utility', 'Local Work', 'Total MH');
$sum = array('MHS av', 'Getworks', 'Found Blocks', 'Accepted', 'Rejected', 'Discarded', 'Stale', 'Utility', 'Local Work', 'Total MH', 'Work Utility', 'Diff1 Shares');
doforeach('summary', 'summary information', $sum, array(), false);
endtable();
otherrow('<td><br><br></td>');

162
mkinstalldirs
View File

@ -1,162 +0,0 @@
#! /bin/sh
# mkinstalldirs --- make directory hierarchy
scriptversion=2009-04-28.21; # UTC
# Original author: Noah Friedman <friedman@prep.ai.mit.edu>
# Created: 1993-05-16
# Public domain.
#
# This file is maintained in Automake, please report
# bugs to <bug-automake@gnu.org> or send patches to
# <automake-patches@gnu.org>.
nl='
'
IFS=" "" $nl"
errstatus=0
dirmode=
usage="\
Usage: mkinstalldirs [-h] [--help] [--version] [-m MODE] DIR ...
Create each directory DIR (with mode MODE, if specified), including all
leading file name components.
Report bugs to <bug-automake@gnu.org>."
# process command line arguments
while test $# -gt 0 ; do
case $1 in
-h | --help | --h*) # -h for help
echo "$usage"
exit $?
;;
-m) # -m PERM arg
shift
test $# -eq 0 && { echo "$usage" 1>&2; exit 1; }
dirmode=$1
shift
;;
--version)
echo "$0 $scriptversion"
exit $?
;;
--) # stop option processing
shift
break
;;
-*) # unknown option
echo "$usage" 1>&2
exit 1
;;
*) # first non-opt arg
break
;;
esac
done
for file
do
if test -d "$file"; then
shift
else
break
fi
done
case $# in
0) exit 0 ;;
esac
# Solaris 8's mkdir -p isn't thread-safe. If you mkdir -p a/b and
# mkdir -p a/c at the same time, both will detect that a is missing,
# one will create a, then the other will try to create a and die with
# a "File exists" error. This is a problem when calling mkinstalldirs
# from a parallel make. We use --version in the probe to restrict
# ourselves to GNU mkdir, which is thread-safe.
case $dirmode in
'')
if mkdir -p --version . >/dev/null 2>&1 && test ! -d ./--version; then
echo "mkdir -p -- $*"
exec mkdir -p -- "$@"
else
# On NextStep and OpenStep, the 'mkdir' command does not
# recognize any option. It will interpret all options as
# directories to create, and then abort because '.' already
# exists.
test -d ./-p && rmdir ./-p
test -d ./--version && rmdir ./--version
fi
;;
*)
if mkdir -m "$dirmode" -p --version . >/dev/null 2>&1 &&
test ! -d ./--version; then
echo "mkdir -m $dirmode -p -- $*"
exec mkdir -m "$dirmode" -p -- "$@"
else
# Clean up after NextStep and OpenStep mkdir.
for d in ./-m ./-p ./--version "./$dirmode";
do
test -d $d && rmdir $d
done
fi
;;
esac
for file
do
case $file in
/*) pathcomp=/ ;;
*) pathcomp= ;;
esac
oIFS=$IFS
IFS=/
set fnord $file
shift
IFS=$oIFS
for d
do
test "x$d" = x && continue
pathcomp=$pathcomp$d
case $pathcomp in
-*) pathcomp=./$pathcomp ;;
esac
if test ! -d "$pathcomp"; then
echo "mkdir $pathcomp"
mkdir "$pathcomp" || lasterr=$?
if test ! -d "$pathcomp"; then
errstatus=$lasterr
else
if test ! -z "$dirmode"; then
echo "chmod $dirmode $pathcomp"
lasterr=
chmod "$dirmode" "$pathcomp" || lasterr=$?
if test ! -z "$lasterr"; then
errstatus=$lasterr
fi
fi
fi
fi
pathcomp=$pathcomp/
done
done
exit $errstatus
# Local Variables:
# mode: shell-script
# sh-indentation: 2
# eval: (add-hook 'write-file-hooks 'time-stamp)
# time-stamp-start: "scriptversion="
# time-stamp-format: "%:y-%02m-%02d.%02H"
# time-stamp-time-zone: "UTC"
# time-stamp-end: "; # UTC"
# End:

15
ocl.c
View File

@ -392,7 +392,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
strstr(vbuff, "851.4") || // Windows 64 bit ""
strstr(vbuff, "831.4") ||
strstr(vbuff, "898.1") || // 12.2 driver SDK
strstr(vbuff, "923.1"))) { // 12.4 driver SDK
strstr(vbuff, "923.1") || // 12.4
strstr(vbuff, "938.1"))) { // SDK 2.7
applog(LOG_INFO, "Selecting diablo kernel");
clState->chosen_kernel = KL_DIABLO;
/* Detect all 7970s, older ATI and NVIDIA and use poclbm */
@ -410,9 +411,9 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
if (clState->chosen_kernel == KL_PHATK &&
(strstr(vbuff, "844.4") || strstr(vbuff, "851.4") ||
strstr(vbuff, "831.4") || strstr(vbuff, "898.1") ||
strstr(vbuff, "923.1"))) {
strstr(vbuff, "923.1") || strstr(vbuff, "938.1"))) {
applog(LOG_WARNING, "WARNING: You have selected the phatk kernel.");
applog(LOG_WARNING, "You are running SDK 2.6 which performs poorly with this kernel.");
applog(LOG_WARNING, "You are running SDK 2.6+ which performs poorly with this kernel.");
applog(LOG_WARNING, "Downgrade your SDK and delete any .bin files before starting again.");
applog(LOG_WARNING, "Or allow cgminer to automatically choose a more suitable kernel.");
}
@ -467,8 +468,12 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
clState->wsize = cgpu->work_size;
else if (strstr(name, "Tahiti"))
clState->wsize = 64;
else
clState->wsize = (clState->max_work_size <= 256 ? clState->max_work_size : 256) / clState->vwidth;
else {
if (strstr(name, "Cypress"))
clState->wsize = 256;
else
clState->wsize = (clState->max_work_size <= 256 ? clState->max_work_size : 256) / clState->vwidth;
}
cgpu->work_size = clState->wsize;
#ifdef USE_SCRYPT

51
phatk120724.cl → phatk120823.cl
View File

@ -164,7 +164,7 @@ void search( const uint state0, const uint state1, const uint state2, const uint
const uint PreW18, const uint PreW19,
const uint PreW31, const uint PreW32,
__global uint * output)
volatile __global uint * output)
{
@ -387,31 +387,48 @@ void search( const uint state0, const uint state1, const uint state2, const uint
W[117] += W[108] + Vals[3] + Vals[7] + P2(124) + P1(124) + Ch((Vals[0] + Vals[4]) + (K[59] + W(59+64)) + s1(64+59)+ ch(59+64),Vals[1],Vals[2]) -
(-(K[60] + H[7]) - S1((Vals[0] + Vals[4]) + (K[59] + W(59+64)) + s1(64+59)+ ch(59+64)));
#define FOUND (0x800)
#define NFLAG (0x7FF)
#define FOUND (0x0F)
#ifdef VECTORS4
bool result = W[117].x & W[117].y & W[117].z & W[117].w;
if (!result) {
if (!W[117].x)
output[FOUND] = output[NFLAG & W[3].x] = W[3].x;
if (!W[117].y)
output[FOUND] = output[NFLAG & W[3].y] = W[3].y;
if (!W[117].z)
output[FOUND] = output[NFLAG & W[3].z] = W[3].z;
if (!W[117].w)
output[FOUND] = output[NFLAG & W[3].w] = W[3].w;
uint found;
if (!W[117].x) {
found = atomic_add(&output[FOUND], 1);
output[found] = W[3].x;
}
if (!W[117].y) {
found = atomic_add(&output[FOUND], 1);
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
bool result = W[117].x & W[117].y;
if (!result) {
if (!W[117].x)
output[FOUND] = output[NFLAG & W[3].x] = W[3].x;
if (!W[117].y)
output[FOUND] = output[NFLAG & W[3].y] = W[3].y;
uint found;
if (!W[117].x) {
found = atomic_add(&output[FOUND], 1);
output[found] = W[3].x;
}
if (!W[117].y) {
found = atomic_add(&output[FOUND], 1);
output[found] = W[3].y;
}
}
#else
if (!W[117])
output[FOUND] = output[NFLAG & W[3]] = W[3];
if (!W[117]) {
uint found = atomic_add(&output[FOUND], 1);
output[found] = W[3];
}
#endif
}

65
poclbm120724.cl → poclbm120823.cl
View File

@ -80,7 +80,7 @@ void search(const uint state0, const uint state1, const uint state2, const uint
const uint D1A, const uint C1addK5, const uint B1addK6,
const uint W16addK16, const uint W17addK17,
const uint PreVal4addT1, const uint Preval0,
__global uint * output)
volatile __global uint * output)
{
u Vals[24];
u *W = &Vals[8];
@ -1311,43 +1311,46 @@ Vals[1]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
Vals[1]+=K[59];
Vals[1]+=Vals[5];
#define FOUND (0x800)
#define NFLAG (0x7FF)
Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
Vals[2]+=W[12];
Vals[2]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
Vals[2]+=W[5];
Vals[2]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
Vals[2]+=Vals[0];
Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
#define FOUND (0x0F)
#if defined(VECTORS2) || defined(VECTORS4)
Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
Vals[2]+=W[12];
Vals[2]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
Vals[2]+=W[5];
Vals[2]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
Vals[2]+=Vals[0];
Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
if (any(Vals[2] == 0x136032edU)) {
if (Vals[2].x == 0x136032edU)
output[FOUND] = output[NFLAG & nonce.x] = nonce.x;
if (Vals[2].y == 0x136032edU)
output[FOUND] = output[NFLAG & nonce.y] = nonce.y;
uint found;
if (Vals[2].x == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.x;
}
if (Vals[2].y == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.y;
}
#if defined(VECTORS4)
if (Vals[2].z == 0x136032edU)
output[FOUND] = output[NFLAG & nonce.z] = nonce.z;
if (Vals[2].w == 0x136032edU)
output[FOUND] = output[NFLAG & nonce.w] = nonce.w;
if (Vals[2].z == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.z;
}
if (Vals[2].w == 0x136032edU) {
found = atomic_add(&output[FOUND], 1);
output[found] = nonce.w;
}
#endif
}
#else
if ((Vals[2]+
Ma(Vals[6],Vals[5],Vals[7])+
(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22))+
W[12]+
(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U))+
W[5]+
(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U))+
Vals[0]+
(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25))+
ch(Vals[1],Vals[4],Vals[3])) == 0x136032edU)
output[FOUND] = output[NFLAG & nonce] = nonce;
if (Vals[2] == 0x136032edU) {
uint found = atomic_add(&output[FOUND], 1);
output[found] = nonce;
}
#endif
}

13
scrypt120724.cl → scrypt120823.cl
View File

@ -682,12 +682,11 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
unshittify(X);
}
#define FOUND (0x800)
#define NFLAG (0x7FF)
#define FOUND (0x0F)
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void search(__global const uint4 * restrict input,
__global uint*restrict output, __global uint4*restrict padcache,
volatile __global uint*restrict output, __global uint4*restrict padcache,
const uint4 midstate0, const uint4 midstate16, const uint target)
{
uint gid = get_global_id(0);
@ -721,9 +720,11 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
SHA256_fixed(&tmp0,&tmp1);
SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(0x80000000U, 0U, 0U, 0U), (uint4)(0U, 0U, 0U, 0x300U));
bool found = (EndianSwap(ostate1.w) <= target);
if (found)
output[FOUND] = output[NFLAG & gid] = gid;
bool result = (EndianSwap(ostate1.w) <= target);
if (result) {
uint found = atomic_add(&output[FOUND], 1);
output[found] = gid;
}
}
/*-

83
windows-build.txt
View File

@ -12,12 +12,13 @@ Most of what is described below (copying files, downloading files, etc.) can be
directly in the MinGW MSYS shell; these instructions do not do so because package
versions and links change over time. The best way is to use your browser, go to the
links directly, and see for yourself which versions you want to install.
Winrar was used to do the extracting of archive files in the making of this guide.
If you think that this documentation was helpful and you wish to donate, you can
If you think that this documentation was helpful and you wish to donate, you can
do so at the following address. 12KaKtrK52iQjPdtsJq7fJ7smC32tXWbWr
**************************************************************************************
* A tip that might help you along the way *
* A tip that might help you along the way *
**************************************************************************************
Enable "QuickEdit Mode" in your Command Prompt Window or MinGW Command Prompt
Window (No need to go into the context menu to choose edit-mark/copy/paste):
@ -36,7 +37,7 @@ what you copied.
Go to this url ==> http://www.mingw.org/wiki/Getting_Started
Click the link that says "Download and run the latest mingw-get-inst version."
Download and run the latest file. Install MinGW in the default directory.
(I downloaded the one labeled "mingw-get-inst-20111118" - note that this could
(I downloaded the one labeled "mingw-get-inst-20120426" - note that this could
be a different version later.)
Make sure to check the option for "Download latest repository catalogs".
I just selected all the check boxes (excluding "Fortran Compiler") so that everything
@ -67,7 +68,7 @@ struct tcp_keepalive
**************************************************************************************
* Run the MSYS shell for the first time to create your user directory *
**************************************************************************************
**************************************************************************************
(Start Icon/keyboard key ==> All Programs ==> MinGW ==> MinGW Shell).
This will create your user directory for you.
@ -83,44 +84,44 @@ window.
**************************************************************************************
* Copy CGMiner source to your MSYS working directory *
**************************************************************************************
Copy CGMiner source code directory into:
Copy CGMiner source code directory into:
\MinGW\msys\1.0\home\(folder with your user name)
**************************************************************************************
* Install AMD APP SDK, latest version (only if you want GPU mining) *
**************************************************************************************
Note: You do not need to install the AMD APP SDK if you are only using Nvidia GPU's
Go to this url for the latest AMD APP SDK:
Go to this url for the latest AMD APP SDK:
http://developer.amd.com/sdks/AMDAPPSDK/downloads/Pages/default.aspx
Go to this url for legacy AMD APP SDK's:
http://developer.amd.com/sdks/AMDAPPSDK/downloads/pages/AMDAPPSDKDownloadArchive.aspx
Download and install whichever version you like best.
Copy the folders in \Program Files (x86)\AMD APP\include to \MinGW\include
Copy the folders in \Program Files (x86)\AMD APP\include to \MinGW\include
Copy \Program Files (x86)\AMD APP\lib\x86\libOpenCL.a to \MinGW\lib
Note: If you are on a 32 bit version of windows "Program Files (x86)" will be
Note: If you are on a 32 bit version of windows "Program Files (x86)" will be
"Program Files".
Note2: If you update your APP SDK later you might want to recopy the above files
Note2: If you update your APP SDK later you might want to recopy the above files
**************************************************************************************
* Install AMD ADL SDK, latest version (only if you want GPU monitoring) *
**************************************************************************************
Note: You do not need to install the AMD ADL SDK if you are only using Nvidia GPU's
Note: You do not need to install the AMD ADL SDK if you are only using Nvidia GPU's
Go to this url ==> http://developer.amd.com/sdks/ADLSDK/Pages/default.aspx
Download and unzip the file you downloaded.
Pull adl_defines.h, adl_sdk.h, and adl_structures.h out of the include folder
Pull adl_defines.h, adl_sdk.h, and adl_structures.h out of the include folder
Put those files into the ADL_SDK folder in your source tree as shown below.
\MinGW\msys\1.0\home\(folder with your user name)\cgminer-x.x.x\ADL_SDK
**************************************************************************************
* Install GTK-WIN, required for Pkg-config in the next step *
**************************************************************************************
Go to this url ==> http://sourceforge.net/projects/gtk-win/
Go to this url ==> http://sourceforge.net/projects/gtk-win/
Download the file.
After you have downloaded the file Double click/run it and this will install GTK+
I chose all the selection boxes when I installed.
Copy libglib-2.0-0.dll and intl.dll from \Program Files (x86)\gtk2-runtime\bin to
Copy libglib-2.0-0.dll and intl.dll from \Program Files (x86)\gtk2-runtime\bin to
\MinGW\bin
Note: If you are on a 32 bit version of windows "Program Files (x86)" will be
Note: If you are on a 32 bit version of windows "Program Files (x86)" will be
"Program Files".
**************************************************************************************
@ -132,28 +133,28 @@ Download the file from the tool link. Extract "pkg-config.exe" from bin and plac
your \MinGW\bin directory.
Download the file from the "Dev" link. Extract "pkg.m4" from share\aclocal and place
in your \MingW\share\aclocal directory.
**************************************************************************************
* Install libcurl *
**************************************************************************************
Go to this url ==> http://curl.haxx.se/download.html#Win32
At the section where it says "Win32 - Generic", Click on the link that indicates
Win32 2000.XP 7.24.0 libcurl SSL and download it.
Win32 2000.XP 7.27.0 libcurl SSL and download it.
The one I downloaded may not be current for you. Choose the latest.
Extract the files that are in the zip (bin, include, and lib) to their respective
locations in MinGW (\MinGW\bin, \MinGW\include, and \MinGW\lib).
Edit the file \MinGW\lib\pkgconfig\libcurl.pc and change "-lcurl" to
Edit the file \MinGW\lib\pkgconfig\libcurl.pc and change "-lcurl" to
"-lcurl -lcurldll".
Ref. http://old.nabble.com/gcc-working-with-libcurl-td20506927.html
**************************************************************************************
* Build cgminer.exe *
**************************************************************************************
Run the MinGW MSYS shell
(Start Icon/keyboard key ==> All Programs ==> MinGW ==> MinGW Shell).
Run the MinGW MSYS shell
(Start Icon/keyboard key ==> All Programs ==> MinGW ==> MinGW Shell).
Change the working directory to your CGMiner project folder.
Example: cd cgminer-2.1.2 [Enter Key] if you are unsure then type "ls -la"
Another way is to type "cd cg" and then press the tab key; It will auto fill.
Another way is to type "cd cg" and then press the tab key; It will auto fill.
Type the lines below one at a time. Look for problems after each one before going on
to the next.
@ -169,16 +170,16 @@ Make a directory and copy the following files into it. This will be your CGMiner
Folder that you use for mining. Remember the .cl filenames could change on later
releases. If you installed a different version of libcurl then some of those dll's
may be different as well.
cgminer.exe from \MinGW\msys\1.0\home\(username)\cgminer-x.x.x
cgminer.exe from \MinGW\msys\1.0\home\(username)\cgminer-x.x.x
*.cl from \MinGW\msys\1.0\home\(username)\cgminer-x.x.x
README from \MinGW\msys\1.0\home\(username)\cgminer-x.x.x
libcurl.dll from \MinGW\bin
libeay32.dll from \MinGW\bin
libidn-11.dll from \MinGW\bin
libssl32.dll from \MinGW\bin
libeay32.dll from \MinGW\bin
ssleay32.dll from \MinGW\bin
libpdcurses.dll from \MinGW\bin
pthreadGC2.dll from \MinGW\bin
**************************************************************************************
* Optional - Install Git into MinGW/MSYS *
**************************************************************************************
@ -187,12 +188,13 @@ Click on the Downloads tab.
Download the latest "Portable" git archive.
Extract the git*.exe files from the bin folder and put them into \MinGW\bin.
Extract the share\git-core folder and place it into \MinGW\share.
After the previous step you should have a folder called \MinGW\share\git-core.
To test if it is working, open a MinGW shell and type the following:
git config -–global core.autocrlf false (note: one time run only)
git clone git://github.com/ckolivas/cgminer.git
If you simply just want to update the source after you have already cloned, type:
git pull git://github.com/ckolivas/cgminer.git
git pull
Now you can get the latest source directly from github.
@ -209,21 +211,38 @@ From now on when your current working directory is the cgminer source directory
You can simply type "adl.sh" and it will place the ADL header files into place
For you. Make sure you never remove the ADL_SDK folder from your home folder.
**************************************************************************************
* Optional - Install libusb if you need auto USB device detection; required for Ztex *
**************************************************************************************
Go to this url ==> http://libusbx.org/
Click on the "Downloads" tab.
Click on "releases".
Click on the latest version. I downloaded 1.0.12; yours may be newer.
Do not download from the link that says "Looking for the latest version?".
Click on "Windows"
Click on the file and download it. I downloaded libusbx-1.0.12-win.7z.
Extract the the following from the file and place in where directed.
Copy libusb.h from include\libusbx-1.0 to \MinGW\include\libusb-1.0\libusb.h
Copy contents of MinGW32\static \MinGW\lib
Copy contents of MinGW32\dll to \MinGW\lib
You will have to copy "libusb-1.0.dll" to your working cgminer binary directory.
**************************************************************************************
* Some ./configure options *
**************************************************************************************
--enable-cpumining Build with cpu mining support(default disabled)
--disable-opencl Override detection and disable building with opencl
--disable-adl Override detection and disable building with adl
--enable-bitforce Compile support for BitForce FPGAs(default disabled)
--enable-icarus Compile support for Icarus Board(default disabled)
--enable-modminer Compile support for ModMiner FPGAs(default disabled)
--enable-ztex Compile support for Ztex Board(default disabled)
--enable-scrypt Compile support for scrypt litecoin mining (default disabled)
--without-curses Compile support for curses TUI (default enabled)
--without-libudev Autodetect FPGAs using libudev (default enabled)
######################################################################################
# #
# Native WIN32 setup and build instructions (on mingw32/Windows) complete #
# #
######################################################################################
Addendum:
Ztex support requires libusb support. The most comprehensive support is
currently available from the libusbx project here:
http://libusbx.org/
######################################################################################