ethash: implement work loop

parent 0cbe8b14
Pipeline #17669470 passed with stages
in 10 minutes and 5 seconds
......@@ -38,15 +38,12 @@ func NewEthash(seedhash []byte) (*Ethash, error) {
return nil, fmt.Errorf("failed to determine ethash dag storage directory")
}
cache := (*(*[]byte)(unsafe.Pointer(light.cache)))[:light.cache_size]
cache := C.GoBytes(unsafe.Pointer(light.cache), C.int(light.cache_size))
fullsize := C.ethash_get_datasize(light.block_number)
full := C.ethash_full_new_internal((*C.char)(unsafe.Pointer(&dir[0])), sh, fullsize, light, nil)
// FIXME full.file_size undefined (type C.ethash_full_t has no field or method file_size)
// why do we need a hacky new struct type when it works just fine for light above ???
fullInternal := (*(*C.struct_ethash_full_internal)(unsafe.Pointer(full)))
dag := (*(*[]byte)(unsafe.Pointer(fullInternal.data)))[:fullInternal.file_size]
dag := C.GoBytes(unsafe.Pointer(C.ethash_full_dag(full)), C.int(C.ethash_full_dag_size(full)/4))
return &Ethash{
cache,
......
......@@ -120,11 +120,6 @@ struct ethash_full {
uint64_t file_size;
node* data;
};
struct ethash_full_internal {
FILE* file;
uint64_t file_size;
node* data;
};
/**
* Allocate and initialize a new ethash_full handler. Internal version.
......
......@@ -393,8 +393,10 @@ __kernel void ethash_search(
if (as_ulong(as_uchar8(state[0]).s76543210) < target)
{
uint slot = min(MAX_OUTPUTS, atomic_inc(&g_output[0]) + 1);
g_output[slot] = gid;
// TODO "error: call to 'min' is ambiguous" with beignet
// uint slot = min(MAX_OUTPUTS, atomic_inc(&g_output[0]) + 1);
// g_output[slot] = gid;
g_output[1] = gid;
}
}
......
......@@ -80,8 +80,7 @@ func (worker *Cryptonight) gpuThread(key []string, cl *cryptonightCL, workChan c
start := time.Now()
results := make([]uint32, 0x100)
err := cl.RunJob(results, work.NextNonce(cl.Intensity))
if err != nil {
if err := cl.RunJob(results, work.NextNonce(cl.Intensity)); err != nil {
log.Errorw("cl error", "error", err)
return
}
......
......@@ -48,7 +48,10 @@ func (worker *Ethash) Configure(config Config) error {
}
func (worker *Ethash) Start() error {
totalThreads := 1
totalThreads := len(worker.config.CLDevices)
for _, c := range worker.config.Processors {
totalThreads += c.Threads
}
workChannels := make([]chan *ethash.Work, totalThreads)
for i := 0; i < totalThreads; i++ {
......@@ -76,6 +79,18 @@ func (worker *Ethash) Start() error {
return err
}
log.Info("DAG initialized")
if len(worker.config.CLDevices) > 0 {
for i, d := range worker.config.CLDevices {
cl, err := newEthashCL(d, worker.hash)
if err != nil {
return err
}
key := []string{"opencl", fmt.Sprintf("%v", d.Device.Platform.Index), fmt.Sprintf("%v", d.Device.Index)}
go worker.clThread(key, cl, workChannels[i])
}
}
}
for _, ch := range workChannels {
......@@ -109,10 +124,45 @@ func (worker *Ethash) thread(key []string, workChan chan *ethash.Work) {
}
}
func (worker *Ethash) clThread(key []string, cl *ethashCL, workChan chan *ethash.Work) {
defer cl.Release()
work := <-workChan
cl.Update(work.Header, work.Target)
var ok bool
var results [2]uint32
for {
select {
case work, ok = <-workChan:
if !ok {
return
}
cl.Update(work.Header, work.Target)
default:
start := time.Now()
worker.lock.RLock()
startNonce := uint64(worker.rand.Uint32())
cl.Run(work.ExtraNonce+startNonce, results)
if results[0] > 0 {
worker.Shares <- ethash.Share{
JobId: work.JobId,
Nonce: startNonce + uint64(results[1]),
}
}
worker.lock.RUnlock()
worker.metrics.IncrCounter(key, float32(10*1024/time.Since(start).Seconds()))
}
}
}
func (w *Ethash) Capabilities() Capabilities {
return Capabilities{
CPU: true,
OpenCL: false,
OpenCL: true,
CUDA: false,
}
}
......@@ -3,7 +3,9 @@ package worker
import (
"fmt"
"math"
"math/big"
"strings"
"unsafe"
"github.com/jgillich/go-opencl/cl"
"github.com/pkg/errors"
......@@ -13,15 +15,18 @@ import (
)
type ethashCL struct {
ctx *cl.Context
queue *cl.CommandQueue
program *cl.Program
cache *cl.MemObject
dag *cl.MemObject
header *cl.MemObject
search *cl.MemObject
searchKernel *cl.Kernel
dagKernel *cl.Kernel
ctx *cl.Context
queue *cl.CommandQueue
program *cl.Program
cache *cl.MemObject
dag *cl.MemObject
header *cl.MemObject
search *cl.MemObject
searchKernel *cl.Kernel
dagKernel *cl.Kernel
localWorkSize int
workgroupSize int
globalWorkSize int
}
func newEthashCL(config CLDeviceConfig, ethash *ethash.Ethash) (*ethashCL, error) {
......@@ -54,12 +59,32 @@ func newEthashCL(config CLDeviceConfig, ethash *ethash.Ethash) (*ethashCL, error
return nil, errors.WithStack(err)
}
// TODO CreateBuffer results in Invalid Host Ptr, might be a bug in the bindings
cache, err := ctx.CreateEmptyBuffer(cl.MemReadOnly, len(ethash.Cache))
if err != nil {
return nil, errors.WithStack(err)
}
if _, err := queue.EnqueueWriteBuffer(cache, true, 0, len(ethash.Cache), unsafe.Pointer(&ethash.Cache[0]), nil); err != nil {
return nil, errors.WithStack(err)
}
dag, err := ctx.CreateEmptyBuffer(cl.MemReadOnly, len(ethash.DAG))
if err != nil {
return nil, errors.WithStack(err)
}
header, err := ctx.CreateEmptyBuffer(cl.MemReadOnly, 32)
if err != nil {
return nil, errors.WithStack(err)
}
program, err := ctx.CreateProgramWithSource([]string{kernel})
if err != nil {
return nil, errors.WithStack(err)
}
options := []string{
fmt.Sprintf("-D%v=%v", "PLATFORM", 0), // TODO 1 for AMD, 2 for NVIDIA
fmt.Sprintf("-D%v=%v", "GROUP_SIZE", workgroupSize),
fmt.Sprintf("-D%v=%v", "DAG_SIZE", len(ethash.DAG)/128),
fmt.Sprintf("-D%v=%v", "LIGHT_SIZE", len(ethash.Cache)/64), // TODO what's the right size?
......@@ -74,16 +99,6 @@ func newEthashCL(config CLDeviceConfig, ethash *ethash.Ethash) (*ethashCL, error
return nil, errors.WithStack(err)
}
cache, err := ctx.CreateBuffer(cl.MemReadOnly, ethash.Cache)
if err != nil {
return nil, errors.WithStack(err)
}
dag, err := ctx.CreateEmptyBuffer(cl.MemReadOnly, len(ethash.DAG))
if err != nil {
return nil, errors.WithStack(err)
}
searchKernel, err := program.CreateKernel("ethash_search")
if err != nil {
return nil, errors.WithStack(err)
......@@ -94,11 +109,6 @@ func newEthashCL(config CLDeviceConfig, ethash *ethash.Ethash) (*ethashCL, error
return nil, errors.WithStack(err)
}
header, err := ctx.CreateEmptyBuffer(cl.MemReadOnly, 32)
if err != nil {
return nil, errors.WithStack(err)
}
if err := searchKernel.SetArgBuffer(1, header); err != nil {
return nil, errors.WithStack(err)
}
......@@ -136,7 +146,7 @@ func newEthashCL(config CLDeviceConfig, ethash *ethash.Ethash) (*ethashCL, error
}
for i := 0; i < fullRuns; i++ {
dagKernel.SetArg(0, i*globalWorkSize)
dagKernel.SetArgUint32(0, uint32(i*globalWorkSize))
if _, err := queue.EnqueueNDRangeKernel(dagKernel, nil, []int{globalWorkSize}, []int{workgroupSize}, nil); err != nil {
return nil, errors.WithStack(err)
}
......@@ -145,5 +155,74 @@ func newEthashCL(config CLDeviceConfig, ethash *ethash.Ethash) (*ethashCL, error
}
}
return &ethashCL{ctx, queue, program, cache, dag, header, search, searchKernel, dagKernel}, nil
return &ethashCL{
ctx: ctx,
queue: queue,
program: program,
cache: cache,
dag: dag,
header: header,
search: search,
searchKernel: searchKernel,
dagKernel: dagKernel,
localWorkSize: localWorkSize,
workgroupSize: workgroupSize,
globalWorkSize: globalWorkSize,
}, nil
}
func (cl *ethashCL) Update(header []byte, target *big.Int) error {
zero := uint32(0)
targetBytes := target.Bytes()
if _, err := cl.queue.EnqueueWriteBuffer(cl.header, false, 0, len(header), unsafe.Pointer(&header[0]), nil); err != nil {
return err
}
if _, err := cl.queue.EnqueueWriteBuffer(cl.search, false, 0, 4, unsafe.Pointer(&zero), nil); err != nil {
return err
}
if err := cl.searchKernel.SetArgBuffer(0, cl.search); err != nil {
return err
}
if err := cl.searchKernel.SetArgUnsafe(0, 64, unsafe.Pointer(&targetBytes[0])); err != nil {
return err
}
return nil
}
func (cl *ethashCL) Run(nonce uint64, results [2]uint32) error {
if _, err := cl.queue.EnqueueReadBuffer(cl.search, true, 0, 4*len(results), unsafe.Pointer(&results[0]), nil); err != nil {
return err
}
if err := cl.searchKernel.SetArgUint64(3, nonce); err != nil {
return err
}
if _, err := cl.queue.EnqueueNDRangeKernel(cl.searchKernel, nil, []int{cl.globalWorkSize}, []int{cl.workgroupSize}, nil); err != nil {
return err
}
if err := cl.queue.Finish(); err != nil {
return err
}
return nil
}
func (cl *ethashCL) Release() {
defer cl.ctx.Release()
defer cl.queue.Release()
defer cl.program.Release()
defer cl.cache.Release()
defer cl.dag.Release()
defer cl.header.Release()
defer cl.search.Release()
defer cl.searchKernel.Release()
defer cl.dagKernel.Release()
}
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment