path: root/Godeps/_workspace/src/github.com/ethereum/ethash/ethash_opencl.go

// Copyright 2014 The go-ethereum Authors
// This file is part of the go-ethereum library.
//
// The go-ethereum library is free software: you can redistribute it and/or modify
// it under the terms of the GNU Lesser General Public License as published by
// the Free Software Foundation, either version 3 of the License, or
// (at your option) any later version.
//
// The go-ethereum library is distributed in the hope that it will be useful,
// but WITHOUT ANY WARRANTY; without even the implied warranty of
// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
// GNU Lesser General Public License for more details.
//
// You should have received a copy of the GNU Lesser General Public License
// along with the go-ethereum library. If not, see <http://www.gnu.org/licenses/>.

// +build opencl

package ethash

//#cgo LDFLAGS: -w
//#include <stdint.h>
//#include <string.h>
//#include "src/libethash/internal.h"
import "C"

import (
    crand "crypto/rand"
    "encoding/binary"
    "fmt"
    "math"
    "math/big"
    mrand "math/rand"
    "strconv"
    "strings"
    "sync"
    "sync/atomic"
    "time"
    "unsafe"

    "github.com/Gustav-Simonsson/go-opencl/cl"
    "github.com/ethereum/go-ethereum/common"
    "github.com/ethereum/go-ethereum/pow"
)

/*

  This code have two main entry points:

  1. The initCL(...)  function configures one or more OpenCL device
     (for now only GPU) and loads the Ethash DAG onto device memory

  2. The Search(...) function loads a Ethash nonce into device(s) memory and
     executes the Ethash OpenCL kernel.

  Throughout the code, we refer to "host memory" and "device memory".
  For most systems (e.g. regular PC GPU miner) the host memory is RAM and
  device memory is the GPU global memory (e.g. GDDR5).

  References mentioned in code comments:

  1. https://github.com/ethereum/wiki/wiki/Ethash
  2. https://github.com/ethereum/cpp-ethereum/blob/develop/libethash-cl/ethash_cl_miner.cpp
  3. https://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/
  4. http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2013/12/AMD_OpenCL_Programming_User_Guide.pdf

*/

type OpenCLDevice struct {
    deviceId int
    device   *cl.Device
    openCL11 bool // OpenCL version 1.1 and 1.2 are handled a bit different
    openCL12 bool

    dagBuf        *cl.MemObject // Ethash full DAG in device mem
    headerBuf     *cl.MemObject // Hash of block-to-mine in device mem
    searchBuffers []*cl.MemObject

    searchKernel *cl.Kernel
    hashKernel   *cl.Kernel

    queue         *cl.CommandQueue
    ctx           *cl.Context
    workGroupSize int

    nonceRand *mrand.Rand // seeded by crypto/rand, see comments where it's initialised
    result    common.Hash
}

type OpenCLMiner struct {
    mu sync.Mutex

    ethash *Ethash // Ethash full DAG & cache in host mem

    deviceIds []int
    devices   []*OpenCLDevice

    dagSize uint64

    hashRate int32 // Go atomics & uint64 have some issues; int32 is supported on all platforms
}

type pendingSearch struct {
    bufIndex   uint32
    startNonce uint64
}

const (
    SIZEOF_UINT32 = 4

    // See [1]
    ethashMixBytesLen = 128
    ethashAccesses    = 64

    // See [4]
    workGroupSize    = 32 // must be multiple of 8
    maxSearchResults = 63
    searchBufSize    = 2
    globalWorkSize   = 1024 * 256
)

func NewCL(deviceIds []int) *OpenCLMiner {
    ids := make([]int, len(deviceIds))
    copy(ids, deviceIds)
    return &OpenCLMiner{
        ethash:    New(),
        dagSize:   0, // to see if we need to update DAG.
        deviceIds: ids,
    }
}

