worker: ethash cl fixes

parent b57890d7
Pipeline #17974282 passed with stages
in 8 minutes
......@@ -393,10 +393,8 @@ __kernel void ethash_search(
if (as_ulong(as_uchar8(state[0]).s76543210) < target)
{
// 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;
uint slot = min(MAX_OUTPUTS, atomic_inc(&g_output[0]) + 1);
g_output[slot] = gid;
}
}
......
......@@ -37,8 +37,9 @@ func (worker *Cryptonight) Start() error {
workChannels := make([]chan *cryptonight.Work, totalThreads)
for i := 0; i < totalThreads; i++ {
workChannels[i] = make(chan *cryptonight.Work, 1)
index := i
defer close(workChannels[index])
defer func(i int) {
close(workChannels[i])
}(i)
}
if len(worker.clDevices) > 0 {
......
......@@ -51,8 +51,9 @@ func (worker *Ethash) Start() error {
workChannels := make([]chan *ethash.Work, totalThreads)
for i := 0; i < totalThreads; i++ {
workChannels[i] = make(chan *ethash.Work, 1)
index := i
defer close(workChannels[index])
defer func(i int) {
close(workChannels[i])
}(i)
}
var light *ethash.Light
......@@ -170,6 +171,8 @@ func (worker *Ethash) clThread(key []string, cl *ethashCL, workChan chan *ethash
var ok bool
var results [2]uint32
nonce := uint64(worker.rand.Uint32())
for {
select {
case work, ok = <-workChan:
......@@ -179,20 +182,22 @@ func (worker *Ethash) clThread(key []string, cl *ethashCL, workChan chan *ethash
if err := cl.Update(work.Header, work.Target); err != nil {
workerError(err)
}
nonce = uint64(worker.rand.Uint32())
default:
start := time.Now()
startNonce := uint64(worker.rand.Uint32())
if err := cl.Run(work.ExtraNonce+startNonce, results); err != nil {
if err := cl.Run(work.ExtraNonce+nonce, results); err != nil {
workerError(err)
}
if results[0] > 0 {
worker.Shares <- ethash.Share{
JobId: work.JobId,
Nonce: startNonce + uint64(results[1]),
Nonce: nonce + uint64(results[1]),
}
}
worker.metrics.IncrCounter(key, float32(10*1024/time.Since(start).Seconds()))
nonce += uint64(cl.globalWorkSize)
worker.metrics.IncrCounter(key, float32(float64(cl.globalWorkSize)/time.Since(start).Seconds()))
}
}
......
......@@ -68,7 +68,7 @@ func newEthashCL(config CLDeviceConfig, light *ethash.Light) (*ethashCL, error)
return nil, errors.WithStack(err)
}
dag, err := ctx.CreateEmptyBuffer(cl.MemReadOnly, light.DataSize)
dag, err := ctx.CreateEmptyBuffer(cl.MemReadOnly, light.DataSize/4)
if err != nil {
return nil, errors.WithStack(err)
}
......@@ -86,10 +86,10 @@ func newEthashCL(config CLDeviceConfig, light *ethash.Light) (*ethashCL, error)
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", light.DataSize/128),
fmt.Sprintf("-D%v=%v", "DAG_SIZE", light.DataSize/128/4),
fmt.Sprintf("-D%v=%v", "LIGHT_SIZE", len(light.Cache)/64), // TODO what's the right size?
//fmt.Sprintf("-D%v=%v", "ACCESSES", workgroupSize), TODO??
fmt.Sprintf("-D%v=%v", "MAX_OUTPUTS", 1),
fmt.Sprintf("-D%v=%v", "MAX_OUTPUTS", "1u"),
// fmt.Sprintf("-D%v=%v", "PLATFORM", workgroupSize), TODO!!
fmt.Sprintf("-D%v=%v", "COMPUTE", 0), // TODO
fmt.Sprintf("-D%v=%v", "THREADS_PER_HASH", 8),
......@@ -127,7 +127,7 @@ func newEthashCL(config CLDeviceConfig, light *ethash.Light) (*ethashCL, error)
return nil, errors.WithStack(err)
}
work := light.DataSize / 128
work := light.DataSize / 128 / 4
fullRuns := work / globalWorkSize
restWork := work % globalWorkSize
if restWork > 0 {
......@@ -176,19 +176,19 @@ func (cl *ethashCL) Update(header []byte, target *big.Int) error {
targetBytes := target.Bytes()
if _, err := cl.queue.EnqueueWriteBuffer(cl.header, false, 0, len(header), unsafe.Pointer(&header[0]), nil); err != nil {
return err
return errors.WithStack(err)
}
if _, err := cl.queue.EnqueueWriteBuffer(cl.search, false, 0, 4, unsafe.Pointer(&zero), nil); err != nil {
return err
return errors.WithStack(err)
}
if err := cl.searchKernel.SetArgBuffer(0, cl.search); err != nil {
return err
return errors.WithStack(err)
}
if err := cl.searchKernel.SetArgUnsafe(0, 64, unsafe.Pointer(&targetBytes[0])); err != nil {
return err
if err := cl.searchKernel.SetArgUnsafe(4, 8, unsafe.Pointer(&targetBytes[0])); err != nil {
return errors.WithStack(err)
}
return nil
......@@ -197,19 +197,19 @@ func (cl *ethashCL) Update(header []byte, target *big.Int) error {
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
return errors.WithStack(err)
}
if err := cl.searchKernel.SetArgUint64(3, nonce); err != nil {
return err
return errors.WithStack(err)
}
if _, err := cl.queue.EnqueueNDRangeKernel(cl.searchKernel, nil, []int{cl.globalWorkSize}, []int{cl.workgroupSize}, nil); err != nil {
return err
return errors.WithStack(err)
}
if err := cl.queue.Finish(); err != nil {
return err
return errors.WithStack(err)
}
return nil
......
......@@ -43,6 +43,6 @@ type CLDeviceConfig struct {
// workerError reports an error and sleeps
func workerError(err error) {
raven.CaptureError(err, nil)
log.Error(err)
log.Errorf("%+v", err)
time.Sleep(time.Minute * 1)
}
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