From 06f71ea0b7781544250e699a1bd920e45bdc32dd Mon Sep 17 00:00:00 2001 From: David Vorick Date: Thu, 18 Jun 2015 00:15:23 -0400 Subject: [PATCH 1/4] cleanup, suggest additional changes via TODO --- sia-gpu-miner.c | 50 +++++++++++++++++++++++++++++++++++++------------ 1 file changed, 38 insertions(+), 12 deletions(-) diff --git a/sia-gpu-miner.c b/sia-gpu-miner.c index ed91599..9717b3f 100644 --- a/sia-gpu-miner.c +++ b/sia-gpu-miner.c @@ -1,3 +1,4 @@ +// TODO: document this #ifdef __linux__ #define _GNU_SOURCE #define _POSIX_SOURCE @@ -13,12 +14,16 @@ #include "network.h" +// TODO: document this #ifdef __APPLE__ #include #else #include #endif +// TODO: explain what these are, not just why they have the values they were +// given. +// // Minimum intensity being less than 8 may break things #define MIN_INTENSITY 8 #define MAX_INTENSITY 32 @@ -27,8 +32,10 @@ // TODO: Might wanna establish a min/max for this, too... #define DEFAULT_CPI 3 +// TODO: Document this #define MAX_SOURCE_SIZE (0x200000) +// TODO: Document these cl_command_queue command_queue = NULL; cl_mem blockHeadermobj = NULL; cl_mem headerHashmobj = NULL; @@ -37,21 +44,32 @@ cl_mem nonceOutmobj = NULL; cl_kernel kernel = NULL; cl_int ret; +// TODO: Document these size_t local_item_size = 256; -unsigned int blocks_mined = 0, intensity = DEFAULT_INTENSITY; +unsigned int blocks_mined = 0; +unsigned int intensity = DEFAULT_INTENSITY; static volatile int quit = 0; int target_corrupt_flag = 0; +// TODO: explain what this is void quitSignal(int __unused) { quit = 1; - printf("\nCaught deadly signal, quitting...\n"); + printf("\nCaught kill signal, quitting...\n"); + // TODO: ??? It just prints a statement? Shouldn't it do + // something more, like call exit() or cleanup()? } +// TODO: This doesn't really count as a docstring. Need to explain what the +// function does at a high level, not a low level. Low level is for reading the +// code. +// // Perform (2^intensity) * cycles_per_iter hashes // Return -1 if a block is found // Else return the hashrate in MH/s double grindNonces(int cycles_per_iter) { // Start timing this iteration + // + // TODO: Explain the difference between the linux and non-linux stuff. #ifdef __linux__ struct timespec begin, end; clock_gettime(CLOCK_REALTIME, &begin); @@ -97,18 +115,25 @@ double grindNonces(int cycles_per_iter) { // Copy input data to the memory buffer ret = clEnqueueWriteBuffer(command_queue, blockHeadermobj, CL_TRUE, 0, 80 * sizeof(uint8_t), blockHeader, 0, NULL, NULL); - if (ret != CL_SUCCESS) { printf("failed to write to blockHeadermobj buffer: %d\n", ret); exit(1); } + if (ret != CL_SUCCESS) { + printf("failed to write to blockHeadermobj buffer: %d\n", ret); exit(1); + } ret = clEnqueueWriteBuffer(command_queue, nonceOutmobj, CL_TRUE, 0, 8 * sizeof(uint8_t), nonceOut, 0, NULL, NULL); - if (ret != CL_SUCCESS) { printf("failed to write to targmobj buffer: %d\n", ret); exit(1); } + if (ret != CL_SUCCESS) { + printf("failed to write to targmobj buffer: %d\n", ret); exit(1); + } + // Run the kernel ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, &globalid_offset, &global_item_size, &local_item_size, 0, NULL, NULL); - if (ret != CL_SUCCESS) { printf("failed to start kernel: %d\n", ret); exit(1); } + if (ret != CL_SUCCESS) { + printf("failed to start kernel: %d\n", ret); exit(1); + } - // Copy result to host + // Copy result to host and see if a block was found. ret = clEnqueueReadBuffer(command_queue, nonceOutmobj, CL_TRUE, 0, 8 * sizeof(uint8_t), nonceOut, 0, NULL, NULL); - if (ret != CL_SUCCESS) { printf("failed to read nonce from buffer: %d\n", ret); exit(1); } - - // Did we find one? + if (ret != CL_SUCCESS) { + printf("failed to read nonce from buffer: %d\n", ret); exit(1); + } if (nonceOut[0] != 0) { // Copy nonce to header. memcpy(blockHeader+32, nonceOut, 8); @@ -121,6 +146,9 @@ double grindNonces(int cycles_per_iter) { } // Hashrate is inaccurate if a block was found + // + // TODO: this comment isn't terribly helpful. Also you need to document the + // ifdef __linux__ stuff again. #ifdef __linux__ clock_gettime(CLOCK_REALTIME, &end); double nanosecondsElapsed = 1e9 * (double)(end.tv_sec - begin.tv_sec) + (double)(end.tv_nsec - begin.tv_nsec); @@ -305,7 +333,7 @@ int main(int argc, char *argv[]) { printf("\t\tOpenCL platforms will likely have different devices available. Default is 0.\n"); printf("\n"); printf("\t C - cycles per iter: Number of kernel executions between Sia API calls and hash rate updates\n"); - printf("\t\tIncrease this if your miner is receiving invalid targets. Default is %ud.\n", DEFAULT_CPI); + printf("\t\tIncrease this if your miner is receiving invalid targets. Default is %u.\n", DEFAULT_CPI); printf("\n"); printPlatformsAndDevices(); exit(0); @@ -511,8 +539,6 @@ int main(int argc, char *argv[]) { ret = clReleaseContext(context); free_network(); - free(source_str); - return 0; } From d808282bc7d754254aed544006f6a192ecbfbcc9 Mon Sep 17 00:00:00 2001 From: JoshVorick Date: Fri, 19 Jun 2015 02:03:36 -0400 Subject: [PATCH 2/4] General linting --- sia-gpu-miner.c | 93 +++++++++++++++++++++++++------------------------ 1 file changed, 48 insertions(+), 45 deletions(-) diff --git a/sia-gpu-miner.c b/sia-gpu-miner.c index 9717b3f..e94ad8e 100644 --- a/sia-gpu-miner.c +++ b/sia-gpu-miner.c @@ -1,4 +1,10 @@ -// TODO: document this +/* + * A cross-platform GPU miner built for Sia + * using OpenCL for interacting with the graphics card + * and libcurl for interacting with the Sia daemon. + */ + +// If using Linux, use a more accurate and reliable timer #ifdef __linux__ #define _GNU_SOURCE #define _POSIX_SOURCE @@ -14,62 +20,63 @@ #include "network.h" -// TODO: document this +// OpenCL header is named differently on Apple devices for some reason #ifdef __APPLE__ #include #else #include #endif -// TODO: explain what these are, not just why they have the values they were -// given. -// -// Minimum intensity being less than 8 may break things +// 2^intensity hashes are calculated each time the kernel is called +// Minimum of 2^8 (256) because our default local_item_size is 256 +// global_item_size (2^intensity) must be a multiple of local_item_size +// Max of 2^32 so that people can't send an hour of work to the GPU at one time #define MIN_INTENSITY 8 #define MAX_INTENSITY 32 #define DEFAULT_INTENSITY 16 -// TODO: Might wanna establish a min/max for this, too... -#define DEFAULT_CPI 3 +// Number of times the GPU kernel is called between updating the command line text +#define MIN_CPI 1 // Must do one call per update +#define MAX_CPI 65536 // 2^16 is a slightly arbitrary max +#define DEFAULT_CPI 3 -// TODO: Document this +// The maximum size of the .cl file we read in and compile #define MAX_SOURCE_SIZE (0x200000) -// TODO: Document these +// Objects needed to call the kernel +// global namespace so our grindNonce function can access them cl_command_queue command_queue = NULL; -cl_mem blockHeadermobj = NULL; -cl_mem headerHashmobj = NULL; -cl_mem targmobj = NULL; -cl_mem nonceOutmobj = NULL; cl_kernel kernel = NULL; cl_int ret; -// TODO: Document these -size_t local_item_size = 256; +// mem objects for storing our kernel parameters +cl_mem blockHeadermobj = NULL; +cl_mem nonceOutmobj = NULL; + +// More gobal variables the grindNonce needs to access +size_t local_item_size = 256; // Size of local work groups. 256 is usually optimal unsigned int blocks_mined = 0; unsigned int intensity = DEFAULT_INTENSITY; static volatile int quit = 0; + +// If we get a corrupt target, we want to remember so that if subsequent curl calls +// reutrn more corrupt targets, we don't spam the cmd line with errors int target_corrupt_flag = 0; -// TODO: explain what this is +// Set quit variable when SIGINT is received so we can do proper cleanup void quitSignal(int __unused) { quit = 1; printf("\nCaught kill signal, quitting...\n"); - // TODO: ??? It just prints a statement? Shouldn't it do - // something more, like call exit() or cleanup()? } -// TODO: This doesn't really count as a docstring. Need to explain what the -// function does at a high level, not a low level. Low level is for reading the -// code. -// -// Perform (2^intensity) * cycles_per_iter hashes -// Return -1 if a block is found -// Else return the hashrate in MH/s +// Given a number of cycles per iter, grind nonces will poll Sia for a block +// then do 2^intensity hashes cycles_per_iter times, checking for a successful +// hash each time +// Returns -1 if it finds a block, otherwise it returns the hash_rate of the GPU double grindNonces(int cycles_per_iter) { + // Start timing this iteration - // - // TODO: Explain the difference between the linux and non-linux stuff. + // If on Linux, use the more reliable/accurate timer, otherwise use default timer #ifdef __linux__ struct timespec begin, end; clock_gettime(CLOCK_REALTIME, &begin); @@ -145,10 +152,8 @@ double grindNonces(int cycles_per_iter) { } } - // Hashrate is inaccurate if a block was found - // - // TODO: this comment isn't terribly helpful. Also you need to document the - // ifdef __linux__ stuff again. + // If on Linux, use the more reliable/accurate timer, otherwise use default timer + // to calculate the hash rate of thie iteration #ifdef __linux__ clock_gettime(CLOCK_REALTIME, &end); double nanosecondsElapsed = 1e9 * (double)(end.tv_sec - begin.tv_sec) + (double)(end.tv_nsec - begin.tv_nsec); @@ -343,16 +348,17 @@ int main(int argc, char *argv[]) { printf("Please pass in a number following your flag (e.g. -I 22)\n"); exit(1); } + // atoi returns 0 on error intensity = atoi(argv[i]); - if (intensity == 0 && argv[i][0] != '0') { + if (intensity == 0 && argv[i][0] != '0') { // Check if atoi returned 0 because of an error printf("Invalid number passed to \'-I\'\n"); exit(1); } if(intensity < MIN_INTENSITY || intensity > MAX_INTENSITY) { printf("intensity must be between %u and %u. %u is invalid\n", MIN_INTENSITY, MAX_INTENSITY, intensity); - printf("Note that the minimum intensity is %d, and the maximum is %d.\n", MIN_INTENSITY, MAX_INTENSITY); - intensity = DEFAULT_INTENSITY; + printf("Note that the default intensity is %d\n", DEFAULT_INTENSITY); + exit(1); } printf("Intensity set to %u\n", intensity); break; @@ -361,8 +367,6 @@ int main(int argc, char *argv[]) { printf("Please pass in a number following your flag (e.g. -p 1)\n"); exit(1); } - // Again, zero return on error. Default is zero. - // I don't see a problem here. platformid = atoi(argv[i]); if (platformid == 0 && argv[i][0] != '0') { printf("Invalid number passed to \'-p\'\n"); @@ -374,7 +378,6 @@ int main(int argc, char *argv[]) { printf("Please pass in a number following your flag (e.g. -d 1)\n"); exit(1); } - // See comment for previous option. deviceidx = atoi(argv[i]); if (deviceidx == 0 && argv[i][0] != '0') { printf("Invalid number passed to \'-d\'\n"); @@ -391,6 +394,12 @@ int main(int argc, char *argv[]) { printf("Invalid number passed to \'-C\'\n"); exit(1); } + + if(cycles_per_iter < MIN_CPI || cycles_per_iter > MAX_CPI) { + printf("cycles per iter must be between %u and %u. %u is invalid\n", MIN_CPI, MAX_CPI, cycles_per_iter); + printf("Note that the default cycles per iter is %d\n", DEFAULT_CPI); + exit(1); + } printf("Cycles per iteration set to %u\n", cycles_per_iter); break; case 'P': @@ -453,10 +462,6 @@ int main(int argc, char *argv[]) { // Create Buffer Objects blockHeadermobj = clCreateBuffer(context, CL_MEM_READ_ONLY, 80 * sizeof(uint8_t), NULL, &ret); if (ret != CL_SUCCESS) { printf("failed to create blockHeadermobj buffer: %d\n", ret); exit(1); } - headerHashmobj = clCreateBuffer(context, CL_MEM_READ_WRITE, 32 * sizeof(uint8_t), NULL, &ret); - if (ret != CL_SUCCESS) { printf("failed to create headerHashmobj buffer: %d\n", ret); exit(1); } - targmobj = clCreateBuffer(context, CL_MEM_READ_ONLY, 32 * sizeof(uint8_t), NULL, &ret); - if (ret != CL_SUCCESS) { printf("failed to create targmobj buffer: %d\n", ret); exit(1); } nonceOutmobj = clCreateBuffer(context, CL_MEM_READ_WRITE, 8 * sizeof(uint8_t), NULL, &ret); if (ret != CL_SUCCESS) { printf("failed to create nonceOutmobj buffer: %d\n", ret); exit(1); } @@ -518,7 +523,7 @@ int main(int argc, char *argv[]) { // Repeat until no block is found do { hash_rate = grindNonces(cycles_per_iter); - } while (hash_rate == -1); + } while (hash_rate == -1 && !quit); if (!quit) { printf("\rMining at %.3f MH/s\t%u blocks mined", hash_rate, blocks_mined); @@ -532,8 +537,6 @@ int main(int argc, char *argv[]) { ret = clReleaseKernel(kernel); ret = clReleaseProgram(program); ret = clReleaseMemObject(blockHeadermobj); - ret = clReleaseMemObject(headerHashmobj); - ret = clReleaseMemObject(targmobj); ret = clReleaseMemObject(nonceOutmobj); ret = clReleaseCommandQueue(command_queue); ret = clReleaseContext(context); From dd4f2236967c5eaaad55883bf616fd3c5df9fc6d Mon Sep 17 00:00:00 2001 From: David Vorick Date: Fri, 19 Jun 2015 13:03:27 -0400 Subject: [PATCH 3/4] add more docstrings --- network.c | 11 +++++++++++ sia-gpu-miner.c | 17 +++++++++++++---- 2 files changed, 24 insertions(+), 4 deletions(-) diff --git a/network.c b/network.c index 082c51a..e7bc6f2 100644 --- a/network.c +++ b/network.c @@ -5,14 +5,17 @@ #include "network.h" +// TODO: document what this does. It should also probably be renamed. struct inData { uint8_t *bytes; size_t len; }; +// TODO: Is it safe to have these as global variables? char *bfw_url, *submit_url; CURL *curl; +// check_http_response int check_http_response(CURL *curl) { long http_code = 0; curl_easy_getinfo(curl, CURLINFO_RESPONSE_CODE, &http_code); @@ -23,6 +26,7 @@ int check_http_response(CURL *curl) { return 0; } +// set_port establishes the port that siad is on. void set_port(char *port) { bfw_url = malloc(29 + strlen(port)); submit_url = malloc(28 + strlen(port)); @@ -31,6 +35,8 @@ void set_port(char *port) { } // Write network data to an array of bytes +// +// TODO: nmemb is not a good name. 'in' is also probably not a good name. size_t writefunc(void *ptr, size_t size, size_t nmemb, struct inData *in) { if (in == NULL) return size*nmemb; @@ -46,6 +52,7 @@ size_t writefunc(void *ptr, size_t size, size_t nmemb, struct inData *in) { return size*nmemb; } +// init_network initializes curl networking. void init_network() { curl = curl_easy_init(); if (!curl) { @@ -53,6 +60,8 @@ void init_network() { } } +// get_header_for_work fetches a block header from siad. This block header is +// ready for nonce grinding. int get_header_for_work(uint8_t *target, uint8_t *header) { if (!curl) { fprintf(stderr, "Invalid curl object passed to get_block_for_work()\n"); @@ -91,6 +100,7 @@ int get_header_for_work(uint8_t *target, uint8_t *header) { return 0; } +// submit_header submits a block header to siad. int submit_header(uint8_t *header) { CURLcode res; @@ -111,6 +121,7 @@ int submit_header(uint8_t *header) { return check_http_response(curl); } +// free_network closes the network connection. void free_network() { curl_easy_cleanup(curl); } diff --git a/sia-gpu-miner.c b/sia-gpu-miner.c index e94ad8e..72c2dd8 100644 --- a/sia-gpu-miner.c +++ b/sia-gpu-miner.c @@ -166,6 +166,8 @@ double grindNonces(int cycles_per_iter) { return hash_rate; } +// selectOCLDevice manages opencl device selection as requested by the command +// line arguments. void selectOCLDevice(cl_platform_id *OCLPlatform, cl_device_id *OCLDevice, cl_uint platformid, cl_uint deviceidx) { cl_uint platformCount, deviceCount; cl_platform_id *platformids; @@ -243,6 +245,8 @@ void selectOCLDevice(cl_platform_id *OCLPlatform, cl_device_id *OCLDevice, cl_ui *OCLDevice = deviceids[deviceidx]; } +// printPlatformsAndDevices prints out a list of opencl platforms and devices +// that were found on the system. void printPlatformsAndDevices() { cl_uint platformCount, deviceCount; cl_platform_id *platformids; @@ -304,7 +308,9 @@ void printPlatformsAndDevices() { } free(platformids); } - + +// main reads the command line arguments and then starts the miner. The program +// will exit if there are any errors. int main(int argc, char *argv[]) { cl_platform_id platform_id = NULL; cl_device_id device_id = NULL; @@ -327,20 +333,23 @@ int main(int argc, char *argv[]) { switch (c) { case 'h': printf("\nUsage:\n\n"); + printf("\t C - cycles per iter: Number of kernel executions between Sia API calls and hash rate updates\n"); + printf("\t\tIncrease this if your miner is receiving invalid targets. Default is %u.\n", DEFAULT_CPI); + printf("\n"); printf("\t I - intensity: This is the amount of work sent to the GPU in one batch.\n"); printf("\t\tInterpretation is 2^intensity; the default is 16. Lower if GPU crashes or\n"); printf("\t\tif more desktop interactivity is desired. Raising it may improve performance.\n"); printf("\n"); + printf("\t P - port: which port to use when talking to the siad api.\n"); + printf("\n"); printf("\t p - OpenCL platform ID: Just what it says on the tin. If you're finding no GPUs,\n"); printf("\t\tyet you're sure they exist, try a value other than 0, like 1, or 2. Default is 0.\n"); printf("\n"); printf("\t d - OpenCL device ID: Self-explanatory; it's the GPU index. Note that different\n"); printf("\t\tOpenCL platforms will likely have different devices available. Default is 0.\n"); printf("\n"); - printf("\t C - cycles per iter: Number of kernel executions between Sia API calls and hash rate updates\n"); - printf("\t\tIncrease this if your miner is receiving invalid targets. Default is %u.\n", DEFAULT_CPI); - printf("\n"); printPlatformsAndDevices(); + printf("\n"); exit(0); break; case 'I': From b34ade16821982bbdf6cbc5db702c27cd20aaf9c Mon Sep 17 00:00:00 2001 From: JoshVorick Date: Fri, 19 Jun 2015 22:23:06 -0400 Subject: [PATCH 4/4] lint network.c --- network.c | 47 ++++++++++++++++++++++++----------------------- 1 file changed, 24 insertions(+), 23 deletions(-) diff --git a/network.c b/network.c index e7bc6f2..47d0b07 100644 --- a/network.c +++ b/network.c @@ -5,14 +5,16 @@ #include "network.h" -// TODO: document what this does. It should also probably be renamed. -struct inData { +// Buffer for receiving data through curl +struct inBuffer { uint8_t *bytes; - size_t len; + size_t len; // Number of bytes read in/allocated }; -// TODO: Is it safe to have these as global variables? +// URL strings for receiving and submitting blocks char *bfw_url, *submit_url; + +// CURL object to connect to siad CURL *curl; // check_http_response @@ -34,22 +36,21 @@ void set_port(char *port) { sprintf(submit_url, "localhost:%s/miner/submitheader", port); } -// Write network data to an array of bytes -// -// TODO: nmemb is not a good name. 'in' is also probably not a good name. -size_t writefunc(void *ptr, size_t size, size_t nmemb, struct inData *in) { - if (in == NULL) - return size*nmemb; - size_t new_len = size*nmemb; - in->bytes = (uint8_t*)malloc(new_len); - if (in->bytes == NULL) { +// Write network data to a buffer (inBuf) +size_t writefunc(void *ptr, size_t size, size_t num_elems, struct inBuffer *inBuf) { + if (inBuf == NULL) { + return size*num_elems; + } + size_t new_len = size*num_elems; + inBuf->bytes = (uint8_t*)malloc(new_len); + if (inBuf->bytes == NULL) { fprintf(stderr, "malloc() failed\n"); exit(EXIT_FAILURE); } - memcpy(in->bytes, ptr, size*nmemb); - in->len = new_len; + memcpy(inBuf->bytes, ptr, size*num_elems); + inBuf->len = new_len; - return size*nmemb; + return size*num_elems; } // init_network initializes curl networking. @@ -69,13 +70,13 @@ int get_header_for_work(uint8_t *target, uint8_t *header) { } CURLcode res; - struct inData in; + struct inBuffer inBuf; // Get data from siad curl_easy_reset(curl); curl_easy_setopt(curl, CURLOPT_URL, bfw_url); curl_easy_setopt(curl, CURLOPT_WRITEFUNCTION, writefunc); - curl_easy_setopt(curl, CURLOPT_WRITEDATA, &in); + curl_easy_setopt(curl, CURLOPT_WRITEDATA, &inBuf); res = curl_easy_perform(curl); if(res != CURLE_OK) { @@ -86,16 +87,16 @@ int get_header_for_work(uint8_t *target, uint8_t *header) { if (check_http_response(curl)) { return 1; } - if (in.len != 112) { - fprintf(stderr, "curl did not receive correct bytes (got %lu, expected 112)\n", in.len); + if (inBuf.len != 112) { + fprintf(stderr, "curl did not receive correct bytes (got %lu, expected 112)\n", inBuf.len); return 1; } // Copy data to return - memcpy(target, in.bytes, 32); - memcpy(header, in.bytes+32, 80); + memcpy(target, inBuf.bytes, 32); + memcpy(header, inBuf.bytes+32, 80); - free(in.bytes); + free(inBuf.bytes); return 0; }