func PrintDevices() {
    fmt.Println("=============================================")
    fmt.Println("============ OpenCL Device Info =============")
    fmt.Println("=============================================")

    var found []*cl.Device

    platforms, err := cl.GetPlatforms()
    if err != nil {
        fmt.Println("Plaform error (check your OpenCL installation): %v", err)
        return
    }

    for i, p := range platforms {
        fmt.Println("Platform id             ", i)
        fmt.Println("Platform Name           ", p.Name())
        fmt.Println("Platform Vendor         ", p.Vendor())
        fmt.Println("Platform Version        ", p.Version())
        fmt.Println("Platform Extensions     ", p.Extensions())
        fmt.Println("Platform Profile        ", p.Profile())
        fmt.Println("")

        devices, err := cl.GetDevices(p, cl.DeviceTypeGPU)
        if err != nil {
            fmt.Println("Device error (check your GPU drivers) :", err)
            return
        }

        for _, d := range devices {
            fmt.Println("Device OpenCL id        ", i)
            fmt.Println("Device id for mining    ", len(found))
            fmt.Println("Device Name             ", d.Name())
            fmt.Println("Vendor                  ", d.Vendor())
            fmt.Println("Version                 ", d.Version())
            fmt.Println("Driver version          ", d.DriverVersion())
            fmt.Println("Address bits            ", d.AddressBits())
            fmt.Println("Max clock freq          ", d.MaxClockFrequency())
            fmt.Println("Global mem size         ", d.GlobalMemSize())
            fmt.Println("Max constant buffer size", d.MaxConstantBufferSize())
            fmt.Println("Max mem alloc size      ", d.MaxMemAllocSize())
            fmt.Println("Max compute units       ", d.MaxComputeUnits())
            fmt.Println("Max work group size     ", d.MaxWorkGroupSize())
            fmt.Println("Max work item sizes     ", d.MaxWorkItemSizes())
            fmt.Println("=============================================")

            found = append(found, d)
        }
    }
    if len(found) == 0 {
        fmt.Println("Found no GPU(s). Check that your OS can see the GPU(s)")
    } else {
        var idsFormat string
        for i := 0; i < len(found); i++ {
            idsFormat += strconv.Itoa(i)
            if i != len(found)-1 {
                idsFormat += ","
            }
        }
        fmt.Printf("Found %v devices. Benchmark first GPU:       geth gpubench 0\n", len(found))
        fmt.Printf("Mine using all GPUs:                        geth --minegpu %v\n", idsFormat)
    }
}

// See [2]. We basically do the same here, but the Go OpenCL bindings
// are at a slightly higher abtraction level.
func InitCL(blockNum uint64, c *OpenCLMiner) error {
    platforms, err := cl.GetPlatforms()
    if err != nil {
        return fmt.Errorf("Plaform error: %v\nCheck your OpenCL installation and then run geth gpuinfo", err)
    }

    var devices []*cl.Device
    for _, p := range platforms {
        ds, err := cl.GetDevices(p, cl.DeviceTypeGPU)
        if err != nil {
            return fmt.Errorf("Devices error: %v\nCheck your GPU drivers and then run geth gpuinfo", err)
        }
        for _, d := range ds {
            devices = append(devices, d)
        }
    }

    pow := New()
    _ = pow.getDAG(blockNum)     // generates DAG if we don't have it
    pow.Light.getCache(blockNum) // and cache

    c.ethash = pow
    dagSize := uint64(C.ethash_get_datasize(C.uint64_t(blockNum)))
    c.dagSize = dagSize

    for _, id := range c.deviceIds {
        if id > len(devices)-1 {
            return fmt.Errorf("Device id not found. See available device ids with: geth gpuinfo")
        } else {
            err := initCLDevice(id, devices[id], c)
            if err != nil {
                return err
            }
        }
    }
    if len(c.devices) == 0 {
        return fmt.Errorf("No GPU devices found")
    }
    return nil
}

