예제 #1
0
// Verify checks whether the block's nonce is valid.
func (l *Light) Verify(block pow.Block) bool {
	// TODO: do ethash_quick_verify before getCache in order
	// to prevent DOS attacks.
	var (
		blockNum   = block.NumberU64()
		difficulty = block.Difficulty()
		cache      = l.getCache(blockNum)
		dagSize    = C.ethash_get_datasize(C.uint64_t(blockNum))
	)
	if l.test {
		dagSize = dagSizeForTesting
	}
	if blockNum >= epochLength*2048 {
		glog.V(logger.Debug).Infof("block number %d too high, limit is %d", epochLength*2048)
		return false
	}
	// Recompute the hash using the cache.
	hash := hashToH256(block.HashNoNonce())
	ret := C.ethash_light_compute_internal(cache.ptr, dagSize, hash, C.uint64_t(block.Nonce()))
	if !ret.success {
		return false
	}
	// Make sure cache is live until after the C call.
	// This is important because a GC might happen and execute
	// the finalizer before the call completes.
	_ = cache
	// The actual check.
	target := new(big.Int).Div(minDifficulty, difficulty)
	return h256ToHash(ret.result).Big().Cmp(target) <= 0
}
예제 #2
0
// Verify checks whether the block's nonce is valid.
func (l *Light) Verify(block pow.Block) bool {
	// TODO: do ethash_quick_verify before getCache in order
	// to prevent DOS attacks.
	blockNum := block.NumberU64()
	if blockNum >= epochLength*2048 {
		glog.V(logger.Debug).Infof("block number %d too high, limit is %d", epochLength*2048)
		return false
	}

	difficulty := block.Difficulty()
	/* Cannot happen if block header diff is validated prior to PoW, but can
		 happen if PoW is checked first due to parallel PoW checking.
		 We could check the minimum valid difficulty but for SoC we avoid (duplicating)
	   Ethereum protocol consensus rules here which are not in scope of Ethash
	*/
	if difficulty.Cmp(common.Big0) == 0 {
		glog.V(logger.Debug).Infof("invalid block difficulty")
		return false
	}

	cache := l.getCache(blockNum)
	dagSize := C.ethash_get_datasize(C.uint64_t(blockNum))

	if l.test {
		dagSize = dagSizeForTesting
	}
	// Recompute the hash using the cache.
	hash := hashToH256(block.HashNoNonce())
	ret := C.ethash_light_compute_internal(cache.ptr, dagSize, hash, C.uint64_t(block.Nonce()))
	if !ret.success {
		return false
	}

	// avoid mixdigest malleability as it's not included in a block's "hashNononce"
	if block.MixDigest() != h256ToHash(ret.mix_hash) {
		return false
	}

	// Make sure cache is live until after the C call.
	// This is important because a GC might happen and execute
	// the finalizer before the call completes.
	_ = cache
	// The actual check.
	target := new(big.Int).Div(minDifficulty, difficulty)
	return h256ToHash(ret.result).Big().Cmp(target) <= 0
}
예제 #3
0
// 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
}
예제 #4
0
// Verify checks whether the block's nonce is valid.
func (l *Light) Verify(block pow.Block) bool {
	// TODO: do ethash_quick_verify before getCache in order
	// to prevent DOS attacks.
	blockNum := block.NumberU64()
	if blockNum >= epochLength*2048 {
		glog.V(logger.Debug).Infof("block number %d too high, limit is %d", epochLength*2048)
		return false
	}

	difficulty := block.Difficulty()
	/* Cannot happen if block header diff is validated prior to PoW, but can
		 happen if PoW is checked first due to parallel PoW checking.
		 We could check the minimum valid difficulty but for SoC we avoid (duplicating)
	   Ethereum protocol consensus rules here which are not in scope of Ethash
	*/
	if difficulty.Cmp(common.Big0) == 0 {
		glog.V(logger.Debug).Infof("invalid block difficulty")
		return false
	}

	cache := l.getCache(blockNum)
	dagSize := C.ethash_get_datasize(C.uint64_t(blockNum))
	if l.test {
		dagSize = dagSizeForTesting
	}
	// Recompute the hash using the cache.
	ok, mixDigest, result := cache.compute(uint64(dagSize), block.HashNoNonce(), block.Nonce())
	if !ok {
		return false
	}

	// avoid mixdigest malleability as it's not included in a block's "hashNononce"
	if block.MixDigest() != mixDigest {
		return false
	}

	// The actual check.
	target := new(big.Int).Div(maxUint256, difficulty)
	return result.Big().Cmp(target) <= 0
}
예제 #5
0
// generate creates the actual DAG. it can be called from multiple
// goroutines. the first call will generate the DAG, subsequent
// calls wait until it is generated.
func (d *dag) generate() {
	d.gen.Do(func() {
		var (
			started   = time.Now()
			seedHash  = makeSeedHash(d.epoch)
			blockNum  = C.uint64_t(d.epoch * epochLength)
			cacheSize = C.ethash_get_cachesize(blockNum)
			dagSize   = C.ethash_get_datasize(blockNum)
		)
		if d.test {
			cacheSize = cacheSizeForTesting
			dagSize = dagSizeForTesting
		}
		if d.dir == "" {
			d.dir = DefaultDir
		}
		glog.V(logger.Info).Infof("Generating DAG for epoch %d (%x)", d.epoch, seedHash)
		// Generate a temporary cache.
		// TODO: this could share the cache with Light
		cache := C.ethash_light_new_internal(cacheSize, (*C.ethash_h256_t)(unsafe.Pointer(&seedHash[0])))
		defer C.ethash_light_delete(cache)
		// Generate the actual DAG.
		d.ptr = C.ethash_full_new_internal(
			C.CString(d.dir),
			hashToH256(seedHash),
			dagSize,
			cache,
			(C.ethash_callback_t)(unsafe.Pointer(C.ethashGoCallback_cgo)),
		)
		if d.ptr == nil {
			panic("ethash_full_new IO or memory error")
		}
		runtime.SetFinalizer(d, freeDAG)
		glog.V(logger.Info).Infof("Done generating DAG for epoch %d, it took %v", d.epoch, time.Since(started))
	})
}
예제 #6
0
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}
}