Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

fixed incorrect hashrate on mac os and fixed some memleaks #47

Open
wants to merge 19 commits into
base: master
Choose a base branch
from
Open

fixed incorrect hashrate on mac os and fixed some memleaks #47

Show file tree
Hide file tree
Changes from 13 commits
Commits
Show all changes
19 commits
Select commit Hold shift + click to select a range
2d4c4a2
fixed mem leak and duplicat free mem
Nov 6, 2016
390939d
fixed a mem leak
Nov 6, 2016
37e1870
use fprintf instead of printf
Nov 6, 2016
669c841
fixed a commented error
Nov 7, 2016
79503ad
checkout nonceOut more precise
Nov 7, 2016
213d872
print more GPU devices info
gjwang Nov 7, 2016
6e4d82e
print more GPU devies info
gjwang Nov 9, 2016
54fd7db
fixed time consume error in APPLE OS
gjwang Nov 9, 2016
2b786a9
refactor get time function
gjwang Nov 9, 2016
0f8cff8
refortor get time function
gjwang Nov 10, 2016
baeac59
print submit_header response
gjwang Nov 10, 2016
1f952a7
fixed sleep in mac os
gjwang Nov 12, 2016
85e58f6
fixed sleep in Windows os
gjwang Nov 12, 2016
4665206
rename get_timestamp_in_us to getTimestampInUs
gjwang Nov 13, 2016
7a10556
reformat: use Tabs for indentation instead of spaces
gjwang Nov 13, 2016
c5ac9c1
reformat printPlatformsAndDevices to more readable
gjwang Nov 13, 2016
fbf742a
print submit work response more nicer
gjwang Nov 13, 2016
b5448e4
prevent mass requests from miner in case of server error
gjwang Nov 14, 2016
8d8af8d
fixed get time consume errro in windows os
gjwang Nov 15, 2016
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
63 changes: 52 additions & 11 deletions network.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,11 +2,13 @@
#include <stdio.h>
#include <string.h>

#ifdef __linux__
#if defined(__linux__) || defined(__APPLE__)
#include <unistd.h>
#endif
#else
#ifdef __WINDOWS__
#include <windows.h>
#define sleep(seconds) Sleep(seconds*1000)
#endif
#endif

#include <curl/curl.h>
Expand All @@ -20,7 +22,9 @@ struct inBuffer {
};

// URL strings for receiving and submitting blocks
char *bfw_url, *submit_url;
#define MAX_HOST_LEN (2048)
char bfw_url[MAX_HOST_LEN];
char submit_url[MAX_HOST_LEN];