func initCLDevice(deviceId int, device *cl.Device, c *OpenCLMiner) error {
    devMaxAlloc := uint64(device.MaxMemAllocSize())
    devGlobalMem := uint64(device.GlobalMemSize())

    // TODO: more fine grained version logic
    if device.Version() == "OpenCL 1.0" {
        fmt.Println("Device OpenCL version not supported: ", device.Version())
        return fmt.Errorf("opencl version not supported")
    }

    var cl11, cl12 bool
    if device.Version() == "OpenCL 1.1" {
        cl11 = true
    }
    if device.Version() == "OpenCL 1.2" {
        cl12 = true
    }

    // log warnings but carry on; some device drivers report inaccurate values
    if c.dagSize > devGlobalMem {
        fmt.Printf("WARNING: device memory may be insufficient: %v. DAG size: %v.\n", devGlobalMem, c.dagSize)
    }

    if c.dagSize > devMaxAlloc {
        fmt.Printf("WARNING: DAG size (%v) larger than device max memory allocation size (%v).\n", c.dagSize, devMaxAlloc)
        fmt.Printf("You probably have to export GPU_MAX_ALLOC_PERCENT=95\n")
    }

    fmt.Printf("Initialising device %v: %v\n", deviceId, device.Name())

    context, err := cl.CreateContext([]*cl.Device{device})
    if err != nil {
        return fmt.Errorf("failed creating context:", err)
    }

    // TODO: test running with CL_QUEUE_PROFILING_ENABLE for profiling?
    queue, err := context.CreateCommandQueue(device, 0)
    if err != nil {
        return fmt.Errorf("command queue err:", err)
    }

    // See [4] section 3.2 and [3] "clBuildProgram".
    // The OpenCL kernel code is compiled at run-time.
    kvs := make(map[string]string, 4)
    kvs["GROUP_SIZE"] = strconv.FormatUint(workGroupSize, 10)
    kvs["DAG_SIZE"] = strconv.FormatUint(c.dagSize/ethashMixBytesLen, 10)
    kvs["ACCESSES"] = strconv.FormatUint(ethashAccesses, 10)
    kvs["MAX_OUTPUTS"] = strconv.FormatUint(maxSearchResults, 10)
    kernelCode := replaceWords(kernel, kvs)

    program, err := context.CreateProgramWithSource([]string{kernelCode})
    if err != nil {
        return fmt.Errorf("program err:", err)
    }

    /* if using AMD OpenCL impl, you can set this to debug on x86 CPU device.
       see AMD OpenCL programming guide section 4.2

       export in shell before running:
       export AMD_OCL_BUILD_OPTIONS_APPEND="-g -O0"
       export CPU_MAX_COMPUTE_UNITS=1

    buildOpts := "-g -cl-opt-disable"

    */
    buildOpts := ""
    err = program.BuildProgram([]*cl.Device{device}, buildOpts)
    if err != nil {
        return fmt.Errorf("program build err:", err)
    }

    var searchKernelName, hashKernelName string
    searchKernelName = "ethash_search"
    hashKernelName = "ethash_hash"

    searchKernel, err := program.CreateKernel(searchKernelName)
    hashKernel, err := program.CreateKernel(hashKernelName)
    if err != nil {
        return fmt.Errorf("kernel err:", err)
    }

    // TODO: when this DAG size appears, patch the Go bindings
    // (context.go) to work with uint64 as size_t
    if c.dagSize > math.MaxInt32 {
        fmt.Println("DAG too large for allocation.")
        return fmt.Errorf("DAG too large for alloc")
    }

    // TODO: patch up Go bindings to work with size_t, will overflow if > maxint32
    // TODO: fuck. shit's gonna overflow around 2017-06-09 12:17:02
    dagBuf := *(new(*cl.MemObject))
    dagBuf, err = context.CreateEmptyBuffer(cl.MemReadOnly, int(c.dagSize))
    if err != nil {
        return fmt.Errorf("allocating dag buf failed: ", err)
    }

    // write DAG to device mem
    dagPtr := unsafe.Pointer(c.ethash.Full.current.ptr.data)
    _, err = queue.EnqueueWriteBuffer(dagBuf, true, 0, int(c.dagSize), dagPtr, nil)
    if err != nil {
        return fmt.Errorf("writing to dag buf failed: ", err)
    }

    searchBuffers := make([]*cl.MemObject, searchBufSize)
    for i := 0; i < searchBufSize; i++ {
        searchBuff, err := context.CreateEmptyBuffer(cl.MemWriteOnly, (1+maxSearchResults)*SIZEOF_UINT32)
        if err != nil {
            return fmt.Errorf("search buffer err:", err)
        }
        searchBuffers[i] = searchBuff
    }

    headerBuf, err := context.CreateEmptyBuffer(cl.MemReadOnly, 32)
    if err != nil {
        return fmt.Errorf("header buffer err:", err)
    }

    // Unique, random nonces are crucial for mining efficieny.
    // While we do not need cryptographically secure PRNG for nonces,
    // we want to have uniform distribution and minimal repetition of nonces.
    // We could guarantee strict uniqueness of nonces by generating unique ranges,
    // but a int64 seed from crypto/rand should be good enough.
    // we then use math/rand for speed and to avoid draining OS entropy pool
    seed, err := crand.Int(crand.Reader, big.NewInt(math.MaxInt64))
    if err != nil {
        return err
    }
    nonceRand := mrand.New(mrand.NewSource(seed.Int64()))

    deviceStruct := &OpenCLDevice{
        deviceId: deviceId,
        device:   device,
        openCL11: cl11,
        openCL12: cl12,

        dagBuf:        dagBuf,
        headerBuf:     headerBuf,
        searchBuffers: searchBuffers,

        searchKernel: searchKernel,
        hashKernel:   hashKernel,

        queue: queue,
        ctx:   context,

        workGroupSize: workGroupSize,

        nonceRand: nonceRand,
    }
    c.devices = append(c.devices, deviceStruct)

    return nil
}

