Example #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
}
Example #2
0
func (cache *cache) compute(dagSize uint64, hash common.Hash, nonce uint64) (ok bool, mixDigest, result common.Hash) {
	ret := C.ethash_light_compute_internal(cache.ptr, C.uint64_t(dagSize), hashToH256(hash), C.uint64_t(nonce))
	// 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
	return bool(ret.success), h256ToHash(ret.mix_hash), h256ToHash(ret.result)
}
Example #3
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
}
Example #4
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}
}