GOSTSec
/
sgminer
mirror of https://github.com/GOSTSec/sgminer
1
0
Fork 0
Code Issues Releases Wiki Activity
OpenCL GPU miner
You can not select more than 25 topics Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
2474 Commits
12 Branches
10 Tags
44 MiB
 
 
 
 
 
Tree: aabc723326
Branches Tags
${ item.name }
Create tag ${ searchTerm }
Create branch ${ searchTerm }
from 'aabc723326'
${ noResults }
sgminer/driver-opencl.c

1522 lines
36 KiB
Raw Blame History

/*
* Copyright 2011-2012 Con Kolivas
* Copyright 2011-2012 Luke Dashjr
* Copyright 2010 Jeff Garzik
*
* 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.
*/
#include "config.h"
#ifdef HAVE_CURSES
#include <curses.h>
#endif
#include <string.h>
#include <stdbool.h>
#include <stdint.h>
#include <sys/types.h>
#ifndef WIN32
#include <sys/resource.h>
#endif
#include <ccan/opt/opt.h>
#include "compat.h"
#include "miner.h"
#include "driver-opencl.h"
#include "findnonce.h"
#include "ocl.h"
#include "adl.h"
/* TODO: cleanup externals ********************/
#ifdef HAVE_CURSES
extern WINDOW *mainwin, *statuswin, *logwin;
extern void enable_curses(void);
#endif
extern int mining_threads;
extern double total_secs;
extern int opt_g_threads;
extern bool ping;
extern bool opt_loginput;
extern char *opt_kernel_path;
extern int gpur_thr_id;
extern bool opt_noadl;
extern bool have_opencl;
extern void *miner_thread(void *userdata);
extern int dev_from_id(int thr_id);
extern void tailsprintf(char *f, const char *fmt, ...);
extern void wlog(const char *f, ...);
extern void decay_time(double *f, double fadd);
/**********************************************/
#ifdef HAVE_ADL
extern float gpu_temp(int gpu);
extern int gpu_fanspeed(int gpu);
extern int gpu_fanpercent(int gpu);
#endif
#ifdef HAVE_OPENCL
char *set_vector(char *arg)
{
int i, val = 0, device = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set vector";
val = atoi(nextptr);
if (val != 1 && val != 2 && val != 4)
return "Invalid value passed to set_vector";
gpus[device++].vwidth = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atoi(nextptr);
if (val != 1 && val != 2 && val != 4)
return "Invalid value passed to set_vector";
gpus[device++].vwidth = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++)
gpus[i].vwidth = gpus[0].vwidth;
}
return NULL;
}
char *set_worksize(char *arg)
{
int i, val = 0, device = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set work size";
val = atoi(nextptr);
if (val < 1 || val > 9999)
return "Invalid value passed to set_worksize";
gpus[device++].work_size = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atoi(nextptr);
if (val < 1 || val > 9999)
return "Invalid value passed to set_worksize";
gpus[device++].work_size = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++)
gpus[i].work_size = gpus[0].work_size;
}
return NULL;
}
static enum cl_kernels select_kernel(char *arg)
{
if (!strcmp(arg, "diablo"))
return KL_DIABLO;
if (!strcmp(arg, "diakgcn"))
return KL_DIAKGCN;
if (!strcmp(arg, "poclbm"))
return KL_POCLBM;
if (!strcmp(arg, "phatk"))
return KL_PHATK;
#ifdef USE_SCRYPT
if (!strcmp(arg, "scrypt"))
return KL_SCRYPT;
#endif
return KL_NONE;
}
char *set_kernel(char *arg)
{
enum cl_kernels kern;
int i, device = 0;
char *nextptr;
if (opt_scrypt)
return "Cannot use sha256 kernel with scrypt";
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set kernel";
kern = select_kernel(nextptr);
if (kern == KL_NONE)
return "Invalid parameter to set_kernel";
gpus[device++].kernel = kern;
while ((nextptr = strtok(NULL, ",")) != NULL) {
kern = select_kernel(nextptr);
if (kern == KL_NONE)
return "Invalid parameter to set_kernel";
gpus[device++].kernel = kern;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++)
gpus[i].kernel = gpus[0].kernel;
}
return NULL;
}
#endif
#ifdef HAVE_ADL
/* This function allows us to map an adl device to an opencl device for when
* simple enumeration has failed to match them. */
char *set_gpu_map(char *arg)
{
int val1 = 0, val2 = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set gpu map";
if (sscanf(arg, "%d:%d", &val1, &val2) != 2)
return "Invalid description for map pair";
if (val1 < 0 || val1 > MAX_GPUDEVICES || val2 < 0 || val2 > MAX_GPUDEVICES)
return "Invalid value passed to set_gpu_map";
gpus[val1].virtual_adl = val2;
gpus[val1].mapped = true;
while ((nextptr = strtok(NULL, ",")) != NULL) {
if (sscanf(nextptr, "%d:%d", &val1, &val2) != 2)
return "Invalid description for map pair";
if (val1 < 0 || val1 > MAX_GPUDEVICES || val2 < 0 || val2 > MAX_GPUDEVICES)
return "Invalid value passed to set_gpu_map";
gpus[val1].virtual_adl = val2;
gpus[val1].mapped = true;
}
return NULL;
}
void get_intrange(char *arg, int *val1, int *val2)
{
if (sscanf(arg, "%d-%d", val1, val2) == 1) {
*val2 = *val1;
*val1 = 0;
}
}
char *set_gpu_engine(char *arg)
{
int i, val1 = 0, val2 = 0, device = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set gpu engine";
get_intrange(nextptr, &val1, &val2);
if (val1 < 0 || val1 > 9999 || val2 < 0 || val2 > 9999)
return "Invalid value passed to set_gpu_engine";
gpus[device].min_engine = val1;
gpus[device].gpu_engine = val2;
device++;
while ((nextptr = strtok(NULL, ",")) != NULL) {
get_intrange(nextptr, &val1, &val2);
if (val1 < 0 || val1 > 9999 || val2 < 0 || val2 > 9999)
return "Invalid value passed to set_gpu_engine";
gpus[device].min_engine = val1;
gpus[device].gpu_engine = val2;
device++;
}
if (device == 1) {
for (i = 1; i < MAX_GPUDEVICES; i++) {
gpus[i].min_engine = gpus[0].min_engine;
gpus[i].gpu_engine = gpus[0].gpu_engine;
}
}
return NULL;
}
char *set_gpu_fan(char *arg)
{
int i, val1 = 0, val2 = 0, device = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set gpu fan";
get_intrange(nextptr, &val1, &val2);
if (val1 < 0 || val1 > 100 || val2 < 0 || val2 > 100)
return "Invalid value passed to set_gpu_fan";
gpus[device].min_fan = val1;
gpus[device].gpu_fan = val2;
device++;
while ((nextptr = strtok(NULL, ",")) != NULL) {
get_intrange(nextptr, &val1, &val2);
if (val1 < 0 || val1 > 100 || val2 < 0 || val2 > 100)
return "Invalid value passed to set_gpu_fan";
gpus[device].min_fan = val1;
gpus[device].gpu_fan = val2;
device++;
}
if (device == 1) {
for (i = 1; i < MAX_GPUDEVICES; i++) {
gpus[i].min_fan = gpus[0].min_fan;
gpus[i].gpu_fan = gpus[0].gpu_fan;
}
}
return NULL;
}
char *set_gpu_memclock(char *arg)
{
int i, val = 0, device = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set gpu memclock";
val = atoi(nextptr);
if (val < 0 || val >= 9999)
return "Invalid value passed to set_gpu_memclock";
gpus[device++].gpu_memclock = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atoi(nextptr);
if (val < 0 || val >= 9999)
return "Invalid value passed to set_gpu_memclock";
gpus[device++].gpu_memclock = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++)
gpus[i].gpu_memclock = gpus[0].gpu_memclock;
}
return NULL;
}
char *set_gpu_memdiff(char *arg)
{
int i, val = 0, device = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set gpu memdiff";
val = atoi(nextptr);
if (val < -9999 || val > 9999)
return "Invalid value passed to set_gpu_memdiff";
gpus[device++].gpu_memdiff = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atoi(nextptr);
if (val < -9999 || val > 9999)
return "Invalid value passed to set_gpu_memdiff";
gpus[device++].gpu_memdiff = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++)
gpus[i].gpu_memdiff = gpus[0].gpu_memdiff;
}
return NULL;
}
char *set_gpu_powertune(char *arg)
{
int i, val = 0, device = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set gpu powertune";
val = atoi(nextptr);
if (val < -99 || val > 99)
return "Invalid value passed to set_gpu_powertune";
gpus[device++].gpu_powertune = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atoi(nextptr);
if (val < -99 || val > 99)
return "Invalid value passed to set_gpu_powertune";
gpus[device++].gpu_powertune = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++)
gpus[i].gpu_powertune = gpus[0].gpu_powertune;
}
return NULL;
}
char *set_gpu_vddc(char *arg)
{
int i, device = 0;
float val = 0;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set gpu vddc";
val = atof(nextptr);
if (val < 0 || val >= 9999)
return "Invalid value passed to set_gpu_vddc";
gpus[device++].gpu_vddc = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atof(nextptr);
if (val < 0 || val >= 9999)
return "Invalid value passed to set_gpu_vddc";
gpus[device++].gpu_vddc = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++)
gpus[i].gpu_vddc = gpus[0].gpu_vddc;
}
return NULL;
}
char *set_temp_overheat(char *arg)
{
int i, val = 0, device = 0, *to;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set temp overheat";
val = atoi(nextptr);
if (val < 0 || val > 200)
return "Invalid value passed to set temp overheat";
to = &gpus[device++].adl.overtemp;
*to = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atoi(nextptr);
if (val < 0 || val > 200)
return "Invalid value passed to set temp overheat";
to = &gpus[device++].adl.overtemp;
*to = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++) {
to = &gpus[i].adl.overtemp;
*to = val;
}
}
return NULL;
}
char *set_temp_target(char *arg)
{
int i, val = 0, device = 0, *tt;
char *nextptr;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set temp target";
val = atoi(nextptr);
if (val < 0 || val > 200)
return "Invalid value passed to set temp target";
tt = &gpus[device++].adl.targettemp;
*tt = val;
while ((nextptr = strtok(NULL, ",")) != NULL) {
val = atoi(nextptr);
if (val < 0 || val > 200)
return "Invalid value passed to set temp target";
tt = &gpus[device++].adl.targettemp;
*tt = val;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++) {
tt = &gpus[i].adl.targettemp;
*tt = val;
}
}
return NULL;
}
#endif
#ifdef HAVE_OPENCL
char *set_intensity(char *arg)
{
int i, device = 0, *tt;
char *nextptr, val = 0;
nextptr = strtok(arg, ",");
if (nextptr == NULL)
return "Invalid parameters for set intensity";
if (!strncasecmp(nextptr, "d", 1))
gpus[device].dynamic = true;
else {
gpus[device].dynamic = false;
val = atoi(nextptr);
if (val < MIN_INTENSITY || val > MAX_INTENSITY)
return "Invalid value passed to set intensity";
tt = &gpus[device].intensity;
*tt = val;
}
device++;
while ((nextptr = strtok(NULL, ",")) != NULL) {
if (!strncasecmp(nextptr, "d", 1))
gpus[device].dynamic = true;
else {
gpus[device].dynamic = false;
val = atoi(nextptr);
if (val < MIN_INTENSITY || val > MAX_INTENSITY)
return "Invalid value passed to set intensity";
tt = &gpus[device].intensity;
*tt = val;
}
device++;
}
if (device == 1) {
for (i = device; i < MAX_GPUDEVICES; i++) {
gpus[i].dynamic = gpus[0].dynamic;
gpus[i].intensity = gpus[0].intensity;
}
}
return NULL;
}
#endif
#ifdef HAVE_OPENCL
struct device_api opencl_api;
char *print_ndevs_and_exit(int *ndevs)
{
opt_log_output = true;
opencl_api.api_detect();
clear_adl(*ndevs);
applog(LOG_INFO, "%i GPU devices max detected", *ndevs);
exit(*ndevs);
}
#endif
struct cgpu_info gpus[MAX_GPUDEVICES]; /* Maximum number apparently possible */
struct cgpu_info *cpus;
#ifdef HAVE_OPENCL
/* In dynamic mode, only the first thread of each device will be in use.
* This potentially could start a thread that was stopped with the start-stop
* options if one were to disable dynamic from the menu on a paused GPU */
void pause_dynamic_threads(int gpu)
{
struct cgpu_info *cgpu = &gpus[gpu];
int i;
for (i = 1; i < cgpu->threads; i++) {
struct thr_info *thr = &thr_info[i];
if (!thr->pause && cgpu->dynamic) {
applog(LOG_WARNING, "Disabling extra threads due to dynamic mode.");
applog(LOG_WARNING, "Tune dynamic intensity with --gpu-dyninterval");
}
thr->pause = cgpu->dynamic;
if (!cgpu->dynamic && cgpu->deven != DEV_DISABLED)
tq_push(thr->q, &ping);
}
}
struct device_api opencl_api;
#endif /* HAVE_OPENCL */
#if defined(HAVE_OPENCL) && defined(HAVE_CURSES)
void manage_gpu(void)
{
struct thr_info *thr;
int selected, gpu, i;
char checkin[40];
char input;
if (!opt_g_threads)
return;
opt_loginput = true;
immedok(logwin, true);
clear_logwin();
retry:
for (gpu = 0; gpu < nDevs; gpu++) {
struct cgpu_info *cgpu = &gpus[gpu];
wlog("GPU %d: %.1f / %.1f Mh/s | A:%d R:%d HW:%d U:%.2f/m I:%d\n",
gpu, cgpu->rolling, cgpu->total_mhashes / total_secs,
cgpu->accepted, cgpu->rejected, cgpu->hw_errors,
cgpu->utility, cgpu->intensity);
#ifdef HAVE_ADL
if (gpus[gpu].has_adl) {
int engineclock = 0, memclock = 0, activity = 0, fanspeed = 0, fanpercent = 0, powertune = 0;
float temp = 0, vddc = 0;
if (gpu_stats(gpu, &temp, &engineclock, &memclock, &vddc, &activity, &fanspeed, &fanpercent, &powertune)) {
char logline[255];
strcpy(logline, ""); // In case it has no data
if (temp != -1)
sprintf(logline, "%.1f C ", temp);
if (fanspeed != -1 || fanpercent != -1) {
tailsprintf(logline, "F: ");
if (fanpercent != -1)
tailsprintf(logline, "%d%% ", fanpercent);
if (fanspeed != -1)
tailsprintf(logline, "(%d RPM) ", fanspeed);
tailsprintf(logline, " ");
}
if (engineclock != -1)
tailsprintf(logline, "E: %d MHz ", engineclock);
if (memclock != -1)
tailsprintf(logline, "M: %d Mhz ", memclock);
if (vddc != -1)
tailsprintf(logline, "V: %.3fV ", vddc);
if (activity != -1)
tailsprintf(logline, "A: %d%% ", activity);
if (powertune != -1)
tailsprintf(logline, "P: %d%%", powertune);
tailsprintf(logline, "\n");
wlog(logline);
}
}
#endif
wlog("Last initialised: %s\n", cgpu->init);
wlog("Intensity: ");
if (gpus[gpu].dynamic)
wlog("Dynamic (only one thread in use)\n");
else
wlog("%d\n", gpus[gpu].intensity);
for (i = 0; i < mining_threads; i++) {
thr = &thr_info[i];
if (thr->cgpu != cgpu)
continue;
get_datestamp(checkin, &thr->last);
wlog("Thread %d: %.1f Mh/s %s ", i, thr->rolling, cgpu->deven != DEV_DISABLED ? "Enabled" : "Disabled");
switch (cgpu->status) {
default:
case LIFE_WELL:
wlog("ALIVE");
break;
case LIFE_SICK:
wlog("SICK reported in %s", checkin);
break;
case LIFE_DEAD:
wlog("DEAD reported in %s", checkin);
break;
case LIFE_INIT:
case LIFE_NOSTART:
wlog("Never started");
break;
}
if (thr->pause)
wlog(" paused");
wlog("\n");
}
wlog("\n");
}
wlogprint("[E]nable [D]isable [I]ntensity [R]estart GPU %s\n",adl_active ? "[C]hange settings" : "");
wlogprint("Or press any other key to continue\n");
input = getch();
if (nDevs == 1)
selected = 0;
else
selected = -1;
if (!strncasecmp(&input, "e", 1)) {
struct cgpu_info *cgpu;
if (selected)
selected = curses_int("Select GPU to enable");
if (selected < 0 || selected >= nDevs) {
wlogprint("Invalid selection\n");
goto retry;
}
if (gpus[selected].deven != DEV_DISABLED) {
wlogprint("Device already enabled\n");
goto retry;
}
gpus[selected].deven = DEV_ENABLED;
for (i = 0; i < mining_threads; ++i) {
thr = &thr_info[i];
cgpu = thr->cgpu;
if (cgpu->api != &opencl_api)
continue;
if (dev_from_id(i) != selected)
continue;
if (cgpu->status != LIFE_WELL) {
wlogprint("Must restart device before enabling it");
goto retry;
}
applog(LOG_DEBUG, "Pushing ping to thread %d", thr->id);
tq_push(thr->q, &ping);
}
goto retry;
} if (!strncasecmp(&input, "d", 1)) {
if (selected)
selected = curses_int("Select GPU to disable");
if (selected < 0 || selected >= nDevs) {
wlogprint("Invalid selection\n");
goto retry;
}
if (gpus[selected].deven == DEV_DISABLED) {
wlogprint("Device already disabled\n");
goto retry;
}
gpus[selected].deven = DEV_DISABLED;
goto retry;
} else if (!strncasecmp(&input, "i", 1)) {
int intensity;
char *intvar;
if (selected)
selected = curses_int("Select GPU to change intensity on");
if (selected < 0 || selected >= nDevs) {
wlogprint("Invalid selection\n");
goto retry;
}
intvar = curses_input("Set GPU scan intensity (d or " _MIN_INTENSITY_STR " -> " _MAX_INTENSITY_STR ")");
if (!intvar) {
wlogprint("Invalid input\n");
goto retry;
}
if (!strncasecmp(intvar, "d", 1)) {
wlogprint("Dynamic mode enabled on gpu %d\n", selected);
gpus[selected].dynamic = true;
pause_dynamic_threads(selected);
free(intvar);
goto retry;
}
intensity = atoi(intvar);
free(intvar);
if (intensity < MIN_INTENSITY || intensity > MAX_INTENSITY) {
wlogprint("Invalid selection\n");
goto retry;
}
gpus[selected].dynamic = false;
gpus[selected].intensity = intensity;
wlogprint("Intensity on gpu %d set to %d\n", selected, intensity);
pause_dynamic_threads(selected);
goto retry;
} else if (!strncasecmp(&input, "r", 1)) {
if (selected)
selected = curses_int("Select GPU to attempt to restart");
if (selected < 0 || selected >= nDevs) {
wlogprint("Invalid selection\n");
goto retry;
}
wlogprint("Attempting to restart threads of GPU %d\n", selected);
reinit_device(&gpus[selected]);
goto retry;
} else if (adl_active && (!strncasecmp(&input, "c", 1))) {
if (selected)
selected = curses_int("Select GPU to change settings on");
if (selected < 0 || selected >= nDevs) {
wlogprint("Invalid selection\n");
goto retry;
}
change_gpusettings(selected);
goto retry;
} else
clear_logwin();
immedok(logwin, false);
opt_loginput = false;
}
#else
void manage_gpu(void)
{
}
#endif
#ifdef HAVE_OPENCL
static _clState *clStates[MAX_GPUDEVICES];
#define CL_SET_BLKARG(blkvar) status |= clSetKernelArg(*kernel, num++, sizeof(uint), (void *)&blk->blkvar)
#define CL_SET_ARG(var) status |= clSetKernelArg(*kernel, num++, sizeof(var), (void *)&var)
#define CL_SET_VARG(args, var) status |= clSetKernelArg(*kernel, num++, args * sizeof(uint), (void *)var)
static cl_int queue_poclbm_kernel(_clState *clState, dev_blk_ctx *blk, cl_uint threads)
{
cl_kernel *kernel = &clState->kernel;
unsigned int num = 0;
cl_int status = 0;
CL_SET_BLKARG(ctx_a);
CL_SET_BLKARG(ctx_b);
CL_SET_BLKARG(ctx_c);
CL_SET_BLKARG(ctx_d);
CL_SET_BLKARG(ctx_e);
CL_SET_BLKARG(ctx_f);
CL_SET_BLKARG(ctx_g);
CL_SET_BLKARG(ctx_h);
CL_SET_BLKARG(cty_b);
CL_SET_BLKARG(cty_c);
CL_SET_BLKARG(cty_f);
CL_SET_BLKARG(cty_g);
CL_SET_BLKARG(cty_h);
if (!clState->goffset) {
cl_uint vwidth = clState->vwidth;
uint *nonces = alloca(sizeof(uint) * vwidth);
unsigned int i;
for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + (i * threads);
CL_SET_VARG(vwidth, nonces);
}
CL_SET_BLKARG(fW0);
CL_SET_BLKARG(fW1);
CL_SET_BLKARG(fW2);
CL_SET_BLKARG(fW3);
CL_SET_BLKARG(fW15);
CL_SET_BLKARG(fW01r);
CL_SET_BLKARG(D1A);
CL_SET_BLKARG(C1addK5);
CL_SET_BLKARG(B1addK6);
CL_SET_BLKARG(W16addK16);
CL_SET_BLKARG(W17addK17);
CL_SET_BLKARG(PreVal4addT1);
CL_SET_BLKARG(PreVal0);
CL_SET_ARG(clState->outputBuffer);
return status;
}
static cl_int queue_phatk_kernel(_clState *clState, dev_blk_ctx *blk,
__maybe_unused cl_uint threads)
{
cl_kernel *kernel = &clState->kernel;
cl_uint vwidth = clState->vwidth;
unsigned int i, num = 0;
cl_int status = 0;
uint *nonces;
CL_SET_BLKARG(ctx_a);
CL_SET_BLKARG(ctx_b);
CL_SET_BLKARG(ctx_c);
CL_SET_BLKARG(ctx_d);
CL_SET_BLKARG(ctx_e);
CL_SET_BLKARG(ctx_f);
CL_SET_BLKARG(ctx_g);
CL_SET_BLKARG(ctx_h);
CL_SET_BLKARG(cty_b);
CL_SET_BLKARG(cty_c);
CL_SET_BLKARG(cty_d);
CL_SET_BLKARG(cty_f);
CL_SET_BLKARG(cty_g);
CL_SET_BLKARG(cty_h);
nonces = alloca(sizeof(uint) * vwidth);
for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + i;
CL_SET_VARG(vwidth, nonces);
CL_SET_BLKARG(W16);
CL_SET_BLKARG(W17);
CL_SET_BLKARG(PreVal4_2);
CL_SET_BLKARG(PreVal0);
CL_SET_BLKARG(PreW18);
CL_SET_BLKARG(PreW19);
CL_SET_BLKARG(PreW31);
CL_SET_BLKARG(PreW32);
CL_SET_ARG(clState->outputBuffer);
return status;
}
static cl_int queue_diakgcn_kernel(_clState *clState, dev_blk_ctx *blk,
__maybe_unused cl_uint threads)
{
cl_kernel *kernel = &clState->kernel;
unsigned int num = 0;
cl_int status = 0;
if (!clState->goffset) {
cl_uint vwidth = clState->vwidth;
uint *nonces = alloca(sizeof(uint) * vwidth);
unsigned int i;
for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + i;
CL_SET_VARG(vwidth, nonces);
}
CL_SET_BLKARG(PreVal0);
CL_SET_BLKARG(PreVal4_2);
CL_SET_BLKARG(cty_h);
CL_SET_BLKARG(D1A);
CL_SET_BLKARG(cty_b);
CL_SET_BLKARG(cty_c);
CL_SET_BLKARG(cty_f);
CL_SET_BLKARG(cty_g);
CL_SET_BLKARG(C1addK5);
CL_SET_BLKARG(B1addK6);
CL_SET_BLKARG(PreVal0addK7);
CL_SET_BLKARG(W16addK16);
CL_SET_BLKARG(W17addK17);
CL_SET_BLKARG(PreW18);
CL_SET_BLKARG(PreW19);
CL_SET_BLKARG(W16);
CL_SET_BLKARG(W17);
CL_SET_BLKARG(PreW31);
CL_SET_BLKARG(PreW32);
CL_SET_BLKARG(ctx_a);
CL_SET_BLKARG(ctx_b);
CL_SET_BLKARG(ctx_c);
CL_SET_BLKARG(ctx_d);
CL_SET_BLKARG(ctx_e);
CL_SET_BLKARG(ctx_f);
CL_SET_BLKARG(ctx_g);
CL_SET_BLKARG(ctx_h);
CL_SET_BLKARG(zeroA);
CL_SET_BLKARG(zeroB);
CL_SET_BLKARG(oneA);
CL_SET_BLKARG(twoA);
CL_SET_BLKARG(threeA);
CL_SET_BLKARG(fourA);
CL_SET_BLKARG(fiveA);
CL_SET_BLKARG(sixA);
CL_SET_BLKARG(sevenA);
CL_SET_ARG(clState->outputBuffer);
return status;
}
static cl_int queue_diablo_kernel(_clState *clState, dev_blk_ctx *blk, cl_uint threads)
{
cl_kernel *kernel = &clState->kernel;
unsigned int num = 0;
cl_int status = 0;
if (!clState->goffset) {
cl_uint vwidth = clState->vwidth;
uint *nonces = alloca(sizeof(uint) * vwidth);
unsigned int i;
for (i = 0; i < vwidth; i++)
nonces[i] = blk->nonce + (i * threads);
CL_SET_VARG(vwidth, nonces);
}
CL_SET_BLKARG(PreVal0);
CL_SET_BLKARG(PreVal0addK7);
CL_SET_BLKARG(PreVal4addT1);
CL_SET_BLKARG(PreW18);
CL_SET_BLKARG(PreW19);
CL_SET_BLKARG(W16);
CL_SET_BLKARG(W17);
CL_SET_BLKARG(W16addK16);
CL_SET_BLKARG(W17addK17);
CL_SET_BLKARG(PreW31);
CL_SET_BLKARG(PreW32);
CL_SET_BLKARG(D1A);
CL_SET_BLKARG(cty_b);
CL_SET_BLKARG(cty_c);
CL_SET_BLKARG(cty_h);
CL_SET_BLKARG(cty_f);
CL_SET_BLKARG(cty_g);
CL_SET_BLKARG(C1addK5);
CL_SET_BLKARG(B1addK6);
CL_SET_BLKARG(ctx_a);
CL_SET_BLKARG(ctx_b);
CL_SET_BLKARG(ctx_c);
CL_SET_BLKARG(ctx_d);
CL_SET_BLKARG(ctx_e);
CL_SET_BLKARG(ctx_f);
CL_SET_BLKARG(ctx_g);
CL_SET_BLKARG(ctx_h);
CL_SET_ARG(clState->outputBuffer);
return status;
}
#ifdef USE_SCRYPT
static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, cl_uint threads)
{
cl_uint4 *midstate = (cl_uint4 *)blk->midstate;
cl_kernel *kernel = &clState->kernel;
unsigned int num = 0;
cl_int status = 0;
int i;
CL_SET_ARG(clState->CLbuffer0);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(clState->padbuffer8);
CL_SET_VARG(4, &midstate[0]);
CL_SET_VARG(4, &midstate[16]);
#if 0
clSetKernelArg(clState->kernel,0,sizeof(cl_mem), &clState->CLbuffer[0]);
clSetKernelArg(clState->kernel,1,sizeof(cl_mem), &clState->CLbuffer[1]);
clSetKernelArg(clState->kernel,2,sizeof(cl_mem), &clState->padbuffer8);
clSetKernelArg(clState->kernel,3,sizeof(cl_uint4), &midstate[0]);
clSetKernelArg(clState->kernel,4,sizeof(cl_uint4), &midstate[16]);
#endif
return status;
}
#endif
static void set_threads_hashes(unsigned int vectors, unsigned int *threads,
int64_t *hashes, size_t *globalThreads,
unsigned int minthreads, int intensity)
{
*threads = 1 << (15 + intensity);
if (*threads < minthreads)
*threads = minthreads;
*globalThreads = *threads;
*hashes = *threads * vectors;
}
#endif /* HAVE_OPENCL */
#ifdef HAVE_OPENCL
/* We have only one thread that ever re-initialises GPUs, thus if any GPU
* init command fails due to a completely wedged GPU, the thread will never
* return, unable to harm other GPUs. If it does return, it means we only had
* a soft failure and then the reinit_gpu thread is ready to tackle another
* GPU */
void *reinit_gpu(void *userdata)
{
struct thr_info *mythr = userdata;
struct cgpu_info *cgpu;
struct thr_info *thr;
struct timeval now;
char name[256];
int thr_id;
int gpu;
pthread_detach(pthread_self());
select_cgpu:
cgpu = tq_pop(mythr->q, NULL);
if (!cgpu)
goto out;
if (clDevicesNum() != nDevs) {
applog(LOG_WARNING, "Hardware not reporting same number of active devices, will not attempt to restart GPU");
goto out;
}
gpu = cgpu->device_id;
for (thr_id = 0; thr_id < mining_threads; ++thr_id) {
thr = &thr_info[thr_id];
cgpu = thr->cgpu;
if (cgpu->api != &opencl_api)
continue;
if (dev_from_id(thr_id) != gpu)
continue;
thr = &thr_info[thr_id];
if (!thr) {
applog(LOG_WARNING, "No reference to thread %d exists", thr_id);
continue;
}
thr->rolling = thr->cgpu->rolling = 0;
/* Reports the last time we tried to revive a sick GPU */
gettimeofday(&thr->sick, NULL);
if (!pthread_cancel(thr->pth)) {
applog(LOG_WARNING, "Thread %d still exists, killing it off", thr_id);
} else
applog(LOG_WARNING, "Thread %d no longer exists", thr_id);
}
for (thr_id = 0; thr_id < mining_threads; ++thr_id) {
int virtual_gpu;
thr = &thr_info[thr_id];
cgpu = thr->cgpu;
if (cgpu->api != &opencl_api)
continue;
if (dev_from_id(thr_id) != gpu)
continue;
virtual_gpu = cgpu->virtual_gpu;
/* Lose this ram cause we may get stuck here! */
//tq_freeze(thr->q);
thr->q = tq_new();
if (!thr->q)
quit(1, "Failed to tq_new in reinit_gpu");
/* Lose this ram cause we may dereference in the dying thread! */
//free(clState);
applog(LOG_INFO, "Reinit GPU thread %d", thr_id);
clStates[thr_id] = initCl(virtual_gpu, name, sizeof(name));
if (!clStates[thr_id]) {
applog(LOG_ERR, "Failed to reinit GPU thread %d", thr_id);
goto select_cgpu;
}
applog(LOG_INFO, "initCl() finished. Found %s", name);
if (unlikely(thr_info_create(thr, NULL, miner_thread, thr))) {
applog(LOG_ERR, "thread %d create failed", thr_id);
return NULL;
}
applog(LOG_WARNING, "Thread %d restarted", thr_id);
}
gettimeofday(&now, NULL);
get_datestamp(cgpu->init, &now);
for (thr_id = 0; thr_id < mining_threads; ++thr_id) {
thr = &thr_info[thr_id];
cgpu = thr->cgpu;
if (cgpu->api != &opencl_api)
continue;
if (dev_from_id(thr_id) != gpu)
continue;
tq_push(thr->q, &ping);
}
goto select_cgpu;
out:
return NULL;
}
#else
void *reinit_gpu(__maybe_unused void *userdata)
{
return NULL;
}
#endif
#ifdef HAVE_OPENCL
struct device_api opencl_api;
static void opencl_detect()
{
int i;
nDevs = clDevicesNum();
if (nDevs < 0) {
applog(LOG_ERR, "clDevicesNum returned error, no GPUs usable");
nDevs = 0;
}
if (!nDevs)
return;
for (i = 0; i < nDevs; ++i) {
struct cgpu_info *cgpu;
cgpu = &gpus[i];
cgpu->deven = DEV_ENABLED;
cgpu->api = &opencl_api;
cgpu->device_id = i;
cgpu->threads = opt_g_threads;
cgpu->virtual_gpu = i;
add_cgpu(cgpu);
}
if (!opt_noadl)
init_adl(nDevs);
}
static void reinit_opencl_device(struct cgpu_info *gpu)
{
tq_push(thr_info[gpur_thr_id].q, gpu);
}
#ifdef HAVE_ADL
static void get_opencl_statline_before(char *buf, struct cgpu_info *gpu)
{
if (gpu->has_adl) {
int gpuid = gpu->device_id;
float gt = gpu_temp(gpuid);
int gf = gpu_fanspeed(gpuid);
int gp;
if (gt != -1)
tailsprintf(buf, "%5.1fC ", gt);
else
tailsprintf(buf, " ", gt);
if (gf != -1)
tailsprintf(buf, "%4dRPM ", gf);
else if ((gp = gpu_fanpercent(gpuid)) != -1)
tailsprintf(buf, "%3d%% ", gp);
else
tailsprintf(buf, " ");
tailsprintf(buf, "| ");
}
}
#endif
static void get_opencl_statline(char *buf, struct cgpu_info *gpu)
{
tailsprintf(buf, " I:%2d", gpu->intensity);
}
struct opencl_thread_data {
cl_int (*queue_kernel_parameters)(_clState *, dev_blk_ctx *, cl_uint);
uint32_t *res;
struct work *last_work;
struct work _last_work;
};
static uint32_t *blank_res;
static bool opencl_thread_prepare(struct thr_info *thr)
{
char name[256];
struct timeval now;
struct cgpu_info *cgpu = thr->cgpu;
int gpu = cgpu->device_id;
int virtual_gpu = cgpu->virtual_gpu;
int i = thr->id;
static bool failmessage = false;
if (!blank_res)
blank_res = calloc(BUFFERSIZE, 1);
if (!blank_res) {
applog(LOG_ERR, "Failed to calloc in opencl_thread_init");
return false;
}
strcpy(name, "");
applog(LOG_INFO, "Init GPU thread %i GPU %i virtual GPU %i", i, gpu, virtual_gpu);
clStates[i] = initCl(virtual_gpu, name, sizeof(name));
if (!clStates[i]) {
if (use_curses)
enable_curses();
applog(LOG_ERR, "Failed to init GPU thread %d, disabling device %d", i, gpu);
if (!failmessage) {
char *buf;
applog(LOG_ERR, "Restarting the GPU from the menu will not fix this.");
applog(LOG_ERR, "Try restarting cgminer.");
failmessage = true;
#ifdef HAVE_CURSES
if (use_curses) {
buf = curses_input("Press enter to continue");
if (buf)
free(buf);
}
#endif
}
cgpu->deven = DEV_DISABLED;
cgpu->status = LIFE_NOSTART;
cgpu->device_last_not_well = time(NULL);
cgpu->device_not_well_reason = REASON_DEV_NOSTART;
cgpu->dev_nostart_count++;
return false;
}
if (!cgpu->name)
cgpu->name = strdup(name);
if (!cgpu->kname)
{
switch (clStates[i]->chosen_kernel) {
case KL_DIABLO:
cgpu->kname = "diablo";
break;
case KL_DIAKGCN:
cgpu->kname = "diakgcn";
break;
case KL_PHATK:
cgpu->kname = "phatk";
break;
#ifdef USE_SCRYPT
case KL_SCRYPT:
cgpu->kname = "scrypt";
break;
#endif
case KL_POCLBM:
cgpu->kname = "poclbm";
break;
default:
break;
}
}
applog(LOG_INFO, "initCl() finished. Found %s", name);
gettimeofday(&now, NULL);
get_datestamp(cgpu->init, &now);
have_opencl = true;
return true;
}
static bool opencl_thread_init(struct thr_info *thr)
{
const int thr_id = thr->id;
struct cgpu_info *gpu = thr->cgpu;
struct opencl_thread_data *thrdata;
_clState *clState = clStates[thr_id];
cl_int status;
thrdata = calloc(1, sizeof(*thrdata));
thr->cgpu_data = thrdata;
if (!thrdata) {
applog(LOG_ERR, "Failed to calloc in opencl_thread_init");
return false;
}
switch (clState->chosen_kernel) {
case KL_POCLBM:
thrdata->queue_kernel_parameters = &queue_poclbm_kernel;
break;
case KL_PHATK:
thrdata->queue_kernel_parameters = &queue_phatk_kernel;
break;
case KL_DIAKGCN:
thrdata->queue_kernel_parameters = &queue_diakgcn_kernel;
break;
#ifdef USE_SCRYPT
case KL_SCRYPT:
thrdata->queue_kernel_parameters = &queue_scrypt_kernel;
break;
#endif
default:
case KL_DIABLO:
thrdata->queue_kernel_parameters = &queue_diablo_kernel;
break;
}
thrdata->res = calloc(BUFFERSIZE, 1);
if (!thrdata->res) {
free(thrdata);
applog(LOG_ERR, "Failed to calloc in opencl_thread_init");
return false;
}
#ifdef USE_SCRYPT
if (opt_scrypt)
status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, BUFFERSIZE, blank_res, 0, NULL,NULL);
#endif
status = clEnqueueWriteBuffer(clState->commandQueue, clState->outputBuffer, CL_TRUE, 0,
BUFFERSIZE, blank_res, 0, NULL, NULL);
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error: clEnqueueWriteBuffer failed.");
return false;
}
gpu->status = LIFE_WELL;
gpu->device_last_well = time(NULL);
return true;
}
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];
clFinish(clState->commandQueue);
if (thrdata->res[FOUND]) {
thrdata->last_work = &thrdata->_last_work;
memcpy(thrdata->last_work, work, sizeof(*thrdata->last_work));
}
}
static bool opencl_prepare_work(struct thr_info __maybe_unused *thr, struct work *work)
{
precalc_hash(&work->blk, (uint32_t *)(work->midstate), (uint32_t *)(work->data + 64));
return true;
}
extern int opt_dynamic_interval;
static int64_t opencl_scanhash(struct thr_info *thr, struct work *work,
int64_t __maybe_unused max_nonce)
{
const int thr_id = thr->id;
struct opencl_thread_data *thrdata = thr->cgpu_data;
struct cgpu_info *gpu = thr->cgpu;
_clState *clState = clStates[thr_id];
const cl_kernel *kernel = &clState->kernel;
const int dynamic_us = opt_dynamic_interval * 1000;
cl_int status;
size_t globalThreads[1];
size_t localThreads[1] = { clState->wsize };
unsigned int threads;
int64_t hashes;
/* This finish flushes the readbuffer set with CL_FALSE later */
clFinish(clState->commandQueue);
gettimeofday(&gpu->tv_gpuend, NULL);
if (gpu->dynamic) {
struct timeval diff;
suseconds_t gpu_us;
timersub(&gpu->tv_gpuend, &gpu->tv_gpustart, &diff);
gpu_us = diff.tv_sec * 1000000 + diff.tv_usec;
if (likely(gpu_us >= 0)) {
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;
}
}
}
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 */
if (thrdata->res[FOUND]) {
/* Clear the buffer again */
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.");
return -1;
}
if (unlikely(thrdata->last_work)) {
applog(LOG_DEBUG, "GPU %d found something in last work?", gpu->device_id);
postcalc_hash_async(thr, thrdata->last_work, thrdata->res);
thrdata->last_work = NULL;
} else {
applog(LOG_DEBUG, "GPU %d found something?", gpu->device_id);
postcalc_hash_async(thr, work, thrdata->res);
}
memset(thrdata->res, 0, BUFFERSIZE);
clFinish(clState->commandQueue);
}
gettimeofday(&gpu->tv_gpustart, NULL);
if (clState->goffset) {
size_t global_work_offset[1];
global_work_offset[0] = work->blk.nonce;
status = clEnqueueNDRangeKernel(clState->commandQueue, *kernel, 1, global_work_offset,
globalThreads, localThreads, 0, NULL, NULL);
} else
status = clEnqueueNDRangeKernel(clState->commandQueue, *kernel, 1, NULL,
globalThreads, localThreads, 0, NULL, NULL);
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)", status);
return -1;
}
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;
}
/* 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 */
work->blk.nonce += gpu->max_hashes;
return hashes;
}
static void opencl_thread_shutdown(struct thr_info *thr)
{
const int thr_id = thr->id;
_clState *clState = clStates[thr_id];
clReleaseCommandQueue(clState->commandQueue);
clReleaseKernel(clState->kernel);
clReleaseProgram(clState->program);
clReleaseContext(clState->context);
}
struct device_api opencl_api = {
.dname = "opencl",
.name = "GPU",
.api_detect = opencl_detect,
.reinit_device = reinit_opencl_device,
#ifdef HAVE_ADL
.get_statline_before = get_opencl_statline_before,
#endif
.get_statline = get_opencl_statline,
.thread_prepare = opencl_thread_prepare,
.thread_init = opencl_thread_init,
.free_work = opencl_free_work,
.prepare_work = opencl_prepare_work,
.scanhash = opencl_scanhash,
.thread_shutdown = opencl_thread_shutdown,
};
#endif