func (c *OpenCLMiner) Search(block pow.Block, stop <-chan struct{}, index int) (uint64, []byte) {
    c.mu.Lock()
    newDagSize := uint64(C.ethash_get_datasize(C.uint64_t(block.NumberU64())))
    if newDagSize > c.dagSize {
        // TODO: clean up buffers from previous DAG?
        err := InitCL(block.NumberU64(), c)
        if err != nil {
            fmt.Println("OpenCL init error: ", err)
            return 0, []byte{0}
        }
    }
    defer c.mu.Unlock()

    // Avoid unneeded OpenCL initialisation if we received stop while running InitCL
    select {
    case <-stop:
        return 0, []byte{0}
    default:
    }

    headerHash := block.HashNoNonce()
    diff := block.Difficulty()
    target256 := new(big.Int).Div(maxUint256, diff)
    target64 := new(big.Int).Rsh(target256, 192).Uint64()
    var zero uint32 = 0

    d := c.devices[index]

    _, err := d.queue.EnqueueWriteBuffer(d.headerBuf, false, 0, 32, unsafe.Pointer(&headerHash[0]), nil)
    if err != nil {
        fmt.Println("Error in Search clEnqueueWriterBuffer : ", err)
        return 0, []byte{0}
    }

    for i := 0; i < searchBufSize; i++ {
        _, err := d.queue.EnqueueWriteBuffer(d.searchBuffers[i], false, 0, 4, unsafe.Pointer(&zero), nil)
        if err != nil {
            fmt.Println("Error in Search clEnqueueWriterBuffer : ", err)
            return 0, []byte{0}
        }
    }

    // wait for all search buffers to complete
    err = d.queue.Finish()
    if err != nil {
        fmt.Println("Error in Search clFinish : ", err)
        return 0, []byte{0}
    }

    err = d.searchKernel.SetArg(1, d.headerBuf)
    if err != nil {
        fmt.Println("Error in Search clSetKernelArg : ", err)
        return 0, []byte{0}
    }

    err = d.searchKernel.SetArg(2, d.dagBuf)
    if err != nil {
        fmt.Println("Error in Search clSetKernelArg : ", err)
        return 0, []byte{0}
    }

    err = d.searchKernel.SetArg(4, target64)
    if err != nil {
        fmt.Println("Error in Search clSetKernelArg : ", err)
        return 0, []byte{0}
    }
    err = d.searchKernel.SetArg(5, uint32(math.MaxUint32))
    if err != nil {
        fmt.Println("Error in Search clSetKernelArg : ", err)
        return 0, []byte{0}
    }

    // wait on this before returning
    var preReturnEvent *cl.Event
    if d.openCL12 {
        preReturnEvent, err = d.ctx.CreateUserEvent()
        if err != nil {
            fmt.Println("Error in Search create CL user event : ", err)
            return 0, []byte{0}
        }
    }

    pending := make([]pendingSearch, 0, searchBufSize)
    var p *pendingSearch
    searchBufIndex := uint32(0)
    var checkNonce uint64
    loops := int64(0)
    prevHashRate := int32(0)
    start := time.Now().UnixNano()
    // we grab a single random nonce and sets this as argument to the kernel search function
    // the device will then add each local threads gid to the nonce, creating a unique nonce
    // for each device computing unit executing in parallel
    initNonce := uint64(d.nonceRand.Int63())
    for nonce := initNonce; ; nonce += uint64(globalWorkSize) {
        select {
        case <-stop:

            /*
                if d.openCL12 {
                    err = cl.WaitForEvents([]*cl.Event{preReturnEvent})
                    if err != nil {
                        fmt.Println("Error in Search WaitForEvents: ", err)
                    }
                }
            */

            atomic.AddInt32(&c.hashRate, -prevHashRate)
            return 0, []byte{0}
        default:
        }

        if (loops % (1 << 7)) == 0 {
            elapsed := time.Now().UnixNano() - start
            // TODO: verify if this is correct hash rate calculation
            hashes := (float64(1e9) / float64(elapsed)) * float64(loops*1024*256)
            hashrateDiff := int32(hashes) - prevHashRate
            prevHashRate = int32(hashes)
            atomic.AddInt32(&c.hashRate, hashrateDiff)
        }
        loops++

        err = d.searchKernel.SetArg(0, d.searchBuffers[searchBufIndex])
        if err != nil {
            fmt.Println("Error in Search clSetKernelArg : ", err)
            return 0, []byte{0}
        }
        err = d.searchKernel.SetArg(3, nonce)
        if err != nil {
            fmt.Println("Error in Search clSetKernelArg : ", err)
            return 0, []byte{0}
        }

        // execute kernel
        _, err := d.queue.EnqueueNDRangeKernel(
            d.searchKernel,
            []int{0},
            []int{globalWorkSize},
            []int{d.workGroupSize},
            nil)
        if err != nil {
            fmt.Println("Error in Search clEnqueueNDRangeKernel : ", err)
            return 0, []byte{0}
        }

        pending = append(pending, pendingSearch{bufIndex: searchBufIndex, startNonce: nonce})
        searchBufIndex = (searchBufIndex + 1) % searchBufSize

        if len(pending) == searchBufSize {
            p = &(pending[searchBufIndex])
            cres, _, err := d.queue.EnqueueMapBuffer(d.searchBuffers[p.bufIndex], true,
                cl.MapFlagRead, 0, (1+maxSearchResults)*SIZEOF_UINT32,
                nil)
            if err != nil {
                fmt.Println("Error in Search clEnqueueMapBuffer: ", err)
                return 0, []byte{0}
            }

            results := cres.ByteSlice()
            nfound := binary.LittleEndian.Uint32(results)
            nfound = uint32(math.Min(float64(nfound), float64(maxSearchResults)))
            // OpenCL returns the offsets from the start nonce
            for i := uint32(0); i < nfound; i++ {
                lo := (i + 1) * SIZEOF_UINT32
                hi := (i + 2) * SIZEOF_UINT32
                upperNonce := uint64(binary.LittleEndian.Uint32(results[lo:hi]))
                checkNonce = p.startNonce + upperNonce
                if checkNonce != 0 {
                    cn := C.uint64_t(checkNonce)
                    ds := C.uint64_t(c.dagSize)
                    // We verify that the nonce is indeed a solution by
                    // executing the Ethash verification function (on the CPU).
                    ret := C.ethash_light_compute_internal(c.ethash.Light.current.ptr, ds, hashToH256(headerHash), cn)
                    // TODO: return result first
                    if ret.success && h256ToHash(ret.result).Big().Cmp(target256) <= 0 {
                        _, err = d.queue.EnqueueUnmapMemObject(d.searchBuffers[p.bufIndex], cres, nil)
                        if err != nil {
                            fmt.Println("Error in Search clEnqueueUnmapMemObject: ", err)
                        }
                        if d.openCL12 {
                            err = cl.WaitForEvents([]*cl.Event{preReturnEvent})
                            if err != nil {
                                fmt.Println("Error in Search WaitForEvents: ", err)
                            }
                        }
                        return checkNonce, C.GoBytes(unsafe.Pointer(&ret.mix_hash), C.int(32))
                    }

                    _, err := d.queue.EnqueueWriteBuffer(d.searchBuffers[p.bufIndex], false, 0, 4, unsafe.Pointer(&zero), nil)
                    if err != nil {
                        fmt.Println("Error in Search cl: EnqueueWriteBuffer", err)
                        return 0, []byte{0}
                    }
                }
            }
            _, err = d.queue.EnqueueUnmapMemObject(d.searchBuffers[p.bufIndex], cres, nil)
            if err != nil {
                fmt.Println("Error in Search clEnqueueUnMapMemObject: ", err)
                return 0, []byte{0}
            }
            pending = append(pending[:searchBufIndex], pending[searchBufIndex+1:]...)
        }
    }
    if d.openCL12 {
        err := cl.WaitForEvents([]*cl.Event{preReturnEvent})
        if err != nil {
            fmt.Println("Error in Search clWaitForEvents: ", err)
            return 0, []byte{0}
        }
    }
    return 0, []byte{0}
}

func (c *OpenCLMiner) Verify(block pow.Block) bool {
    return c.ethash.Light.Verify(block)
}
func (c *OpenCLMiner) GetHashrate() int64 {
    return int64(atomic.LoadInt32(&c.hashRate))
}
func (c *OpenCLMiner) Turbo(on bool) {
    // This is GPU mining. Always be turbo.
}

func replaceWords(text string, kvs map[string]string) string {
    for k, v := range kvs {
        text = strings.Replace(text, k, v, -1)
    }
    return text
}

func logErr(err error) {
    if err != nil {
        fmt.Println("Error in OpenCL call:", err)
    }
}

func argErr(err error) error {
    return fmt.Errorf("arg err: %v", err)
}