// CURL object to connect to siad
CURL *curl;
Expand All @@ -41,12 +45,46 @@ int check_http_response(CURL *curl) {

// set_host establishes the hostname and port that siad is on.
void set_host(char *host, char *port) {
bfw_url = malloc(21 + strlen(host) + strlen(port));
submit_url = malloc(20 + strlen(host) + strlen(port));
size_t host_len = 21 + strlen(host) + strlen(port);
if (host_len >= MAX_HOST_LEN) {
fprintf(stderr, "Error: host is over of size, host_len=%zu > MAX_HOST_LEN=%d\n", host_len, MAX_HOST_LEN);
exit(1);
}

sprintf(bfw_url, "%s%s/miner/header", host, port);
sprintf(submit_url, "%s%s/miner/header", host, port);
}

static void printMem(const uint8_t *mem, size_t size, const char *format, int num_in_line){
if (format == NULL) {
format = "%02x ";
}

if (num_in_line == 0) {
num_in_line = 16;
}

for (int i=0; i<size; i++) {
printf(format, mem[i]);
if ((i+1)%num_in_line == 0) {
printf("\n");
}
}
printf("\n");
}

//static void printMemHex(const uint8_t *mem, size_t size){
// printMem(mem, size, "%02x ", 16);
//}
//
//static void printMemChar(const uint8_t *mem, size_t size){
// printMem(mem, size, "%c", 16);
//}
//
//static void printMemDec(const uint8_t *mem, size_t size){
// printMem(mem, size, "%d ", 16);
//}

// 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) {
Expand Down Expand Up @@ -92,12 +130,7 @@ int get_header_for_work(uint8_t *target, uint8_t *header) {
fprintf(stderr, "Failed to get header from %s, curl_easy_perform() failed: %s\n", bfw_url, curl_easy_strerror(res));
fprintf(stderr, "Are you sure that siad is running?\n");
// Pause in order to prevent spamming the console
#ifdef __linux__
sleep(3); // 3 seconds
#endif
#ifdef __WINDOWS__
Sleep(3000); // 3 seconds
#endif
}

if (check_http_response(curl)) {
Expand All @@ -120,6 +153,7 @@ int get_header_for_work(uint8_t *target, uint8_t *header) {
// submit_header submits a block header to siad.
int submit_header(uint8_t *header) {
CURLcode res;
struct inBuffer inBuf;

curl_easy_reset(curl);
curl_easy_setopt(curl, CURLOPT_USERAGENT, "Sia-Agent");
Expand All @@ -129,13 +163,20 @@ int submit_header(uint8_t *header) {
curl_easy_setopt(curl, CURLOPT_POSTFIELDS, header);
// Prevent printing to stdout
curl_easy_setopt(curl, CURLOPT_WRITEFUNCTION, writefunc);
curl_easy_setopt(curl, CURLOPT_WRITEDATA, NULL);
curl_easy_setopt(curl, CURLOPT_WRITEDATA, &inBuf);

res = curl_easy_perform(curl);
if (res != CURLE_OK) {
fprintf(stderr, "Failed to submit block, curl_easy_perform() failed: %s\n", curl_easy_strerror(res));
return 1;
}

if (inBuf.bytes) {
printMem(inBuf.bytes, inBuf.len, "%c", INT_MAX);

free(inBuf.bytes);
}

return check_http_response(curl);
}

Expand Down
82 changes: 58 additions & 24 deletions sia-gpu-miner.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,7 @@
// OpenCL headers are different for Apple.
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#include <sys/time.h>
#else
#include <CL/cl.h>
#endif
Expand Down Expand Up @@ -70,19 +71,22 @@ void quitSignal(int unused) {
printf("\nCaught kill signal, quitting...\n");
}

inline uint64_t get_timestamp_in_us()
{
#if defined(__linux__) || defined(__APPLE__)
struct timeval now;
gettimeofday(&now, NULL);
return now.tv_sec * 1000000LL + now.tv_usec;
#else
return clock()/ CLOCKS_PER_SEC * 1000000LL;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

seems like there are a lot of ways to get the current timestamp. This post describes the differences between various functions. I also found this gist by jbenet which shows how to get the current time on OS X. Not sure what the situation is like for Windows.

#endif
}

// 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.
#ifdef __linux__
struct timespec begin, end;
clock_gettime(CLOCK_REALTIME, &begin);
#else
clock_t startTime = clock();
#endif

uint8_t blockHeader[80];
uint8_t target[32] = {255};
uint8_t nonceOut[8] = {0};
Expand Down Expand Up @@ -112,6 +116,9 @@ double grindNonces(int cycles_per_iter) {
blockHeader[i + 32] = target[7-i];
}

// Start timing this iteration.
int64_t beginInUs = get_timestamp_in_us();

// By doing a bunch of low intensity calls, we prevent freezing
// By splitting them up inside this function, we also avoid calling
// get_block_for_work too often.
Expand Down Expand Up @@ -141,7 +148,7 @@ double grindNonces(int cycles_per_iter) {
if (ret != CL_SUCCESS) {
printf("failed to read nonce from buffer: %d\n", ret); exit(1);
}
if (nonceOut[0] != 0) {
if (*(uint64_t *)&nonceOut != 0) {
// Copy nonce to header.
memcpy(blockHeader+32, nonceOut, 8);
if (!submit_header(blockHeader)) {
Expand All @@ -153,17 +160,12 @@ double grindNonces(int cycles_per_iter) {
}

// Get the time elapsed this function.
#ifdef __linux__
clock_gettime(CLOCK_REALTIME, &end);
double nsElapsed = 1e9 * (double)(end.tv_sec - begin.tv_sec) + (double)(end.tv_nsec - begin.tv_nsec);
double run_time_seconds = nsElapsed * 1e-9;
#else
double run_time_seconds = (double)(clock() - startTime) / CLOCKS_PER_SEC;
#endif

// Calculate the hash rate of thie iteration.
double hash_rate = cycles_per_iter * global_item_size / (run_time_seconds*1000000);
return hash_rate;
int64_t endInUs = get_timestamp_in_us();
int64_t run_time_us = endInUs - beginInUs;

// Calculate the hash rate of thie iteration.
double hash_rate = cycles_per_iter * global_item_size / (double)run_time_us;
return hash_rate;
}

// selectOCLDevice manages opencl device selection as requested by the command
Expand Down Expand Up @@ -243,8 +245,17 @@ void selectOCLDevice(cl_platform_id *OCLPlatform, cl_device_id *OCLDevice, cl_ui
// Done. Return the platform ID and device ID object desired, free lists, and return.
*OCLPlatform = platformids[platformid];
*OCLDevice = deviceids[deviceidx];

if (platformids) {
free(platformids);
}

if (deviceids) {
free(deviceids);
}
}


// printPlatformsAndDevices prints out a list of opencl platforms and devices
// that were found on the system.
void printPlatformsAndDevices() {
Expand Down Expand Up @@ -302,10 +313,33 @@ void printPlatformsAndDevices() {
ret = clGetDeviceInfo(deviceids[j], CL_DEVICE_NAME, 80, str, NULL);
if (ret != CL_SUCCESS) {
printf("\tError while getting device info.\n");
free(deviceids);
continue;
}
printf("\tDevice %d: %s\n", j, str);
} else {
printf("\tDevice %d: %s\n", j, str);
}
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

if we can't read the device name, does it make sense to continue printing the endianness, alloc size, etc.? I figure if one fails, the rest will probably fail too. I'm not very familiar of OpenCL, though; maybe not all devices have names.

Copy link
Member

@lukechampine lukechampine Nov 12, 2016 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Also, if we can assume that all devices have names, then the rest of the info should be printed with a double tab instead of repeating the device id. For example:

Found 1 platform(s) on your computer.
    Device 0 (MyDevice):
        Endianness:    Little-endian
        Address space: 64-bit

instead of:

Found 1 platform(s) on your computer.
    Device 0: MyDevice
    Device 0: Endianness: Little-endian
    Device 0: Address space: 64-bit

I suppose having a uniform prefix is handy if you want to pipe the output to grep or something, but I imagine you only need to do that if you have lots of devices.


cl_bool isLitlle = CL_TRUE;
ret = clGetDeviceInfo(deviceids[j], CL_DEVICE_ENDIAN_LITTLE, sizeof(isLitlle), &isLitlle, NULL);
if (ret != CL_SUCCESS) {
printf("\tError while getting device CL_DEVICE_ENDIAN_LITTLE.\n");
} else {
printf("\tDevice %d: isLitlle: %d\n", j, isLitlle);
}
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

could be more descriptive here. How about:

else if (isLittle) {
    printf("\t\tEndianness: Little-endian\n");
} else {
    printf("\t\tEndianness: Big-endian\n");
}

also, please ensure that your new code is properly indented.


size_t bits = 0;
ret = clGetDeviceInfo(deviceids[j], CL_DEVICE_ADDRESS_BITS, sizeof(bits), &bits, NULL);
if (ret != CL_SUCCESS) {
printf("\tError while getting device CL_DEVICE_ADDRESS_BITS.\n");
} else {
printf("\tDevice %d: bits: %zu\n", j, bits);
}

size_t maxMemSize = 0;
ret = clGetDeviceInfo(deviceids[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(maxMemSize), &maxMemSize, NULL);
if (ret != CL_SUCCESS) {
printf("\tError while getting device CL_DEVICE_MAX_MEM_ALLOC_SIZE.\n");
} else {
printf("\tDevice %d: maxMemSize: %zu\n", j, maxMemSize);
}
}
free(deviceids);
}
Expand Down
2 changes: 1 addition & 1 deletion sia-gpu-miner.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@ __constant static const uchar blake2b_sigma[12][16] = {
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } ,
{ 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 } };

// Target is passed in via headerIn[32 - 29]
// Target is passed in via headerIn[32 - 39]
__kernel void nonceGrind(__global ulong *headerIn, __global ulong *nonceOut) {
ulong target = headerIn[4];
ulong m[16] = { headerIn[0], headerIn[1],
Expand Down