Updated ethash

This commit is contained in:
obscuren 2015-03-14 16:42:05 +01:00
parent 387f6bba3e
commit 95711bc2dd
30 changed files with 1254 additions and 597 deletions

4
Godeps/Godeps.json generated
View File

@ -22,8 +22,8 @@
}, },
{ {
"ImportPath": "github.com/ethereum/ethash", "ImportPath": "github.com/ethereum/ethash",
"Comment": "v17-64-ga323708", "Comment": "v23-11-g5376ec8",
"Rev": "a323708b8c4d253b8567bf6c72727d1aec302225" "Rev": "5376ec8816d6bf787d4fc91a08b4527bc5e1f469"
}, },
{ {
"ImportPath": "github.com/ethereum/serpent-go", "ImportPath": "github.com/ethereum/serpent-go",

View File

@ -8,3 +8,5 @@ pyethash.egg-info/
*.so *.so
*~ *~
*.swp *.swp
MANIFEST
dist/

View File

@ -1,7 +1,7 @@
cmake_minimum_required(VERSION 2.8.7) cmake_minimum_required(VERSION 2.8.7)
project(ethash) project(ethash)
set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/Modules/") set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/modules/")
set(ETHHASH_LIBS ethash) set(ETHHASH_LIBS ethash)
if (WIN32 AND WANT_CRYPTOPP) if (WIN32 AND WANT_CRYPTOPP)
@ -10,6 +10,12 @@ endif()
add_subdirectory(src/libethash) add_subdirectory(src/libethash)
# bin2h.cmake doesn't work # bin2h.cmake doesn't work
#add_subdirectory(src/libethash-cl EXCLUDE_FROM_ALL) if (NOT OpenCL_FOUND)
find_package(OpenCL)
endif()
if (OpenCL_FOUND)
add_subdirectory(src/libethash-cl)
endif()
add_subdirectory(src/benchmark EXCLUDE_FROM_ALL) add_subdirectory(src/benchmark EXCLUDE_FROM_ALL)
add_subdirectory(test/c EXCLUDE_FROM_ALL) add_subdirectory(test/c EXCLUDE_FROM_ALL)

View File

@ -0,0 +1,17 @@
include setup.py
# C sources
include src/libethash/internal.c
include src/libethash/sha3.c
include src/libethash/util.c
include src/python/core.c
# Headers
include src/libethash/compiler.h
include src/libethash/data_sizes.h
include src/libethash/endian.h
include src/libethash/ethash.h
include src/libethash/fnv.h
include src/libethash/internal.h
include src/libethash/sha3.h
include src/libethash/util.h

View File

@ -1,3 +1,6 @@
.PHONY: clean .PHONY: clean test
test:
./test/test.sh
clean: clean:
rm -rf *.so pyethash.egg-info/ build/ test/python/python-virtual-env/ test/c/build/ pyethash/*.{so,pyc} rm -rf *.so pyethash.egg-info/ build/ test/python/python-virtual-env/ test/c/build/ pyethash.so test/python/*.pyc dist/ MANIFEST

View File

@ -8,6 +8,6 @@ file(GLOB SOURCE "../../cryptopp/*.cpp")
add_library(${LIBRARY} ${HEADERS} ${SOURCE}) add_library(${LIBRARY} ${HEADERS} ${SOURCE})
set(CRYPTOPP_INCLUDE_DIRS "../.." PARENT_SCOPE) set(CRYPTOPP_INCLUDE_DIRS "../.." "../../../" PARENT_SCOPE)
set(CRYPTOPP_LIBRARIES ${LIBRARY} PARENT_SCOPE) set(CRYPTOPP_LIBRARIES ${LIBRARY} PARENT_SCOPE)
set(CRYPTOPP_FOUND TRUE PARENT_SCOPE) set(CRYPTOPP_FOUND TRUE PARENT_SCOPE)

View File

@ -1,3 +1,13 @@
/*
###################################################################################
###################################################################################
#################### ####################
#################### EDIT AND YOU SHALL FEEL MY WRATH - jeff ####################
#################### ####################
###################################################################################
###################################################################################
*/
package ethash package ethash
/* /*
@ -11,8 +21,8 @@ import "C"
import ( import (
"bytes" "bytes"
"encoding/binary" "encoding/binary"
"fmt"
"io/ioutil" "io/ioutil"
"log"
"math/big" "math/big"
"math/rand" "math/rand"
"os" "os"
@ -21,25 +31,26 @@ import (
"time" "time"
"unsafe" "unsafe"
"github.com/ethereum/go-ethereum/crypto"
"github.com/ethereum/go-ethereum/ethutil" "github.com/ethereum/go-ethereum/ethutil"
"github.com/ethereum/go-ethereum/logger" "github.com/ethereum/go-ethereum/logger"
"github.com/ethereum/go-ethereum/pow" "github.com/ethereum/go-ethereum/pow"
) )
var tt256 = new(big.Int).Exp(big.NewInt(2), big.NewInt(256), big.NewInt(0)) var minDifficulty = new(big.Int).Exp(big.NewInt(2), big.NewInt(256), big.NewInt(0))
var powlogger = logger.NewLogger("POW") var powlogger = logger.NewLogger("POW")
type DAG struct { type ParamsAndCache struct {
SeedBlockNum uint64 params *C.ethash_params
dag unsafe.Pointer // full GB of memory for dag cache *C.ethash_cache
file bool Epoch uint64
} }
type ParamsAndCache struct { type DAG struct {
params *C.ethash_params dag unsafe.Pointer // full GB of memory for dag
cache *C.ethash_cache file bool
SeedBlockNum uint64 paramsAndCache *ParamsAndCache
} }
type Ethash struct { type Ethash struct {
@ -48,10 +59,9 @@ type Ethash struct {
chainManager pow.ChainManager chainManager pow.ChainManager
dag *DAG dag *DAG
paramsAndCache *ParamsAndCache paramsAndCache *ParamsAndCache
nextdag unsafe.Pointer
ret *C.ethash_return_value ret *C.ethash_return_value
dagMutex *sync.Mutex dagMutex *sync.RWMutex
cacheMutex *sync.Mutex cacheMutex *sync.RWMutex
} }
func parseNonce(nonce []byte) (uint64, error) { func parseNonce(nonce []byte) (uint64, error) {
@ -65,151 +75,212 @@ func parseNonce(nonce []byte) (uint64, error) {
const epochLength uint64 = 30000 const epochLength uint64 = 30000
func GetSeedBlockNum(blockNum uint64) uint64 { func makeParamsAndCache(chainManager pow.ChainManager, blockNum uint64) (*ParamsAndCache, error) {
var seedBlockNum uint64 = 0 if blockNum >= epochLength*2048 {
if blockNum > epochLength { return nil, fmt.Errorf("block number is out of bounds (value %v, limit is %v)", blockNum, epochLength*2048)
seedBlockNum = ((blockNum - 1) / epochLength) * epochLength
} }
return seedBlockNum
}
/*
XXX THIS DOESN'T WORK!! NEEDS FIXING
blockEpoch will underflow and wrap around causing massive issues
func GetSeedBlockNum(blockNum uint64) uint64 {
var seedBlockNum uint64 = 0
if blockNum > epochLength {
seedBlockNum = ((blockNum - 1) / epochLength) * epochLength
}
return seedBlockNum
}
*/
func makeParamsAndCache(chainManager pow.ChainManager, blockNum uint64) *ParamsAndCache {
seedBlockNum := GetSeedBlockNum(blockNum)
paramsAndCache := &ParamsAndCache{ paramsAndCache := &ParamsAndCache{
params: new(C.ethash_params), params: new(C.ethash_params),
cache: new(C.ethash_cache), cache: new(C.ethash_cache),
SeedBlockNum: seedBlockNum, Epoch: blockNum / epochLength,
} }
C.ethash_params_init(paramsAndCache.params, C.uint32_t(seedBlockNum)) C.ethash_params_init(paramsAndCache.params, C.uint32_t(uint32(blockNum)))
paramsAndCache.cache.mem = C.malloc(paramsAndCache.params.cache_size) paramsAndCache.cache.mem = C.malloc(paramsAndCache.params.cache_size)
seedHash := chainManager.GetBlockByNumber(seedBlockNum).SeedHash()
log.Println("Making Cache") seedHash, err := GetSeedHash(blockNum)
if err != nil {
return nil, err
}
powlogger.Infoln("Making Cache")
start := time.Now() start := time.Now()
C.ethash_mkcache(paramsAndCache.cache, paramsAndCache.params, (*C.uint8_t)(unsafe.Pointer(&seedHash[0]))) C.ethash_mkcache(paramsAndCache.cache, paramsAndCache.params, (*C.uint8_t)(unsafe.Pointer(&seedHash[0])))
log.Println("Took:", time.Since(start)) powlogger.Infoln("Took:", time.Since(start))
return paramsAndCache return paramsAndCache, nil
} }
func (pow *Ethash) updateCache() { func (pow *Ethash) UpdateCache(force bool) error {
pow.cacheMutex.Lock() pow.cacheMutex.Lock()
seedNum := GetSeedBlockNum(pow.chainManager.CurrentBlock().NumberU64()) thisEpoch := pow.chainManager.CurrentBlock().NumberU64()
if pow.paramsAndCache.SeedBlockNum != seedNum { if force || pow.paramsAndCache.Epoch != thisEpoch {
pow.paramsAndCache = makeParamsAndCache(pow.chainManager, pow.chainManager.CurrentBlock().NumberU64()) var err error
pow.paramsAndCache, err = makeParamsAndCache(pow.chainManager, pow.chainManager.CurrentBlock().NumberU64())
if err != nil {
panic(err)
}
} }
pow.cacheMutex.Unlock() pow.cacheMutex.Unlock()
return nil
} }
func makeDAG(p *ParamsAndCache) *DAG { func makeDAG(p *ParamsAndCache) *DAG {
d := &DAG{ d := &DAG{
dag: C.malloc(p.params.full_size), dag: C.malloc(p.params.full_size),
SeedBlockNum: p.SeedBlockNum, file: false,
paramsAndCache: p,
} }
donech := make(chan string)
go func() {
t := time.NewTicker(5 * time.Second)
tstart := time.Now()
done:
for {
select {
case <-t.C:
powlogger.Infof("... still generating DAG (%v) ...\n", time.Since(tstart).Seconds())
case str := <-donech:
powlogger.Infof("... %s ...\n", str)
break done
}
}
}()
C.ethash_compute_full_data(d.dag, p.params, p.cache) C.ethash_compute_full_data(d.dag, p.params, p.cache)
donech <- "DAG generation completed"
return d return d
} }
func (pow *Ethash) writeDagToDisk(dag *DAG, seedNum uint64) *os.File { func (pow *Ethash) writeDagToDisk(dag *DAG, epoch uint64) *os.File {
data := C.GoBytes(unsafe.Pointer(dag.dag), C.int(pow.paramsAndCache.params.full_size)) if epoch > 2048 {
panic(fmt.Errorf("Epoch must be less than 2048 (is %v)", epoch))
}
data := C.GoBytes(unsafe.Pointer(dag.dag), C.int(dag.paramsAndCache.params.full_size))
file, err := os.Create("/tmp/dag") file, err := os.Create("/tmp/dag")
if err != nil { if err != nil {
panic(err) panic(err)
} }
num := make([]byte, 8) dataEpoch := make([]byte, 8)
binary.BigEndian.PutUint64(num, seedNum) binary.BigEndian.PutUint64(dataEpoch, epoch)
file.Write(num) file.Write(dataEpoch)
file.Write(data) file.Write(data)
return file return file
} }
func (pow *Ethash) UpdateDAG() { func (pow *Ethash) UpdateDAG() {
pow.cacheMutex.Lock() blockNum := pow.chainManager.CurrentBlock().NumberU64()
pow.dagMutex.Lock() if blockNum >= epochLength*2048 {
// This will crash in the 2030s or 2040s
panic(fmt.Errorf("Current block number is out of bounds (value %v, limit is %v)", blockNum, epochLength*2048))
}
seedNum := GetSeedBlockNum(pow.chainManager.CurrentBlock().NumberU64()) pow.dagMutex.Lock()
if pow.dag == nil || pow.dag.SeedBlockNum != seedNum { defer pow.dagMutex.Unlock()
thisEpoch := blockNum / epochLength
if pow.dag == nil || pow.dag.paramsAndCache.Epoch != thisEpoch {
if pow.dag != nil && pow.dag.dag != nil { if pow.dag != nil && pow.dag.dag != nil {
C.free(pow.dag.dag) C.free(pow.dag.dag)
pow.dag.dag = nil pow.dag.dag = nil
} }
if pow.dag != nil && pow.dag.paramsAndCache.cache.mem != nil {
C.free(pow.dag.paramsAndCache.cache.mem)
pow.dag.paramsAndCache.cache.mem = nil
}
// Make the params and cache for the DAG
paramsAndCache, err := makeParamsAndCache(pow.chainManager, blockNum)
if err != nil {
panic(err)
}
// TODO: On non-SSD disks, loading the DAG from disk takes longer than generating it in memory
pow.paramsAndCache = paramsAndCache
path := path.Join("/", "tmp", "dag") path := path.Join("/", "tmp", "dag")
pow.dag = nil pow.dag = nil
log.Println("Generating dag") powlogger.Infoln("Retrieving DAG")
start := time.Now() start := time.Now()
file, err := os.Open(path) file, err := os.Open(path)
if err != nil { if err != nil {
log.Printf("No dag found in '%s'. Generating new dago(takes a while)...", path) powlogger.Infof("No DAG found. Generating new DAG in '%s' (this takes a while)...\n", path)
pow.dag = makeDAG(pow.paramsAndCache) pow.dag = makeDAG(paramsAndCache)
file = pow.writeDagToDisk(pow.dag, seedNum) file = pow.writeDagToDisk(pow.dag, thisEpoch)
pow.dag.file = true
} else { } else {
data, err := ioutil.ReadAll(file) data, err := ioutil.ReadAll(file)
if err != nil { if err != nil {
panic(err) powlogger.Infof("DAG load err: %v\n", err)
} }
num := binary.BigEndian.Uint64(data[0:8]) if len(data) < 8 {
if num < seedNum { powlogger.Infof("DAG in '%s' is less than 8 bytes, it must be corrupted. Generating new DAG (this takes a while)...\n", path)
log.Printf("Old found. Generating new dag (takes a while)...") pow.dag = makeDAG(paramsAndCache)
pow.dag = makeDAG(pow.paramsAndCache) file = pow.writeDagToDisk(pow.dag, thisEpoch)
file = pow.writeDagToDisk(pow.dag, seedNum) pow.dag.file = true
} else { } else {
data = data[8:] dataEpoch := binary.BigEndian.Uint64(data[0:8])
pow.dag = &DAG{ if dataEpoch < thisEpoch {
dag: unsafe.Pointer(&data[0]), powlogger.Infof("DAG in '%s' is stale. Generating new DAG (this takes a while)...\n", path)
file: true, pow.dag = makeDAG(paramsAndCache)
SeedBlockNum: pow.paramsAndCache.SeedBlockNum, file = pow.writeDagToDisk(pow.dag, thisEpoch)
pow.dag.file = true
} else if dataEpoch > thisEpoch {
// FIXME
panic(fmt.Errorf("Saved DAG in '%s' reports to be from future epoch %v (current epoch is %v)\n", path, dataEpoch, thisEpoch))
} else if len(data) != (int(paramsAndCache.params.full_size) + 8) {
powlogger.Infof("DAG in '%s' is corrupted. Generating new DAG (this takes a while)...\n", path)
pow.dag = makeDAG(paramsAndCache)
file = pow.writeDagToDisk(pow.dag, thisEpoch)
pow.dag.file = true
} else {
data = data[8:]
pow.dag = &DAG{
dag: unsafe.Pointer(&data[0]),
file: true,
paramsAndCache: paramsAndCache,
}
} }
} }
} }
log.Println("Took:", time.Since(start)) powlogger.Infoln("Took:", time.Since(start))
file.Close() file.Close()
} }
pow.dagMutex.Unlock()
pow.cacheMutex.Unlock()
} }
func New(chainManager pow.ChainManager) *Ethash { func New(chainManager pow.ChainManager) *Ethash {
paramsAndCache, err := makeParamsAndCache(chainManager, chainManager.CurrentBlock().NumberU64())
if err != nil {
panic(err)
}
return &Ethash{ return &Ethash{
turbo: true, turbo: true,
paramsAndCache: makeParamsAndCache(chainManager, chainManager.CurrentBlock().NumberU64()), paramsAndCache: paramsAndCache,
chainManager: chainManager, chainManager: chainManager,
dag: nil, dag: nil,
cacheMutex: new(sync.Mutex), cacheMutex: new(sync.RWMutex),
dagMutex: new(sync.Mutex), dagMutex: new(sync.RWMutex),
} }
} }
func (pow *Ethash) DAGSize() uint64 { func (pow *Ethash) DAGSize() uint64 {
return uint64(pow.paramsAndCache.params.full_size) return uint64(pow.dag.paramsAndCache.params.full_size)
} }
func (pow *Ethash) CacheSize() uint64 { func (pow *Ethash) CacheSize() uint64 {
return uint64(pow.paramsAndCache.params.cache_size) return uint64(pow.paramsAndCache.params.cache_size)
} }
func (pow *Ethash) GetSeedHash(blockNum uint64) []byte { func GetSeedHash(blockNum uint64) ([]byte, error) {
seednum := GetSeedBlockNum(blockNum) if blockNum >= epochLength*2048 {
return pow.chainManager.GetBlockByNumber(seednum).SeedHash() return nil, fmt.Errorf("block number is out of bounds (value %v, limit is %v)", blockNum, epochLength*2048)
}
epoch := blockNum / epochLength
seedHash := make([]byte, 32)
var i uint64
for i = 0; i < 32; i++ {
seedHash[i] = 0
}
for i = 0; i < epoch; i++ {
seedHash = crypto.Sha3(seedHash)
}
return seedHash, nil
} }
func (pow *Ethash) Stop() { func (pow *Ethash) Stop() {
@ -224,17 +295,18 @@ func (pow *Ethash) Stop() {
if pow.dag.dag != nil && !pow.dag.file { if pow.dag.dag != nil && !pow.dag.file {
C.free(pow.dag.dag) C.free(pow.dag.dag)
} }
if pow.dag != nil && pow.dag.paramsAndCache != nil && pow.dag.paramsAndCache.cache.mem != nil {
C.free(pow.dag.paramsAndCache.cache.mem)
pow.dag.paramsAndCache.cache.mem = nil
}
pow.dag.dag = nil pow.dag.dag = nil
} }
func (pow *Ethash) Search(block pow.Block, stop <-chan struct{}) (uint64, []byte, []byte) { func (pow *Ethash) Search(block pow.Block, stop <-chan struct{}) (uint64, []byte, []byte) {
//pow.UpdateDAG() pow.UpdateDAG()
// Not very elegant, multiple mining instances are not supported pow.dagMutex.RLock()
//pow.dagMutex.Lock() defer pow.dagMutex.RUnlock()
//pow.cacheMutex.Lock()
//defer pow.cacheMutex.Unlock()
//defer pow.dagMutex.Unlock()
r := rand.New(rand.NewSource(time.Now().UnixNano())) r := rand.New(rand.NewSource(time.Now().UnixNano()))
miningHash := block.HashNoNonce() miningHash := block.HashNoNonce()
@ -246,7 +318,7 @@ func (pow *Ethash) Search(block pow.Block, stop <-chan struct{}) (uint64, []byte
nonce := uint64(r.Int63()) nonce := uint64(r.Int63())
cMiningHash := (*C.uint8_t)(unsafe.Pointer(&miningHash[0])) cMiningHash := (*C.uint8_t)(unsafe.Pointer(&miningHash[0]))
target := new(big.Int).Div(tt256, diff) target := new(big.Int).Div(minDifficulty, diff)
var ret C.ethash_return_value var ret C.ethash_return_value
for { for {
@ -262,14 +334,17 @@ func (pow *Ethash) Search(block pow.Block, stop <-chan struct{}) (uint64, []byte
hashes := ((float64(1e9) / float64(elapsed)) * float64(i-starti)) / 1000 hashes := ((float64(1e9) / float64(elapsed)) * float64(i-starti)) / 1000
pow.HashRate = int64(hashes) pow.HashRate = int64(hashes)
C.ethash_full(&ret, pow.dag.dag, pow.paramsAndCache.params, cMiningHash, C.uint64_t(nonce)) C.ethash_full(&ret, pow.dag.dag, pow.dag.paramsAndCache.params, cMiningHash, C.uint64_t(nonce))
result := ethutil.Bytes2Big(C.GoBytes(unsafe.Pointer(&ret.result[0]), C.int(32))) result := ethutil.Bytes2Big(C.GoBytes(unsafe.Pointer(&ret.result[0]), C.int(32)))
// TODO: disagrees with the spec https://github.com/ethereum/wiki/wiki/Ethash#mining
if result.Cmp(target) <= 0 { if result.Cmp(target) <= 0 {
mixDigest := C.GoBytes(unsafe.Pointer(&ret.mix_hash[0]), C.int(32)) mixDigest := C.GoBytes(unsafe.Pointer(&ret.mix_hash[0]), C.int(32))
seedHash, err := GetSeedHash(block.NumberU64()) // This seedhash is useless
return nonce, mixDigest, pow.GetSeedHash(block.NumberU64()) if err != nil {
panic(err)
}
return nonce, mixDigest, seedHash
} }
nonce += 1 nonce += 1
@ -283,32 +358,41 @@ func (pow *Ethash) Search(block pow.Block, stop <-chan struct{}) (uint64, []byte
} }
func (pow *Ethash) Verify(block pow.Block) bool { func (pow *Ethash) Verify(block pow.Block) bool {
// Make sure the SeedHash is set correctly
if bytes.Compare(block.SeedHash(), pow.GetSeedHash(block.NumberU64())) != 0 {
return false
}
return pow.verify(block.HashNoNonce(), block.MixDigest(), block.Difficulty(), block.NumberU64(), block.Nonce()) return pow.verify(block.HashNoNonce(), block.MixDigest(), block.Difficulty(), block.NumberU64(), block.Nonce())
} }
func (pow *Ethash) verify(hash []byte, mixDigest []byte, difficulty *big.Int, blockNum uint64, nonce uint64) bool { func (pow *Ethash) verify(hash []byte, mixDigest []byte, difficulty *big.Int, blockNum uint64, nonce uint64) bool {
// First check: make sure header, mixDigest, nonce are correct without hitting the DAG // Make sure the block num is valid
if blockNum >= epochLength*2048 {
powlogger.Infoln(fmt.Sprintf("Block number exceeds limit, invalid (value is %v, limit is %v)",
blockNum, epochLength*2048))
return false
}
// First check: make sure header, mixDigest, nonce are correct without hitting the cache
// This is to prevent DOS attacks // This is to prevent DOS attacks
chash := (*C.uint8_t)(unsafe.Pointer(&hash[0])) chash := (*C.uint8_t)(unsafe.Pointer(&hash[0]))
cnonce := C.uint64_t(nonce) cnonce := C.uint64_t(nonce)
target := new(big.Int).Div(tt256, difficulty) target := new(big.Int).Div(minDifficulty, difficulty)
var pAc *ParamsAndCache var pAc *ParamsAndCache
// If its an old block (doesn't use the current cache) // If its an old block (doesn't use the current cache)
// get the cache for it but don't update (so we don't need the mutex) // get the cache for it but don't update (so we don't need the mutex)
// Otherwise, it's the current block or a future. // Otherwise, it's the current block or a future block.
// If current, updateCache will do nothing. // If current, updateCache will do nothing.
if GetSeedBlockNum(blockNum) < pow.paramsAndCache.SeedBlockNum { if blockNum/epochLength < pow.paramsAndCache.Epoch {
pAc = makeParamsAndCache(pow.chainManager, blockNum) var err error
// If we can't make the params for some reason, this block is invalid
pAc, err = makeParamsAndCache(pow.chainManager, blockNum)
if err != nil {
powlogger.Infoln(err)
return false
}
} else { } else {
pow.updateCache() pow.UpdateCache(false)
pow.cacheMutex.Lock() pow.cacheMutex.RLock()
defer pow.cacheMutex.Unlock() defer pow.cacheMutex.RUnlock()
pAc = pow.paramsAndCache pAc = pow.paramsAndCache
} }

View File

@ -1,2 +0,0 @@
pyethash.egg-info/
*.so

View File

@ -1,3 +0,0 @@
import pyethash.core
core = pyethash.core
EPOCH_LENGTH = 30000

View File

@ -1,21 +1,33 @@
#!/usr/bin/env python #!/usr/bin/env python
from distutils.core import setup, Extension from distutils.core import setup, Extension
pyethash_core = Extension('pyethash.core', pyethash = Extension('pyethash',
sources = [ sources = [
'src/python/core.c', 'src/python/core.c',
'src/libethash/util.c', 'src/libethash/util.c',
'src/libethash/internal.c', 'src/libethash/internal.c',
'src/libethash/sha3.c' 'src/libethash/sha3.c'],
depends = [
'src/libethash/ethash.h',
'src/libethash/compiler.h',
'src/libethash/data_sizes.h',
'src/libethash/endian.h',
'src/libethash/ethash.h',
'src/libethash/fnv.h',
'src/libethash/internal.h',
'src/libethash/sha3.h',
'src/libethash/util.h'
], ],
extra_compile_args = ["-std=gnu99"]) extra_compile_args = ["-Isrc/", "-std=gnu99", "-Wall"])
setup ( setup (
name = 'pyethash', name = 'pyethash',
author = "Matthew Wampler-Doty", author = "Matthew Wampler-Doty",
author_email = "matthew.wampler.doty@gmail.com", author_email = "matthew.wampler.doty@gmail.com",
license = 'GPL', license = 'GPL',
version = '1.0', version = '23',
url = 'https://github.com/ethereum/ethash',
download_url = 'https://github.com/ethereum/ethash/tarball/v23',
description = 'Python wrappers for ethash, the ethereum proof of work hashing function', description = 'Python wrappers for ethash, the ethereum proof of work hashing function',
ext_modules = [pyethash_core], ext_modules = [pyethash],
) )

View File

@ -31,7 +31,7 @@
#include <algorithm> #include <algorithm>
#ifdef WITH_CRYPTOPP #ifdef WITH_CRYPTOPP
#include <libethash/SHA3_cryptopp.h> #include <libethash/sha3_cryptopp.h>
#include <string> #include <string>
#else #else

View File

@ -2,14 +2,39 @@ set(LIBRARY ethash-cl)
set(CMAKE_BUILD_TYPE Release) set(CMAKE_BUILD_TYPE Release)
include(bin2h.cmake) include(bin2h.cmake)
bin2h(SOURCE_FILE ethash_cl_miner_kernel.cl VARIABLE_NAME ethash_cl_miner_kernel HEADER_FILE ${CMAKE_CURRENT_BINARY_DIR}/ethash_cl_miner_kernel.h) bin2h(SOURCE_FILE ethash_cl_miner_kernel.cl
VARIABLE_NAME ethash_cl_miner_kernel
HEADER_FILE ${CMAKE_CURRENT_BINARY_DIR}/ethash_cl_miner_kernel.h)
if (NOT MSVC)
# Initialize CXXFLAGS for c++11
set(CMAKE_CXX_FLAGS "-Wall -std=c++11")
set(CMAKE_CXX_FLAGS_DEBUG "-O0 -g")
set(CMAKE_CXX_FLAGS_MINSIZEREL "-Os -DNDEBUG")
set(CMAKE_CXX_FLAGS_RELEASE "-O4 -DNDEBUG")
set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "-O2 -g")
# Compiler-specific C++11 activation.
if ("${CMAKE_CXX_COMPILER_ID}" MATCHES "GNU")
execute_process(
COMMAND ${CMAKE_CXX_COMPILER} -dumpversion OUTPUT_VARIABLE GCC_VERSION)
if (NOT (GCC_VERSION VERSION_GREATER 4.7 OR GCC_VERSION VERSION_EQUAL 4.7))
message(FATAL_ERROR "${PROJECT_NAME} requires g++ 4.7 or greater.")
endif ()
elseif ("${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -stdlib=libc++")
else ()
message(FATAL_ERROR "Your C++ compiler does not support C++11.")
endif ()
endif()
if (NOT OpenCL_FOUND) if (NOT OpenCL_FOUND)
find_package(OpenCL) find_package(OpenCL)
endif() endif()
if (OpenCL_FOUND) if (OpenCL_FOUND)
include_directories(${OpenCL_INCLUDE_DIRS} ${CMAKE_CURRENT_BINARY_DIR}) include_directories(${OpenCL_INCLUDE_DIRS} ${CMAKE_CURRENT_BINARY_DIR})
include_directories(..) include_directories(..)
add_library(${LIBRARY} ethash_cl_miner.cpp ethash_cl_miner.h) add_library(${LIBRARY} ethash_cl_miner.cpp ethash_cl_miner.h cl.hpp)
TARGET_LINK_LIBRARIES(${LIBRARY} ${OpenCL_LIBRARIES} ethash) TARGET_LINK_LIBRARIES(${LIBRARY} ${OpenCL_LIBRARIES} ethash)
endif() endif()

View File

@ -68,8 +68,7 @@ function(BIN2H)
string(REGEX REPLACE ", $" "" arrayValues ${arrayValues}) string(REGEX REPLACE ", $" "" arrayValues ${arrayValues})
# converts the variable name into proper C identifier # converts the variable name into proper C identifier
# TODO: fix for legacy cmake IF (${CMAKE_VERSION} GREATER 2.8.10) # fix for legacy cmake
IF (${CMAKE_VERSION} GREATER 2.8.10)
string(MAKE_C_IDENTIFIER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME) string(MAKE_C_IDENTIFIER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME)
ENDIF() ENDIF()
string(TOUPPER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME) string(TOUPPER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME)

View File

@ -24,15 +24,22 @@
#include <assert.h> #include <assert.h>
#include <queue> #include <queue>
#include <vector>
#include "ethash_cl_miner.h" #include "ethash_cl_miner.h"
#include "ethash_cl_miner_kernel.h" #include "ethash_cl_miner_kernel.h"
#include <libethash/util.h> #include <libethash/util.h>
#define ETHASH_BYTES 32
// workaround lame platforms
#if !CL_VERSION_1_2
#define CL_MAP_WRITE_INVALIDATE_REGION CL_MAP_WRITE
#define CL_MEM_HOST_READ_ONLY 0
#endif
#undef min #undef min
#undef max #undef max
#define HASH_BYTES 32
static void add_definition(std::string& source, char const* id, unsigned value) static void add_definition(std::string& source, char const* id, unsigned value)
{ {
char buf[256]; char buf[256];
@ -41,6 +48,7 @@ static void add_definition(std::string& source, char const* id, unsigned value)
} }
ethash_cl_miner::ethash_cl_miner() ethash_cl_miner::ethash_cl_miner()
: m_opencl_1_1()
{ {
} }
@ -71,11 +79,23 @@ bool ethash_cl_miner::init(ethash_params const& params, const uint8_t seed[32],
} }
// use default device // use default device
cl::Device& device = devices[0]; unsigned device_num = 0;
debugf("Using device: %s\n", device.getInfo<CL_DEVICE_NAME>().c_str()); cl::Device& device = devices[device_num];
std::string device_version = device.getInfo<CL_DEVICE_VERSION>();
debugf("Using device: %s (%s)\n", device.getInfo<CL_DEVICE_NAME>().c_str(),device_version.c_str());
if (strncmp("OpenCL 1.0", device_version.c_str(), 10) == 0)
{
debugf("OpenCL 1.0 is not supported.\n");
return false;
}
if (strncmp("OpenCL 1.1", device_version.c_str(), 10) == 0)
{
m_opencl_1_1 = true;
}
// create context // create context
m_context = cl::Context({device}); m_context = cl::Context(std::vector<cl::Device>(&device, &device+1));
m_queue = cl::CommandQueue(m_context, device); m_queue = cl::CommandQueue(m_context, device);
// use requested workgroup size, but we require multiple of 8 // use requested workgroup size, but we require multiple of 8
@ -120,7 +140,7 @@ bool ethash_cl_miner::init(ethash_params const& params, const uint8_t seed[32],
ethash_mkcache(&cache, &params, seed); ethash_mkcache(&cache, &params, seed);
// if this throws then it's because we probably need to subdivide the dag uploads for compatibility // if this throws then it's because we probably need to subdivide the dag uploads for compatibility
void* dag_ptr = m_queue.enqueueMapBuffer(m_dag, true, CL_MAP_WRITE_INVALIDATE_REGION, 0, params.full_size); void* dag_ptr = m_queue.enqueueMapBuffer(m_dag, true, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, params.full_size);
ethash_compute_full_data(dag_ptr, &params, &cache); ethash_compute_full_data(dag_ptr, &params, &cache);
m_queue.enqueueUnmapMemObject(m_dag, dag_ptr); m_queue.enqueueUnmapMemObject(m_dag, dag_ptr);
@ -130,7 +150,7 @@ bool ethash_cl_miner::init(ethash_params const& params, const uint8_t seed[32],
// create mining buffers // create mining buffers
for (unsigned i = 0; i != c_num_buffers; ++i) for (unsigned i = 0; i != c_num_buffers; ++i)
{ {
m_hash_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY | CL_MEM_HOST_READ_ONLY, 32*c_hash_batch_size); m_hash_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY | (!m_opencl_1_1 ? CL_MEM_HOST_READ_ONLY : 0), 32*c_hash_batch_size);
m_search_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_max_search_results + 1) * sizeof(uint32_t)); m_search_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_max_search_results + 1) * sizeof(uint32_t));
} }
return true; return true;
@ -176,7 +196,6 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce,
m_hash_kernel.setArg(0, m_hash_buf[buf]); m_hash_kernel.setArg(0, m_hash_buf[buf]);
// execute it! // execute it!
clock_t start_time = clock();
m_queue.enqueueNDRangeKernel( m_queue.enqueueNDRangeKernel(
m_hash_kernel, m_hash_kernel,
cl::NullRange, cl::NullRange,
@ -196,8 +215,8 @@ void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce,
pending_batch const& batch = pending.front(); pending_batch const& batch = pending.front();
// could use pinned host pointer instead, but this path isn't that important. // could use pinned host pointer instead, but this path isn't that important.
uint8_t* hashes = (uint8_t*)m_queue.enqueueMapBuffer(m_hash_buf[batch.buf], true, CL_MAP_READ, 0, batch.count * HASH_BYTES); uint8_t* hashes = (uint8_t*)m_queue.enqueueMapBuffer(m_hash_buf[batch.buf], true, CL_MAP_READ, 0, batch.count * ETHASH_BYTES);
memcpy(ret + batch.base*HASH_BYTES, hashes, batch.count*HASH_BYTES); memcpy(ret + batch.base*ETHASH_BYTES, hashes, batch.count*ETHASH_BYTES);
m_queue.enqueueUnmapMemObject(m_hash_buf[batch.buf], hashes); m_queue.enqueueUnmapMemObject(m_hash_buf[batch.buf], hashes);
pending.pop(); pending.pop();
@ -223,8 +242,19 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
{ {
m_queue.enqueueWriteBuffer(m_search_buf[i], false, 0, 4, &c_zero); m_queue.enqueueWriteBuffer(m_search_buf[i], false, 0, 4, &c_zero);
} }
#if CL_VERSION_1_2
cl::Event pre_return_event; cl::Event pre_return_event;
m_queue.enqueueBarrierWithWaitList(NULL, &pre_return_event); if (!m_opencl_1_1)
{
m_queue.enqueueBarrierWithWaitList(NULL, &pre_return_event);
}
else
#else
{
m_queue.finish();
}
#endif
/* /*
__kernel void ethash_combined_search( __kernel void ethash_combined_search(
@ -284,6 +314,11 @@ void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook
} }
// not safe to return until this is ready // not safe to return until this is ready
pre_return_event.wait(); #if CL_VERSION_1_2
if (!m_opencl_1_1)
{
pre_return_event.wait();
}
#endif
} }

View File

@ -40,4 +40,5 @@ private:
cl::Buffer m_hash_buf[c_num_buffers]; cl::Buffer m_hash_buf[c_num_buffers];
cl::Buffer m_search_buf[c_num_buffers]; cl::Buffer m_search_buf[c_num_buffers];
unsigned m_workgroup_size; unsigned m_workgroup_size;
bool m_opencl_1_1;
}; };

View File

@ -2,7 +2,6 @@ set(LIBRARY ethash)
if (CPPETHEREUM) if (CPPETHEREUM)
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fPIC") set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fPIC")
#else ()
endif () endif ()
set(CMAKE_BUILD_TYPE Release) set(CMAKE_BUILD_TYPE Release)

View File

@ -20,8 +20,6 @@
* @date 2015 * @date 2015
*/ */
// TODO: Update this after ~3.5 years
#pragma once #pragma once
#include <stdint.h> #include <stdint.h>
@ -33,23 +31,24 @@ extern "C" {
#include <stdint.h> #include <stdint.h>
// 2048 Epochs worth of tabulated DAG sizes // 2048 Epochs (~20 years) worth of tabulated DAG sizes
// Generated with the following Mathematica Code: // Generated with the following Mathematica Code:
// GetDataSizes[n_] := Module[{ // GetCacheSizes[n_] := Module[{
// DataSetSizeBytesInit = 2^30, // CacheSizeBytesInit = 2^24,
// MixBytes = 128, // CacheGrowth = 2^17,
// DataSetGrowth = 2^23, // HashBytes = 64,
// j = 0}, // j = 0},
// Reap[ // Reap[
// While[j < n, // While[j < n,
// Module[{i = // Module[{i =
// Floor[(DataSetSizeBytesInit + DataSetGrowth * j) / MixBytes]}, // Floor[(CacheSizeBytesInit + CacheGrowth * j) / HashBytes]},
// While[! PrimeQ[i], i--]; // While[! PrimeQ[i], i--];
// Sow[i*MixBytes]; j++]]]][[2]][[1]] // Sow[i*HashBytes]; j++]]]][[2]][[1]]
static const size_t dag_sizes[] = {
static const size_t dag_sizes[2048] = {
1073739904U, 1082130304U, 1090514816U, 1098906752U, 1107293056U, 1073739904U, 1082130304U, 1090514816U, 1098906752U, 1107293056U,
1115684224U, 1124070016U, 1132461952U, 1140849536U, 1149232768U, 1115684224U, 1124070016U, 1132461952U, 1140849536U, 1149232768U,
1157627776U, 1166013824U, 1174404736U, 1182786944U, 1191180416U, 1157627776U, 1166013824U, 1174404736U, 1182786944U, 1191180416U,
@ -478,300 +477,334 @@ static const size_t dag_sizes[] = {
// While[! PrimeQ[i], i--]; // While[! PrimeQ[i], i--];
// Sow[i*HashBytes]; j++]]]][[2]][[1]] // Sow[i*HashBytes]; j++]]]][[2]][[1]]
const size_t cache_sizes[] = { const size_t cache_sizes[2048] = {
1048384U, 1055552U, 1064512U, 1072832U, 1080896U, 1089344U, 1096768U, 16776896U, 16907456U, 17039296U, 17170112U, 17301056U, 17432512U, 17563072U,
1104448U, 1113664U, 1121216U, 1130176U, 1138624U, 1146304U, 1155008U, 17693888U, 17824192U, 17955904U, 18087488U, 18218176U, 18349504U, 18481088U,
1162816U, 1171264U, 1179328U, 1187392U, 1195456U, 1203392U, 1210816U, 18611392U, 18742336U, 18874304U, 19004224U, 19135936U, 19267264U, 19398208U,
1220416U, 1227712U, 1236416U, 1244608U, 1253312U, 1261376U, 1268416U, 19529408U, 19660096U, 19791424U, 19922752U, 20053952U, 20184896U, 20315968U,
1277632U, 1285696U, 1294016U, 1302208U, 1310656U, 1318336U, 1326784U, 20446912U, 20576576U, 20709184U, 20840384U, 20971072U, 21102272U, 21233216U,
1334848U, 1342912U, 1350848U, 1359808U, 1366208U, 1376192U, 1383488U, 21364544U, 21494848U, 21626816U, 21757376U, 21887552U, 22019392U, 22151104U,
1392448U, 1400384U, 1408832U, 1416512U, 1425344U, 1433408U, 1440704U, 22281536U, 22412224U, 22543936U, 22675264U, 22806464U, 22935872U, 23068096U,
1449664U, 1458112U, 1466048U, 1474496U, 1482688U, 1490752U, 1498688U, 23198272U, 23330752U, 23459008U, 23592512U, 23723968U, 23854912U, 23986112U,
1507136U, 1515328U, 1523264U, 1531456U, 1539904U, 1547584U, 1556288U, 24116672U, 24247616U, 24378688U, 24509504U, 24640832U, 24772544U, 24903488U,
1564352U, 1572544U, 1580608U, 1588544U, 1596992U, 1605568U, 1612096U, 25034432U, 25165376U, 25296704U, 25427392U, 25558592U, 25690048U, 25820096U,
1621952U, 1630144U, 1637696U, 1645888U, 1654336U, 1662784U, 1671104U, 25951936U, 26081728U, 26214208U, 26345024U, 26476096U, 26606656U, 26737472U,
1679168U, 1686848U, 1695296U, 1702208U, 1711168U, 1720256U, 1727552U, 26869184U, 26998208U, 27131584U, 27262528U, 27393728U, 27523904U, 27655744U,
1736128U, 1744576U, 1751488U, 1760576U, 1769408U, 1777472U, 1785664U, 27786688U, 27917888U, 28049344U, 28179904U, 28311488U, 28441792U, 28573504U,
1793984U, 1801664U, 1810112U, 1818304U, 1826624U, 1834816U, 1842752U, 28700864U, 28835648U, 28966208U, 29096768U, 29228608U, 29359808U, 29490752U,
1851328U, 1858112U, 1867456U, 1875904U, 1883968U, 1892288U, 1899712U, 29621824U, 29752256U, 29882816U, 30014912U, 30144448U, 30273728U, 30406976U,
1908416U, 1916608U, 1924544U, 1932992U, 1940672U, 1948736U, 1956928U, 30538432U, 30670784U, 30799936U, 30932672U, 31063744U, 31195072U, 31325248U,
1965632U, 1973824U, 1982144U, 1989824U, 1998784U, 2006848U, 2014784U, 31456192U, 31588288U, 31719232U, 31850432U, 31981504U, 32110784U, 32243392U,
2022848U, 2031424U, 2038976U, 2047424U, 2055616U, 2064064U, 2072384U, 32372672U, 32505664U, 32636608U, 32767808U, 32897344U, 33029824U, 33160768U,
2080448U, 2088512U, 2095936U, 2104768U, 2113472U, 2121664U, 2127808U, 33289664U, 33423296U, 33554368U, 33683648U, 33816512U, 33947456U, 34076992U,
2137792U, 2146112U, 2153408U, 2162624U, 2170304U, 2178496U, 2186944U, 34208704U, 34340032U, 34471744U, 34600256U, 34734016U, 34864576U, 34993984U,
2195392U, 2203456U, 2211136U, 2219968U, 2227648U, 2236096U, 2244416U, 35127104U, 35258176U, 35386688U, 35518528U, 35650624U, 35782336U, 35910976U,
2250944U, 2260928U, 2268736U, 2276672U, 2283328U, 2293696U, 2301632U, 36044608U, 36175808U, 36305728U, 36436672U, 36568384U, 36699968U, 36830656U,
2309312U, 2317888U, 2325952U, 2334656U, 2342848U, 2350144U, 2358848U, 36961984U, 37093312U, 37223488U, 37355072U, 37486528U, 37617472U, 37747904U,
2366656U, 2375488U, 2383552U, 2391616U, 2400064U, 2407616U, 2415808U, 37879232U, 38009792U, 38141888U, 38272448U, 38403392U, 38535104U, 38660672U,
2424256U, 2432704U, 2439616U, 2448704U, 2457152U, 2464064U, 2473792U, 38795584U, 38925632U, 39059264U, 39190336U, 39320768U, 39452096U, 39581632U,
2482112U, 2489792U, 2497472U, 2506432U, 2514752U, 2522816U, 2531264U, 39713984U, 39844928U, 39974848U, 40107968U, 40238144U, 40367168U, 40500032U,
2539456U, 2547136U, 2555456U, 2564032U, 2572096U, 2578496U, 2587712U, 40631744U, 40762816U, 40894144U, 41023552U, 41155904U, 41286208U, 41418304U,
2595776U, 2604736U, 2613056U, 2620736U, 2629184U, 2637632U, 2645824U, 41547712U, 41680448U, 41811904U, 41942848U, 42073792U, 42204992U, 42334912U,
2653888U, 2662208U, 2670016U, 2678464U, 2686912U, 2694464U, 2703296U, 42467008U, 42597824U, 42729152U, 42860096U, 42991552U, 43122368U, 43253696U,
2710976U, 2719424U, 2727104U, 2736064U, 2743232U, 2752192U, 2760512U, 43382848U, 43515712U, 43646912U, 43777088U, 43907648U, 44039104U, 44170432U,
2768704U, 2777024U, 2785088U, 2792512U, 2800576U, 2809024U, 2817856U, 44302144U, 44433344U, 44564288U, 44694976U, 44825152U, 44956864U, 45088448U,
2826176U, 2833984U, 2840896U, 2850752U, 2858048U, 2867008U, 2875328U, 45219008U, 45350464U, 45481024U, 45612608U, 45744064U, 45874496U, 46006208U,
2883392U, 2891584U, 2899648U, 2908096U, 2915648U, 2924224U, 2932672U, 46136768U, 46267712U, 46399424U, 46529344U, 46660672U, 46791488U, 46923328U,
2940736U, 2948672U, 2956736U, 2964928U, 2973248U, 2981824U, 2988992U, 47053504U, 47185856U, 47316928U, 47447872U, 47579072U, 47710144U, 47839936U,
2997184U, 3005248U, 3013952U, 3022144U, 3030592U, 3037376U, 3046976U, 47971648U, 48103232U, 48234176U, 48365248U, 48496192U, 48627136U, 48757312U,
3055552U, 3063616U, 3070784U, 3079744U, 3087808U, 3096512U, 3103808U, 48889664U, 49020736U, 49149248U, 49283008U, 49413824U, 49545152U, 49675712U,
3111872U, 3121088U, 3128896U, 3137216U, 3144896U, 3153856U, 3161152U, 49807168U, 49938368U, 50069056U, 50200256U, 50331584U, 50462656U, 50593472U,
3169984U, 3178432U, 3186496U, 3194816U, 3203008U, 3210176U, 3218624U, 50724032U, 50853952U, 50986048U, 51117632U, 51248576U, 51379904U, 51510848U,
3227072U, 3235264U, 3243712U, 3250496U, 3259456U, 3268544U, 3276736U, 51641792U, 51773248U, 51903296U, 52035136U, 52164032U, 52297664U, 52427968U,
3283648U, 3292736U, 3301184U, 3308224U, 3317696U, 3324736U, 3333184U, 52557376U, 52690112U, 52821952U, 52952896U, 53081536U, 53213504U, 53344576U,
3342272U, 3348544U, 3357248U, 3365312U, 3374912U, 3383104U, 3390784U, 53475776U, 53608384U, 53738816U, 53870528U, 54000832U, 54131776U, 54263744U,
3399488U, 3407296U, 3414976U, 3424192U, 3432256U, 3440576U, 3448768U, 54394688U, 54525248U, 54655936U, 54787904U, 54918592U, 55049152U, 55181248U,
3456832U, 3464896U, 3473216U, 3480128U, 3489344U, 3497408U, 3505856U, 55312064U, 55442752U, 55574336U, 55705024U, 55836224U, 55967168U, 56097856U,
3514048U, 3521344U, 3530432U, 3538624U, 3546304U, 3555008U, 3563072U, 56228672U, 56358592U, 56490176U, 56621888U, 56753728U, 56884928U, 57015488U,
3571648U, 3579712U, 3587392U, 3595456U, 3603904U, 3612352U, 3620416U, 57146816U, 57278272U, 57409216U, 57540416U, 57671104U, 57802432U, 57933632U,
3628864U, 3636928U, 3645248U, 3652928U, 3660992U, 3669184U, 3677888U, 58064576U, 58195264U, 58326976U, 58457408U, 58588864U, 58720192U, 58849984U,
3685952U, 3694528U, 3702592U, 3710528U, 3719104U, 3727168U, 3735488U, 58981696U, 59113024U, 59243456U, 59375552U, 59506624U, 59637568U, 59768512U,
3742784U, 3751232U, 3759424U, 3765184U, 3775808U, 3783872U, 3792832U, 59897792U, 60030016U, 60161984U, 60293056U, 60423872U, 60554432U, 60683968U,
3800768U, 3808832U, 3816256U, 3825344U, 3832768U, 3841856U, 3849536U, 60817216U, 60948032U, 61079488U, 61209664U, 61341376U, 61471936U, 61602752U,
3857344U, 3866432U, 3874496U, 3882304U, 3890752U, 3899072U, 3907264U, 61733696U, 61865792U, 61996736U, 62127808U, 62259136U, 62389568U, 62520512U,
3914816U, 3923008U, 3930688U, 3939904U, 3947968U, 3956416U, 3964736U, 62651584U, 62781632U, 62910784U, 63045056U, 63176128U, 63307072U, 63438656U,
3972544U, 3981248U, 3988928U, 3997376U, 4005824U, 4012864U, 4020928U, 63569216U, 63700928U, 63831616U, 63960896U, 64093888U, 64225088U, 64355392U,
4030144U, 4038592U, 4045504U, 4054592U, 4063168U, 4071104U, 4079552U, 64486976U, 64617664U, 64748608U, 64879424U, 65009216U, 65142464U, 65273792U,
4087232U, 4095808U, 4103872U, 4111168U, 4120384U, 4127936U, 4136512U, 65402816U, 65535424U, 65666752U, 65797696U, 65927744U, 66060224U, 66191296U,
4144832U, 4153024U, 4160704U, 4169408U, 4177216U, 4186048U, 4193344U, 66321344U, 66453056U, 66584384U, 66715328U, 66846656U, 66977728U, 67108672U,
4202048U, 4210496U, 4217536U, 4227008U, 4235072U, 4243264U, 4251584U, 67239104U, 67370432U, 67501888U, 67631296U, 67763776U, 67895104U, 68026304U,
4259392U, 4267712U, 4275776U, 4284352U, 4291904U, 4300096U, 4307648U, 68157248U, 68287936U, 68419264U, 68548288U, 68681408U, 68811968U, 68942912U,
4316992U, 4325056U, 4333376U, 4341056U, 4349888U, 4357568U, 4366016U, 69074624U, 69205568U, 69337024U, 69467584U, 69599168U, 69729472U, 69861184U,
4374464U, 4382528U, 4390208U, 4398656U, 4407232U, 4413632U, 4423616U, 69989824U, 70122944U, 70253888U, 70385344U, 70515904U, 70647232U, 70778816U,
4431808U, 4439744U, 4447936U, 4455872U, 4463296U, 4472128U, 4480576U, 70907968U, 71040832U, 71171648U, 71303104U, 71432512U, 71564992U, 71695168U,
4489024U, 4497344U, 4505152U, 4512448U, 4520896U, 4530112U, 4537664U, 71826368U, 71958464U, 72089536U, 72219712U, 72350144U, 72482624U, 72613568U,
4546496U, 4554688U, 4562752U, 4570816U, 4579264U, 4586944U, 4595648U, 72744512U, 72875584U, 73006144U, 73138112U, 73268672U, 73400128U, 73530944U,
4603712U, 4611392U, 4619072U, 4628032U, 4635584U, 4643776U, 4652864U, 73662272U, 73793344U, 73924544U, 74055104U, 74185792U, 74316992U, 74448832U,
4660672U, 4669376U, 4677056U, 4684096U, 4693184U, 4702144U, 4710208U, 74579392U, 74710976U, 74841664U, 74972864U, 75102784U, 75233344U, 75364544U,
4718528U, 4726336U, 4734272U, 4742464U, 4750784U, 4759232U, 4767296U, 75497024U, 75627584U, 75759296U, 75890624U, 76021696U, 76152256U, 76283072U,
4775872U, 4783808U, 4791872U, 4797376U, 4808512U, 4816192U, 4825024U, 76414144U, 76545856U, 76676672U, 76806976U, 76937792U, 77070016U, 77200832U,
4832704U, 4841024U, 4849472U, 4856512U, 4865984U, 4874176U, 4882112U, 77331392U, 77462464U, 77593664U, 77725376U, 77856448U, 77987776U, 78118336U,
4889792U, 4898752U, 4906688U, 4913984U, 4922816U, 4931008U, 4938944U, 78249664U, 78380992U, 78511424U, 78642496U, 78773056U, 78905152U, 79033664U,
4946624U, 4955584U, 4964032U, 4972096U, 4980032U, 4988864U, 4997056U, 79166656U, 79297472U, 79429568U, 79560512U, 79690816U, 79822784U, 79953472U,
5004992U, 5012288U, 5020096U, 5029312U, 5037632U, 5045696U, 5052224U, 80084672U, 80214208U, 80346944U, 80477632U, 80608576U, 80740288U, 80870848U,
5062592U, 5070784U, 5078848U, 5086784U, 5095232U, 5100736U, 5111488U, 81002048U, 81133504U, 81264448U, 81395648U, 81525952U, 81657536U, 81786304U,
5119936U, 5127104U, 5136064U, 5143616U, 5151424U, 5160256U, 5168704U, 81919808U, 82050112U, 82181312U, 82311616U, 82443968U, 82573376U, 82705984U,
5175232U, 5185472U, 5192384U, 5199296U, 5209664U, 5218112U, 5225536U, 82835776U, 82967744U, 83096768U, 83230528U, 83359552U, 83491264U, 83622464U,
5233472U, 5242816U, 5250496U, 5258944U, 5267264U, 5274944U, 5283776U, 83753536U, 83886016U, 84015296U, 84147776U, 84277184U, 84409792U, 84540608U,
5290048U, 5300032U, 5308096U, 5316544U, 5323328U, 5331904U, 5340736U, 84672064U, 84803008U, 84934336U, 85065152U, 85193792U, 85326784U, 85458496U,
5349056U, 5356864U, 5365312U, 5372096U, 5381696U, 5390272U, 5398336U, 85589312U, 85721024U, 85851968U, 85982656U, 86112448U, 86244416U, 86370112U,
5405888U, 5413696U, 5422784U, 5430976U, 5439424U, 5446976U, 5455808U, 86506688U, 86637632U, 86769344U, 86900672U, 87031744U, 87162304U, 87293632U,
5463616U, 5471168U, 5480128U, 5488064U, 5494592U, 5504704U, 5513152U, 87424576U, 87555392U, 87687104U, 87816896U, 87947968U, 88079168U, 88211264U,
5521216U, 5529536U, 5536576U, 5544256U, 5554112U, 5559616U, 5570368U, 88341824U, 88473152U, 88603712U, 88735424U, 88862912U, 88996672U, 89128384U,
5577664U, 5586752U, 5594944U, 5603008U, 5611456U, 5619392U, 5627584U, 89259712U, 89390272U, 89521984U, 89652544U, 89783872U, 89914816U, 90045376U,
5634368U, 5643328U, 5651264U, 5659328U, 5667008U, 5675584U, 5684416U, 90177088U, 90307904U, 90438848U, 90569152U, 90700096U, 90832832U, 90963776U,
5692864U, 5701568U, 5709632U, 5717056U, 5725376U, 5734336U, 5740096U, 91093696U, 91223744U, 91356992U, 91486784U, 91618496U, 91749824U, 91880384U,
5750336U, 5758912U, 5766848U, 5775296U, 5782976U, 5790784U, 5799616U, 92012224U, 92143552U, 92273344U, 92405696U, 92536768U, 92666432U, 92798912U,
5807936U, 5815232U, 5823808U, 5832256U, 5840192U, 5848768U, 5856832U, 92926016U, 93060544U, 93192128U, 93322816U, 93453632U, 93583936U, 93715136U,
5864896U, 5873344U, 5879872U, 5888576U, 5897792U, 5905216U, 5914432U, 93845056U, 93977792U, 94109504U, 94240448U, 94371776U, 94501184U, 94632896U,
5920448U, 5930944U, 5938624U, 5947328U, 5955392U, 5963456U, 5971648U, 94764224U, 94895552U, 95023424U, 95158208U, 95287744U, 95420224U, 95550016U,
5979328U, 5988032U, 5995712U, 6003904U, 6012736U, 6021056U, 6029248U, 95681216U, 95811904U, 95943872U, 96075328U, 96203584U, 96337856U, 96468544U,
6037184U, 6045632U, 6053312U, 6061376U, 6070208U, 6077504U, 6086464U, 96599744U, 96731072U, 96860992U, 96992576U, 97124288U, 97254848U, 97385536U,
6094784U, 6101696U, 6110912U, 6118592U, 6127168U, 6135616U, 6143296U, 97517248U, 97647808U, 97779392U, 97910464U, 98041408U, 98172608U, 98303168U,
6150208U, 6158912U, 6168128U, 6175808U, 6182464U, 6192832U, 6201152U, 98434496U, 98565568U, 98696768U, 98827328U, 98958784U, 99089728U, 99220928U,
6209344U, 6217664U, 6224576U, 6233408U, 6241472U, 6249664U, 6258496U, 99352384U, 99482816U, 99614272U, 99745472U, 99876416U, 100007104U,
6266816U, 6275008U, 6281152U, 6291136U, 6299456U, 6306752U, 6314816U, 100138048U, 100267072U, 100401088U, 100529984U, 100662592U, 100791872U,
6323776U, 6332096U, 6339392U, 6348224U, 6356288U, 6364096U, 6373184U, 100925248U, 101056064U, 101187392U, 101317952U, 101449408U, 101580608U,
6381376U, 6389696U, 6397504U, 6404416U, 6413632U, 6421952U, 6430016U, 101711296U, 101841728U, 101973824U, 102104896U, 102235712U, 102366016U,
6437824U, 6446912U, 6454592U, 6463168U, 6471616U, 6478144U, 6487232U, 102498112U, 102628672U, 102760384U, 102890432U, 103021888U, 103153472U,
6496192U, 6504128U, 6511936U, 6520256U, 6528832U, 6536896U, 6544576U, 103284032U, 103415744U, 103545152U, 103677248U, 103808576U, 103939648U,
6553408U, 6561472U, 6569792U, 6577216U, 6586304U, 6592448U, 6601024U, 104070976U, 104201792U, 104332736U, 104462528U, 104594752U, 104725952U,
6610624U, 6619072U, 6627136U, 6634816U, 6643264U, 6650816U, 6659776U, 104854592U, 104988608U, 105118912U, 105247808U, 105381184U, 105511232U,
6667712U, 6675904U, 6682688U, 6691904U, 6700864U, 6709184U, 6717376U, 105643072U, 105774784U, 105903296U, 106037056U, 106167872U, 106298944U,
6724544U, 6733504U, 6741824U, 6749888U, 6756032U, 6766528U, 6773056U, 106429504U, 106561472U, 106691392U, 106822592U, 106954304U, 107085376U,
6782912U, 6790976U, 6798016U, 6807488U, 6815168U, 6823744U, 6832064U, 107216576U, 107346368U, 107478464U, 107609792U, 107739712U, 107872192U,
6840128U, 6847552U, 6855872U, 6864064U, 6872128U, 6880576U, 6889408U, 108003136U, 108131392U, 108265408U, 108396224U, 108527168U, 108657344U,
6897472U, 6905792U, 6913472U, 6920896U, 6930368U, 6938432U, 6946624U, 108789568U, 108920384U, 109049792U, 109182272U, 109312576U, 109444928U,
6953536U, 6963136U, 6971072U, 6979136U, 6986944U, 6995392U, 7003712U, 109572928U, 109706944U, 109837888U, 109969088U, 110099648U, 110230976U,
7012288U, 7019072U, 7028416U, 7036352U, 7044416U, 7051712U, 7060672U, 110362432U, 110492992U, 110624704U, 110755264U, 110886208U, 111017408U,
7069376U, 7077568U, 7085504U, 7092544U, 7102016U, 7110592U, 7118656U, 111148864U, 111279296U, 111410752U, 111541952U, 111673024U, 111803456U,
7126208U, 7135168U, 7143104U, 7150912U, 7159744U, 7167808U, 7175744U, 111933632U, 112066496U, 112196416U, 112328512U, 112457792U, 112590784U,
7184192U, 7191232U, 7200448U, 7207744U, 7216576U, 7224128U, 7233472U, 112715968U, 112852672U, 112983616U, 113114944U, 113244224U, 113376448U,
7241536U, 7249856U, 7256512U, 7264832U, 7274048U, 7282112U, 7290176U, 113505472U, 113639104U, 113770304U, 113901376U, 114031552U, 114163264U,
7298752U, 7306688U, 7315136U, 7322816U, 7331392U, 7339456U, 7347776U, 114294592U, 114425536U, 114556864U, 114687424U, 114818624U, 114948544U,
7356224U, 7364288U, 7371712U, 7380928U, 7387456U, 7396544U, 7404352U, 115080512U, 115212224U, 115343296U, 115473472U, 115605184U, 115736128U,
7413568U, 7421632U, 7429696U, 7436864U, 7446464U, 7454144U, 7461952U, 115867072U, 115997248U, 116128576U, 116260288U, 116391488U, 116522944U,
7470784U, 7478336U, 7487296U, 7495616U, 7503424U, 7511872U, 7520192U, 116652992U, 116784704U, 116915648U, 117046208U, 117178304U, 117308608U,
7527616U, 7536448U, 7544512U, 7551424U, 7560128U, 7568576U, 7577536U, 117440192U, 117569728U, 117701824U, 117833024U, 117964096U, 118094656U,
7583552U, 7592512U, 7600448U, 7610048U, 7618496U, 7626176U, 7634752U, 118225984U, 118357312U, 118489024U, 118617536U, 118749632U, 118882112U,
7642816U, 7651264U, 7659328U, 7667008U, 7675456U, 7683136U, 7691584U, 119012416U, 119144384U, 119275328U, 119406016U, 119537344U, 119668672U,
7700416U, 7707584U, 7716416U, 7724224U, 7733056U, 7740608U, 7749184U, 119798464U, 119928896U, 120061376U, 120192832U, 120321728U, 120454336U,
7756096U, 7765952U, 7774016U, 7781824U, 7790528U, 7798592U, 7805888U, 120584512U, 120716608U, 120848192U, 120979136U, 121109056U, 121241408U,
7814336U, 7822784U, 7831232U, 7839296U, 7847104U, 7855552U, 7863616U, 121372352U, 121502912U, 121634752U, 121764416U, 121895744U, 122027072U,
7872448U, 7880128U, 7888576U, 7896256U, 7905088U, 7912768U, 7920448U, 122157632U, 122289088U, 122421184U, 122550592U, 122682944U, 122813888U,
7928768U, 7937344U, 7945792U, 7953728U, 7959488U, 7970752U, 7978816U, 122945344U, 123075776U, 123207488U, 123338048U, 123468736U, 123600704U,
7987136U, 7994816U, 8003392U, 8011712U, 8019904U, 8027456U, 8035264U, 123731264U, 123861952U, 123993664U, 124124608U, 124256192U, 124386368U,
8044352U, 8052544U, 8060224U, 8069056U, 8076736U, 8084672U, 8093504U, 124518208U, 124649024U, 124778048U, 124911296U, 125041088U, 125173696U,
8101312U, 8110016U, 8117696U, 8125888U, 8134592U, 8142016U, 8149952U, 125303744U, 125432896U, 125566912U, 125696576U, 125829056U, 125958592U,
8159168U, 8166976U, 8175296U, 8183488U, 8191808U, 8199616U, 8207296U, 126090304U, 126221248U, 126352832U, 126483776U, 126615232U, 126746432U,
8216128U, 8224576U, 8232256U, 8241088U, 8248256U, 8257472U, 8264128U, 126876608U, 127008704U, 127139392U, 127270336U, 127401152U, 127532224U,
8273728U, 8281792U, 8290112U, 8297152U, 8305216U, 8314816U, 8322752U, 127663552U, 127794752U, 127925696U, 128055232U, 128188096U, 128319424U,
8330944U, 8339392U, 8347072U, 8355392U, 8363968U, 8371904U, 8379328U, 128449856U, 128581312U, 128712256U, 128843584U, 128973632U, 129103808U,
8388544U, 8394944U, 8404544U, 8412736U, 8421184U, 8429504U, 8437696U, 129236288U, 129365696U, 129498944U, 129629888U, 129760832U, 129892288U,
8445376U, 8452544U, 8460736U, 8470208U, 8478016U, 8486848U, 8494144U, 130023104U, 130154048U, 130283968U, 130416448U, 130547008U, 130678336U,
8503232U, 8511296U, 8519488U, 8527424U, 8534464U, 8543936U, 8552384U, 130807616U, 130939456U, 131071552U, 131202112U, 131331776U, 131464384U,
8558912U, 8568128U, 8575936U, 8584256U, 8593216U, 8601536U, 8608832U, 131594048U, 131727296U, 131858368U, 131987392U, 132120256U, 132250816U,
8616896U, 8625728U, 8634176U, 8641856U, 8649664U, 8658112U, 8666176U, 132382528U, 132513728U, 132644672U, 132774976U, 132905792U, 133038016U,
8674112U, 8682944U, 8691136U, 8699456U, 8707648U, 8716096U, 8724416U, 133168832U, 133299392U, 133429312U, 133562048U, 133692992U, 133823296U,
8732608U, 8740672U, 8748352U, 8756032U, 8764864U, 8773568U, 8781376U, 133954624U, 134086336U, 134217152U, 134348608U, 134479808U, 134607296U,
8789824U, 8796992U, 8806208U, 8814272U, 8822336U, 8830912U, 8838848U, 134741056U, 134872384U, 135002944U, 135134144U, 135265472U, 135396544U,
8847296U, 8854336U, 8863552U, 8871488U, 8879296U, 8887616U, 8894528U, 135527872U, 135659072U, 135787712U, 135921472U, 136052416U, 136182848U,
8904512U, 8911424U, 8920768U, 8928704U, 8936128U, 8944576U, 8953664U, 136313792U, 136444864U, 136576448U, 136707904U, 136837952U, 136970048U,
8960576U, 8970176U, 8977984U, 8986304U, 8994112U, 9002432U, 9011008U, 137099584U, 137232064U, 137363392U, 137494208U, 137625536U, 137755712U,
9018176U, 9026624U, 9035584U, 9043904U, 9052096U, 9059264U, 9068096U, 137887424U, 138018368U, 138149824U, 138280256U, 138411584U, 138539584U,
9075904U, 9084224U, 9092288U, 9100352U, 9108928U, 9116992U, 9125824U, 138672832U, 138804928U, 138936128U, 139066688U, 139196864U, 139328704U,
9133504U, 9141824U, 9150272U, 9157952U, 9164608U, 9174848U, 9182912U, 139460032U, 139590208U, 139721024U, 139852864U, 139984576U, 140115776U,
9190976U, 9199552U, 9205312U, 9215936U, 9222592U, 9232192U, 9240512U, 140245696U, 140376512U, 140508352U, 140640064U, 140769856U, 140902336U,
9248704U, 9256256U, 9264832U, 9272896U, 9281344U, 9288896U, 9297088U, 141032768U, 141162688U, 141294016U, 141426496U, 141556544U, 141687488U,
9305536U, 9313984U, 9322304U, 9329728U, 9337792U, 9346112U, 9355072U, 141819584U, 141949888U, 142080448U, 142212544U, 142342336U, 142474432U,
9363136U, 9371072U, 9378752U, 9387712U, 9395648U, 9404224U, 9411008U, 142606144U, 142736192U, 142868288U, 142997824U, 143129408U, 143258944U,
9420608U, 9428416U, 9436864U, 9445312U, 9453376U, 9460928U, 9468736U, 143392448U, 143523136U, 143653696U, 143785024U, 143916992U, 144045632U,
9477824U, 9485248U, 9493696U, 9502144U, 9509056U, 9518528U, 9527104U, 144177856U, 144309184U, 144440768U, 144570688U, 144701888U, 144832448U,
9535424U, 9543616U, 9551296U, 9559744U, 9568192U, 9576256U, 9584576U, 144965056U, 145096384U, 145227584U, 145358656U, 145489856U, 145620928U,
9591872U, 9600704U, 9608384U, 9615808U, 9624512U, 9633472U, 9641536U, 145751488U, 145883072U, 146011456U, 146144704U, 146275264U, 146407232U,
9649856U, 9658048U, 9665728U, 9674432U, 9682496U, 9691072U, 9699136U, 146538176U, 146668736U, 146800448U, 146931392U, 147062336U, 147193664U,
9707072U, 9715136U, 9722176U, 9732032U, 9740096U, 9747904U, 9756352U, 147324224U, 147455936U, 147586624U, 147717056U, 147848768U, 147979456U,
9764288U, 9771584U, 9780544U, 9789376U, 9796928U, 9804224U, 9813952U, 148110784U, 148242368U, 148373312U, 148503232U, 148635584U, 148766144U,
9822016U, 9829696U, 9838016U, 9845824U, 9852992U, 9863104U, 9870656U, 148897088U, 149028416U, 149159488U, 149290688U, 149420224U, 149551552U,
9878464U, 9887552U, 9895744U, 9903808U, 9912128U, 9920192U, 9927616U, 149683136U, 149814976U, 149943616U, 150076352U, 150208064U, 150338624U,
9936064U, 9944768U, 9952576U, 9960128U, 9969472U, 9977152U, 9985216U, 150470464U, 150600256U, 150732224U, 150862784U, 150993088U, 151125952U,
9994048U, 10001216U, 10007744U, 10018496U, 10026944U, 10035136U, 10042432U, 151254976U, 151388096U, 151519168U, 151649728U, 151778752U, 151911104U,
10051264U, 10059584U, 10067648U, 10075712U, 10083904U, 10091456U, 10100672U, 152042944U, 152174144U, 152304704U, 152435648U, 152567488U, 152698816U,
10108864U, 10116928U, 10124864U, 10133056U, 10140736U, 10149824U, 10156736U, 152828992U, 152960576U, 153091648U, 153222976U, 153353792U, 153484096U,
10165952U, 10173376U, 10182208U, 10190528U, 10198336U, 10206272U, 10213696U, 153616192U, 153747008U, 153878336U, 154008256U, 154139968U, 154270912U,
10223296U, 10231744U, 10238656U, 10247488U, 10256192U, 10263872U, 10272448U, 154402624U, 154533824U, 154663616U, 154795712U, 154926272U, 155057984U,
10280896U, 10288448U, 10296512U, 10305088U, 10313536U, 10321088U, 10330048U, 155188928U, 155319872U, 155450816U, 155580608U, 155712064U, 155843392U,
10337984U, 10346176U, 10354112U, 10362304U, 10369088U, 10377152U, 10386752U, 155971136U, 156106688U, 156237376U, 156367424U, 156499264U, 156630976U,
10394816U, 10403648U, 10411712U, 10418624U, 10427968U, 10436032U, 10444736U, 156761536U, 156892352U, 157024064U, 157155008U, 157284416U, 157415872U,
10452928U, 10459712U, 10468672U, 10476608U, 10484416U, 10491328U, 10501952U, 157545536U, 157677248U, 157810496U, 157938112U, 158071744U, 158203328U,
10509376U, 10517824U, 10526528U, 10534336U, 10542656U, 10549696U, 10559168U, 158334656U, 158464832U, 158596288U, 158727616U, 158858048U, 158988992U,
10566592U, 10575808U, 10583488U, 10590656U, 10599488U, 10607936U, 10616768U, 159121216U, 159252416U, 159381568U, 159513152U, 159645632U, 159776192U,
10624832U, 10630336U, 10640576U, 10649536U, 10655168U, 10665152U, 10674112U, 159906496U, 160038464U, 160169536U, 160300352U, 160430656U, 160563008U,
10682176U, 10690496U, 10698176U, 10705216U, 10715072U, 10722752U, 10731328U, 160693952U, 160822208U, 160956352U, 161086784U, 161217344U, 161349184U,
10739264U, 10746688U, 10754752U, 10761664U, 10770752U, 10779712U, 10787776U, 161480512U, 161611456U, 161742272U, 161873216U, 162002752U, 162135872U,
10796608U, 10803392U, 10812352U, 10821056U, 10828736U, 10837952U, 10846144U, 162266432U, 162397888U, 162529216U, 162660032U, 162790976U, 162922048U,
10853824U, 10861376U, 10869952U, 10877248U, 10887104U, 10895296U, 10903232U, 163052096U, 163184576U, 163314752U, 163446592U, 163577408U, 163707968U,
10910912U, 10918976U, 10927936U, 10935872U, 10944448U, 10952384U, 10960832U, 163839296U, 163969984U, 164100928U, 164233024U, 164364224U, 164494912U,
10968512U, 10977088U, 10985024U, 10992832U, 11000896U, 11009984U, 11018048U, 164625856U, 164756672U, 164887616U, 165019072U, 165150016U, 165280064U,
11026112U, 11034304U, 11042624U, 11050432U, 11058368U, 11064512U, 11075392U, 165412672U, 165543104U, 165674944U, 165805888U, 165936832U, 166067648U,
11083712U, 11091776U, 11099584U, 11107904U, 11115968U, 11124416U, 11131712U, 166198336U, 166330048U, 166461248U, 166591552U, 166722496U, 166854208U,
11141056U, 11148608U, 11157184U, 11165248U, 11173312U, 11180992U, 11189056U, 166985408U, 167116736U, 167246656U, 167378368U, 167508416U, 167641024U,
11197376U, 11206592U, 11214656U, 11222336U, 11230784U, 11238464U, 11246528U, 167771584U, 167903168U, 168034112U, 168164032U, 168295744U, 168427456U,
11254976U, 11263552U, 11271872U, 11279552U, 11288512U, 11296576U, 11304256U, 168557632U, 168688448U, 168819136U, 168951616U, 169082176U, 169213504U,
11312192U, 11320768U, 11329216U, 11336384U, 11345216U, 11352512U, 11362112U, 169344832U, 169475648U, 169605952U, 169738048U, 169866304U, 169999552U,
11369408U, 11378624U, 11386688U, 11394496U, 11402816U, 11411264U, 11418688U, 170131264U, 170262464U, 170393536U, 170524352U, 170655424U, 170782016U,
11427776U, 11435584U, 11444032U, 11452096U, 11459648U, 11467072U, 11476928U, 170917696U, 171048896U, 171179072U, 171310784U, 171439936U, 171573184U,
11484992U, 11493184U, 11500352U, 11509312U, 11517248U, 11524928U, 11534144U, 171702976U, 171835072U, 171966272U, 172097216U, 172228288U, 172359232U,
11542208U, 11550272U, 11556416U, 11566784U, 11574208U, 11581376U, 11589568U, 172489664U, 172621376U, 172747712U, 172883264U, 173014208U, 173144512U,
11599552U, 11607104U, 11616064U, 11623616U, 11632576U, 11639872U, 11648704U, 173275072U, 173407424U, 173539136U, 173669696U, 173800768U, 173931712U,
11657024U, 11664704U, 11672896U, 11681216U, 11689792U, 11697856U, 11705536U, 174063424U, 174193472U, 174325696U, 174455744U, 174586816U, 174718912U,
11714368U, 11722688U, 11730496U, 11737408U, 11745728U, 11754304U, 11763008U, 174849728U, 174977728U, 175109696U, 175242688U, 175374272U, 175504832U,
11770816U, 11779648U, 11788096U, 11795776U, 11804608U, 11812544U, 11820992U, 175636288U, 175765696U, 175898432U, 176028992U, 176159936U, 176291264U,
11829184U, 11837248U, 11844928U, 11852096U, 11860928U, 11869888U, 11878336U, 176422592U, 176552512U, 176684864U, 176815424U, 176946496U, 177076544U,
11886272U, 11894336U, 11902144U, 11910848U, 11919296U, 11925952U, 11934784U, 177209152U, 177340096U, 177470528U, 177600704U, 177731648U, 177864256U,
11943616U, 11951552U, 11960128U, 11968192U, 11976512U, 11983168U, 11992768U, 177994816U, 178126528U, 178257472U, 178387648U, 178518464U, 178650176U,
12000832U, 12008896U, 12016832U, 12025408U, 12033856U, 12042176U, 12049984U, 178781888U, 178912064U, 179044288U, 179174848U, 179305024U, 179436736U,
12058048U, 12066112U, 12073792U, 12082624U, 12091328U, 12098752U, 12106816U, 179568448U, 179698496U, 179830208U, 179960512U, 180092608U, 180223808U,
12115904U, 12124096U, 12131776U, 12140224U, 12148672U, 12156736U, 12164032U, 180354752U, 180485696U, 180617152U, 180748096U, 180877504U, 181009984U,
12173248U, 12181184U, 12186176U, 12197824U, 12205888U, 12213952U, 12218944U, 181139264U, 181272512U, 181402688U, 181532608U, 181663168U, 181795136U,
12230336U, 12238784U, 12246592U, 12254272U, 12262336U, 12269888U, 12279104U, 181926592U, 182057536U, 182190016U, 182320192U, 182451904U, 182582336U,
12287936U, 12295744U, 12304064U, 12312512U, 12319936U, 12328768U, 12337088U, 182713792U, 182843072U, 182976064U, 183107264U, 183237056U, 183368384U,
12344896U, 12352832U, 12361408U, 12368704U, 12377152U, 12384832U, 12394432U, 183494848U, 183631424U, 183762752U, 183893824U, 184024768U, 184154816U,
12402496U, 12409024U, 12417728U, 12426688U, 12433216U, 12443584U, 12450752U, 184286656U, 184417984U, 184548928U, 184680128U, 184810816U, 184941248U,
12459968U, 12468032U, 12475712U, 12484544U, 12492608U, 12500416U, 12508352U, 185072704U, 185203904U, 185335616U, 185465408U, 185596352U, 185727296U,
12517184U, 12525376U, 12532288U, 12541888U, 12549568U, 12556864U, 12565568U, 185859904U, 185989696U, 186121664U, 186252992U, 186383552U, 186514112U,
12574528U, 12582208U, 12590528U, 12598592U, 12607424U, 12615488U, 12623552U, 186645952U, 186777152U, 186907328U, 187037504U, 187170112U, 187301824U,
12631744U, 12638656U, 12647744U, 12656576U, 12664768U, 12672832U, 12680896U, 187429184U, 187562048U, 187693504U, 187825472U, 187957184U, 188087104U,
12688576U, 12697408U, 12704192U, 12713408U, 12721216U, 12729664U, 12738496U, 188218304U, 188349376U, 188481344U, 188609728U, 188743616U, 188874304U,
12745792U, 12754496U, 12762688U, 12769472U, 12779456U, 12787648U, 12795712U, 189005248U, 189136448U, 189265088U, 189396544U, 189528128U, 189660992U,
12804032U, 12812224U, 12819008U, 12828352U, 12836672U, 12844736U, 12851648U, 189791936U, 189923264U, 190054208U, 190182848U, 190315072U, 190447424U,
12859456U, 12868672U, 12877504U, 12885568U, 12892864U, 12902336U, 12909376U, 190577984U, 190709312U, 190840768U, 190971328U, 191102656U, 191233472U,
12918208U, 12926656U, 12934976U, 12942784U, 12951104U, 12959552U, 12967744U, 191364032U, 191495872U, 191626816U, 191758016U, 191888192U, 192020288U,
12976064U, 12984256U, 12991936U, 12999488U, 13007936U, 13016768U, 13021504U, 192148928U, 192282176U, 192413504U, 192542528U, 192674752U, 192805952U,
13033024U, 13041472U, 13049408U, 13057472U, 13065664U, 13072064U, 13081408U, 192937792U, 193068608U, 193198912U, 193330496U, 193462208U, 193592384U,
13089344U, 13098688U, 13107008U, 13115072U, 13122752U, 13130944U, 13139648U, 193723456U, 193854272U, 193985984U, 194116672U, 194247232U, 194379712U,
13147712U, 13155776U, 13162432U, 13172672U, 13180864U, 13188928U, 13196992U, 194508352U, 194641856U, 194772544U, 194900672U, 195035072U, 195166016U,
13203392U, 13213504U, 13219264U, 13228736U, 13236928U, 13244992U, 13253056U, 195296704U, 195428032U, 195558592U, 195690304U, 195818176U, 195952576U,
13262528U, 13269952U, 13278784U, 13285952U, 13295552U, 13303616U, 13311808U, 196083392U, 196214336U, 196345792U, 196476736U, 196607552U, 196739008U,
13319744U, 13328192U, 13336256U, 13344704U, 13352384U, 13360576U, 13369024U, 196869952U, 197000768U, 197130688U, 197262784U, 197394368U, 197523904U,
13377344U, 13385408U, 13393216U, 13401664U, 13410112U, 13418176U, 13426496U, 197656384U, 197787584U, 197916608U, 198049472U, 198180544U, 198310208U,
13434688U, 13442368U, 13451072U, 13459136U, 13466944U, 13475648U, 13482944U, 198442432U, 198573632U, 198705088U, 198834368U, 198967232U, 199097792U,
13491904U, 13500352U, 13508288U, 13516736U, 13524416U, 13532224U, 13541312U, 199228352U, 199360192U, 199491392U, 199621696U, 199751744U, 199883968U,
13549504U, 13556288U, 13564736U, 13573184U, 13581376U, 13587008U, 13598656U, 200014016U, 200146624U, 200276672U, 200408128U, 200540096U, 200671168U,
13605952U, 13612864U, 13622464U, 13631168U, 13639616U, 13647808U, 13655104U, 200801984U, 200933312U, 201062464U, 201194944U, 201326144U, 201457472U,
13663424U, 13671872U, 13680064U, 13688768U, 13696576U, 13705024U, 13712576U, 201588544U, 201719744U, 201850816U, 201981632U, 202111552U, 202244032U,
13721536U, 13729216U, 13737664U, 13746112U, 13753024U, 13759552U, 13770304U, 202374464U, 202505152U, 202636352U, 202767808U, 202898368U, 203030336U,
13777856U, 13786688U, 13793984U, 13802176U, 13811264U, 13819328U, 13827904U, 203159872U, 203292608U, 203423296U, 203553472U, 203685824U, 203816896U,
13835456U, 13844416U, 13851584U, 13860544U, 13868992U, 13877056U, 13884608U, 203947712U, 204078272U, 204208192U, 204341056U, 204472256U, 204603328U,
13893184U, 13901248U, 13909696U, 13917632U, 13925056U, 13934528U, 13942336U, 204733888U, 204864448U, 204996544U, 205125568U, 205258304U, 205388864U,
13950784U, 13959104U, 13966912U, 13975232U, 13982656U, 13991872U, 13999936U, 205517632U, 205650112U, 205782208U, 205913536U, 206044736U, 206176192U,
14007872U, 14016064U, 14024512U, 14032064U, 14040896U, 14049088U, 14057408U, 206307008U, 206434496U, 206569024U, 206700224U, 206831168U, 206961856U,
14065088U, 14072896U, 14081344U, 14089664U, 14097856U, 14106304U, 14114752U, 207093056U, 207223616U, 207355328U, 207486784U, 207616832U, 207749056U,
14122688U, 14130752U, 14138816U, 14147008U, 14155072U, 14163904U, 14170432U, 207879104U, 208010048U, 208141888U, 208273216U, 208404032U, 208534336U,
14180288U, 14187328U, 14196032U, 14204864U, 14212672U, 14220736U, 14229056U, 208666048U, 208796864U, 208927424U, 209059264U, 209189824U, 209321792U,
14237504U, 14245568U, 14253632U, 14261824U, 14269888U, 14278592U, 14286656U, 209451584U, 209582656U, 209715136U, 209845568U, 209976896U, 210106432U,
14293696U, 14302784U, 14309696U, 14317504U, 14326336U, 14335936U, 14343232U, 210239296U, 210370112U, 210501568U, 210630976U, 210763712U, 210894272U,
14352064U, 14359232U, 14368064U, 14376512U, 14384576U, 14393024U, 14401472U, 211024832U, 211156672U, 211287616U, 211418176U, 211549376U, 211679296U,
14409536U, 14416832U, 14424512U, 14433856U, 14440768U, 14449984U, 14458816U, 211812032U, 211942592U, 212074432U, 212204864U, 212334016U, 212467648U,
14465728U, 14474816U, 14482112U, 14491328U, 14499392U, 14506816U, 14516032U, 212597824U, 212727616U, 212860352U, 212991424U, 213120832U, 213253952U,
14524352U, 14531392U, 14540224U, 14547392U, 14556992U, 14565184U, 14573248U, 213385024U, 213515584U, 213645632U, 213777728U, 213909184U, 214040128U,
14580928U, 14588864U, 14596928U, 14606272U, 14613824U, 14622656U, 14630464U, 214170688U, 214302656U, 214433728U, 214564544U, 214695232U, 214826048U,
14638912U, 14646976U, 14655296U, 14661952U, 14671808U, 14679872U, 14687936U, 214956992U, 215089088U, 215219776U, 215350592U, 215482304U, 215613248U,
14696384U, 14704576U, 14710336U, 14720192U, 14729152U, 14736448U, 14745152U, 215743552U, 215874752U, 216005312U, 216137024U, 216267328U, 216399296U,
14752448U, 14761792U, 14769856U, 14777024U, 14785984U, 14792384U, 14802752U, 216530752U, 216661696U, 216790592U, 216923968U, 217054528U, 217183168U,
14810816U, 14819264U, 14827328U, 14835136U, 14843072U, 14851264U, 14860096U, 217316672U, 217448128U, 217579072U, 217709504U, 217838912U, 217972672U,
14867648U, 14876096U, 14884544U, 14892736U, 14900672U, 14907968U, 14917312U, 218102848U, 218233024U, 218364736U, 218496832U, 218627776U, 218759104U,
14924864U, 14933824U, 14939968U, 14950336U, 14957632U, 14966464U, 14974912U, 218888896U, 219021248U, 219151936U, 219281728U, 219413056U, 219545024U,
14982592U, 14991296U, 14999104U, 15006272U, 15015232U, 15023936U, 15031616U, 219675968U, 219807296U, 219938624U, 220069312U, 220200128U, 220331456U,
15040448U, 15047488U, 15055552U, 15063616U, 15073216U, 15079744U, 15088064U, 220461632U, 220592704U, 220725184U, 220855744U, 220987072U, 221117888U,
15097664U, 15105344U, 15113792U, 15122368U, 15130048U, 15137728U, 15146176U, 221249216U, 221378368U, 221510336U, 221642048U, 221772736U, 221904832U,
15154112U, 15162688U, 15171392U, 15179456U, 15187264U, 15194176U, 15204032U, 222031808U, 222166976U, 222297536U, 222428992U, 222559936U, 222690368U,
15212224U, 15220544U, 15227456U, 15237056U, 15245248U, 15253184U, 15261632U, 222820672U, 222953152U, 223083968U, 223213376U, 223345984U, 223476928U,
15269824U, 15277376U, 15285824U, 15293888U, 15301568U, 15310784U, 15318848U, 223608512U, 223738688U, 223869376U, 224001472U, 224132672U, 224262848U,
15325504U, 15335104U, 15343168U, 15350848U, 15359936U, 15367232U, 15373376U, 224394944U, 224524864U, 224657344U, 224788288U, 224919488U, 225050432U,
15384256U, 15392576U, 15400384U, 15408832U, 15417152U, 15424832U, 15433024U, 225181504U, 225312704U, 225443776U, 225574592U, 225704768U, 225834176U,
15441344U, 15449152U, 15457088U, 15466432U, 15474112U, 15482816U, 15488576U, 225966784U, 226097216U, 226229824U, 226360384U, 226491712U, 226623424U,
15499072U, 15505856U, 15514816U, 15523264U, 15531584U, 15540032U, 15547328U, 226754368U, 226885312U, 227015104U, 227147456U, 227278528U, 227409472U,
15553984U, 15564608U, 15571904U, 15579968U, 15589312U, 15597376U, 15605696U, 227539904U, 227669696U, 227802944U, 227932352U, 228065216U, 228196288U,
15612992U, 15621824U, 15630016U, 15638464U, 15646144U, 15654592U, 15662912U, 228326464U, 228457792U, 228588736U, 228720064U, 228850112U, 228981056U,
15671104U, 15677248U, 15686848U, 15693376U, 15701696U, 15712064U, 15720256U, 229113152U, 229243328U, 229375936U, 229505344U, 229636928U, 229769152U,
15728576U, 15736384U, 15744704U, 15752512U, 15761344U, 15769024U, 15777728U, 229894976U, 230030272U, 230162368U, 230292416U, 230424512U, 230553152U,
15785152U, 15793984U, 15802048U, 15809984U, 15817024U, 15825856U, 15834944U, 230684864U, 230816704U, 230948416U, 231079616U, 231210944U, 231342016U,
15843008U, 15849664U, 15859136U, 15866432U, 15876032U, 15884096U, 15892288U, 231472448U, 231603776U, 231733952U, 231866176U, 231996736U, 232127296U,
15900608U, 15908416U, 15916864U, 15924928U, 15930176U, 15941056U, 15949504U, 232259392U, 232388672U, 232521664U, 232652608U, 232782272U, 232914496U,
15957824U, 15965632U, 15973952U, 15982528U, 15990592U, 15998272U, 16006976U, 233043904U, 233175616U, 233306816U, 233438528U, 233569984U, 233699776U,
16012736U, 16023104U, 16031296U, 16039616U, 16048064U, 16055744U, 16064192U, 233830592U, 233962688U, 234092224U, 234221888U, 234353984U, 234485312U,
16071488U, 16080832U, 16088768U, 16097216U, 16104896U, 16112704U, 16121792U, 234618304U, 234749888U, 234880832U, 235011776U, 235142464U, 235274048U,
16129856U, 16138048U, 16146112U, 16154176U, 16162624U, 16170688U, 16177856U, 235403456U, 235535936U, 235667392U, 235797568U, 235928768U, 236057152U,
16186816U, 16195136U, 16202176U, 16211648U, 16220096U, 16228288U, 16235584U, 236190272U, 236322752U, 236453312U, 236583616U, 236715712U, 236846528U,
16244672U, 16252864U, 16260544U, 16269248U, 16277056U, 16285504U, 16291648U, 236976448U, 237108544U, 237239104U, 237371072U, 237501632U, 237630784U,
16301632U, 16309312U, 16318144U, 16326208U, 16333888U, 16342336U, 16351168U, 237764416U, 237895232U, 238026688U, 238157632U, 238286912U, 238419392U,
16359232U, 16367552U, 16375616U, 16383296U, 16391744U, 16398016U, 16407616U, 238548032U, 238681024U, 238812608U, 238941632U, 239075008U, 239206336U,
16415936U, 16424896U, 16432448U, 16440896U, 16449088U, 16457024U, 16465472U, 239335232U, 239466944U, 239599168U, 239730496U, 239861312U, 239992384U,
16474048U, 16481216U, 16490048U, 16498624U, 16505792U, 16513984U, 16523072U, 240122816U, 240254656U, 240385856U, 240516928U, 240647872U, 240779072U,
16531136U, 16538944U, 16547264U, 16555328U, 16563776U, 16570816U, 16578112U, 240909632U, 241040704U, 241171904U, 241302848U, 241433408U, 241565248U,
16587712U, 16596544U, 16604992U, 16613312U, 16620608U, 16629568U, 16637888U, 241696192U, 241825984U, 241958848U, 242088256U, 242220224U, 242352064U,
16645696U, 16653632U, 16661696U, 16669888U, 16677568U, 16686272U, 16695232U, 242481856U, 242611648U, 242744896U, 242876224U, 243005632U, 243138496U,
16703168U, 16710464U, 16719424U, 16726592U, 16733888U, 16744384U, 16752448U, 243268672U, 243400384U, 243531712U, 243662656U, 243793856U, 243924544U,
16760768U, 16768448U, 16776896U, 16785344U, 16793536U, 16801216U, 16809664U, 244054592U, 244187072U, 244316608U, 244448704U, 244580032U, 244710976U,
16818112U, 16826176U, 16833472U, 16842688U, 16850752U, 16859072U, 16866368U, 244841536U, 244972864U, 245104448U, 245233984U, 245365312U, 245497792U,
16875328U, 16883392U, 16891712U, 16899776U, 16907456U, 16915264U, 16924352U, 245628736U, 245759936U, 245889856U, 246021056U, 246152512U, 246284224U,
16931776U, 16940608U, 16949056U, 16957376U, 16965056U, 16973248U, 16981696U, 246415168U, 246545344U, 246675904U, 246808384U, 246939584U, 247070144U,
16990144U, 16997056U, 17005888U, 17014208U, 17021504U, 17031104U, 17039296U, 247199552U, 247331648U, 247463872U, 247593536U, 247726016U, 247857088U,
17046976U, 17055424U, 17062592U, 17070016U, 17079488U, 17087936U, 17096512U, 247987648U, 248116928U, 248249536U, 248380736U, 248512064U, 248643008U,
17104576U, 17113024U, 17121088U, 17129408U, 17136832U, 17145664U, 17152832U, 248773312U, 248901056U, 249036608U, 249167552U, 249298624U, 249429184U,
17161792U, 17170112U, 17177792U, 17186368U, 17194304U, 17202496U, 17211328U, 249560512U, 249692096U, 249822784U, 249954112U, 250085312U, 250215488U,
17218624U, 17227712U, 17233984U, 17243584U, 17251904U, 17259712U, 17266624U, 250345792U, 250478528U, 250608704U, 250739264U, 250870976U, 251002816U,
17276608U, 17284672U, 17292224U, 17301056U, 17309632U, 17317568U, 17326016U, 251133632U, 251263552U, 251395136U, 251523904U, 251657792U, 251789248U,
17333824U, 17342272U, 17350208U, 17358784U, 17366848U, 17374912U, 17382592U, 251919424U, 252051392U, 252182464U, 252313408U, 252444224U, 252575552U,
17390656U, 17399488U, 17406784U, 17413952U, 17423936U, 17432512U, 17440448U, 252706624U, 252836032U, 252968512U, 253099712U, 253227584U, 253361728U,
17447744U, 17456704U, 17464768U, 17472064U, 17481536U, 17489344U, 17495488U, 253493056U, 253623488U, 253754432U, 253885504U, 254017216U, 254148032U,
17505728U, 17513792U, 17522368U, 17530816U, 17538112U, 17546944U, 17555264U, 254279488U, 254410432U, 254541376U, 254672576U, 254803264U, 254933824U,
17563072U, 17569856U, 17579456U, 17587904U, 17596352U, 17603776U, 17611712U, 255065792U, 255196736U, 255326528U, 255458752U, 255589952U, 255721408U,
17620672U, 17628992U, 17637184U, 17645504U, 17653568U, 17661632U, 17669824U, 255851072U, 255983296U, 256114624U, 256244416U, 256374208U, 256507712U,
17677376U, 17686208U, 17693888U, 17702336U, 17710144U, 17718208U, 17726528U, 256636096U, 256768832U, 256900544U, 257031616U, 257162176U, 257294272U,
17734336U, 17743808U, 17751872U, 17759936U, 17766592U, 17776448U, 17784512U, 257424448U, 257555776U, 257686976U, 257818432U, 257949632U, 258079552U,
17791936U, 17801152U, 17809216U, 17817152U 258211136U, 258342464U, 258473408U, 258603712U, 258734656U, 258867008U,
258996544U, 259127744U, 259260224U, 259391296U, 259522112U, 259651904U,
259784384U, 259915328U, 260045888U, 260175424U, 260308544U, 260438336U,
260570944U, 260700992U, 260832448U, 260963776U, 261092672U, 261226304U,
261356864U, 261487936U, 261619648U, 261750592U, 261879872U, 262011968U,
262143424U, 262274752U, 262404416U, 262537024U, 262667968U, 262799296U,
262928704U, 263061184U, 263191744U, 263322944U, 263454656U, 263585216U,
263716672U, 263847872U, 263978944U, 264108608U, 264241088U, 264371648U,
264501184U, 264632768U, 264764096U, 264895936U, 265024576U, 265158464U,
265287488U, 265418432U, 265550528U, 265681216U, 265813312U, 265943488U,
266075968U, 266206144U, 266337728U, 266468032U, 266600384U, 266731072U,
266862272U, 266993344U, 267124288U, 267255616U, 267386432U, 267516992U,
267648704U, 267777728U, 267910592U, 268040512U, 268172096U, 268302784U,
268435264U, 268566208U, 268696256U, 268828096U, 268959296U, 269090368U,
269221312U, 269352256U, 269482688U, 269614784U, 269745856U, 269876416U,
270007616U, 270139328U, 270270272U, 270401216U, 270531904U, 270663616U,
270791744U, 270924736U, 271056832U, 271186112U, 271317184U, 271449536U,
271580992U, 271711936U, 271843136U, 271973056U, 272105408U, 272236352U,
272367296U, 272498368U, 272629568U, 272759488U, 272891456U, 273022784U,
273153856U, 273284672U, 273415616U, 273547072U, 273677632U, 273808448U,
273937088U, 274071488U, 274200896U, 274332992U, 274463296U, 274595392U,
274726208U, 274857536U, 274988992U, 275118656U, 275250496U, 275382208U,
275513024U, 275643968U, 275775296U, 275906368U, 276037184U, 276167872U,
276297664U, 276429376U, 276560576U, 276692672U, 276822976U, 276955072U,
277085632U, 277216832U, 277347008U, 277478848U, 277609664U, 277740992U,
277868608U, 278002624U, 278134336U, 278265536U, 278395328U, 278526784U,
278657728U, 278789824U, 278921152U, 279052096U, 279182912U, 279313088U,
279443776U, 279576256U, 279706048U, 279838528U, 279969728U, 280099648U,
280230976U, 280361408U, 280493632U, 280622528U, 280755392U, 280887104U,
281018176U, 281147968U, 281278912U, 281411392U, 281542592U, 281673152U,
281803712U, 281935552U, 282066496U, 282197312U, 282329024U, 282458816U,
282590272U, 282720832U, 282853184U, 282983744U, 283115072U, 283246144U,
283377344U, 283508416U, 283639744U, 283770304U, 283901504U, 284032576U,
284163136U, 284294848U, 284426176U, 284556992U, 284687296U, 284819264U,
284950208U, 285081536U
}; };
#ifdef __cplusplus #ifdef __cplusplus

View File

@ -1,18 +1,18 @@
/* /*
This file is part of cpp-ethereum. This file is part of ethash.
cpp-ethereum is free software: you can redistribute it and/or modify ethash is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or the Free Software Foundation, either version 3 of the License, or
(at your option) any later version. (at your option) any later version.
cpp-ethereum is distributed in the hope that it will be useful, ethash is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details. GNU General Public License for more details.
You should have received a copy of the GNU General Public License You should have received a copy of the GNU General Public License
along with cpp-ethereum. If not, see <http://www.gnu.org/licenses/>. along with ethash. If not, see <http://www.gnu.org/licenses/>.
*/ */
/** @file ethash.h /** @file ethash.h
@ -26,13 +26,15 @@
#include <stddef.h> #include <stddef.h>
#include "compiler.h" #include "compiler.h"
#define REVISION 20 #define REVISION 23
#define DAGSIZE_BYTES_INIT 1073741824U // 2**30 #define DATASET_BYTES_INIT 1073741824U // 2**30
#define DAG_GROWTH 8388608U // 2**23 #define DATASET_BYTES_GROWTH 8388608U // 2**23
#define CACHE_MULTIPLIER 1024 #define CACHE_BYTES_INIT 1073741824U // 2**24
#define CACHE_BYTES_GROWTH 131072U // 2**17
#define EPOCH_LENGTH 30000U #define EPOCH_LENGTH 30000U
#define MIX_BYTES 128 #define MIX_BYTES 128
#define DAG_PARENTS 256 #define HASH_BYTES 64
#define DATASET_PARENTS 256
#define CACHE_ROUNDS 3 #define CACHE_ROUNDS 3
#define ACCESSES 64 #define ACCESSES 64
@ -60,19 +62,38 @@ static inline void ethash_params_init(ethash_params *params, const uint32_t bloc
} }
typedef struct ethash_cache { typedef struct ethash_cache {
void *mem; void *mem;
} ethash_cache; } ethash_cache;
void ethash_mkcache(ethash_cache *cache, ethash_params const *params, const uint8_t seed[32]); void ethash_mkcache(ethash_cache *cache, ethash_params const *params, const uint8_t seed[32]);
void ethash_compute_full_data(void *mem, ethash_params const *params, ethash_cache const *cache); void ethash_compute_full_data(void *mem, ethash_params const *params, ethash_cache const *cache);
void ethash_full(ethash_return_value *ret, void const *full_mem, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce); void ethash_full(ethash_return_value *ret, void const *full_mem, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce);
void ethash_light(ethash_return_value *ret, ethash_cache const *cache, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce); void ethash_light(ethash_return_value *ret, ethash_cache const *cache, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce);
void ethash_get_seedhash(uint8_t seedhash[32], const uint32_t block_number);
static inline void ethash_prep_light(void *cache, ethash_params const *params, const uint8_t seed[32]) { ethash_cache c; c.mem = cache; ethash_mkcache(&c, params, seed); } static inline void ethash_prep_light(void *cache, ethash_params const *params, const uint8_t seed[32]) {
static inline void ethash_compute_light(ethash_return_value *ret, void const *cache, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce) { ethash_cache c; c.mem = (void*)cache; ethash_light(ret, &c, params, header_hash, nonce); } ethash_cache c;
static inline void ethash_prep_full(void *full, ethash_params const *params, void const *cache) { ethash_cache c; c.mem = (void*)cache; ethash_compute_full_data(full, params, &c); } c.mem = cache;
static inline void ethash_compute_full(ethash_return_value *ret, void const *full, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce) { ethash_full(ret, full, params, header_hash, nonce); } ethash_mkcache(&c, params, seed);
}
static inline void ethash_compute_light(ethash_return_value *ret, void const *cache, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce) {
ethash_cache c;
c.mem = (void *) cache;
ethash_light(ret, &c, params, header_hash, nonce);
}
static inline void ethash_prep_full(void *full, ethash_params const *params, void const *cache) {
ethash_cache c;
c.mem = (void *) cache;
ethash_compute_full_data(full, params, &c);
}
static inline void ethash_compute_full(ethash_return_value *ret, void const *full, ethash_params const *params, const uint8_t header_hash[32], const uint64_t nonce) {
ethash_full(ret, full, params, header_hash, nonce);
}
// Returns if hash is less than or equal to difficulty
static inline int ethash_check_difficulty( static inline int ethash_check_difficulty(
const uint8_t hash[32], const uint8_t hash[32],
const uint8_t difficulty[32]) { const uint8_t difficulty[32]) {
@ -81,7 +102,7 @@ static inline int ethash_check_difficulty(
if (hash[i] == difficulty[i]) continue; if (hash[i] == difficulty[i]) continue;
return hash[i] < difficulty[i]; return hash[i] < difficulty[i];
} }
return 0; return 1;
} }
int ethash_quick_check_difficulty( int ethash_quick_check_difficulty(

View File

@ -1,12 +1,12 @@
/* /*
This file is part of cpp-ethereum. This file is part of ethash.
cpp-ethereum is free software: you can redistribute it and/or modify ethash is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or the Free Software Foundation, either version 3 of the License, or
(at your option) any later version. (at your option) any later version.
cpp-ethereum is distributed in the hope that it will be useful, ethash is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details. GNU General Public License for more details.
@ -14,7 +14,7 @@
You should have received a copy of the GNU General Public License You should have received a copy of the GNU General Public License
along with cpp-ethereum. If not, see <http://www.gnu.org/licenses/>. along with cpp-ethereum. If not, see <http://www.gnu.org/licenses/>.
*/ */
/** @file dash.cpp /** @file internal.c
* @author Tim Hughes <tim@twistedfury.com> * @author Tim Hughes <tim@twistedfury.com>
* @author Matthew Wampler-Doty * @author Matthew Wampler-Doty
* @date 2015 * @date 2015
@ -38,12 +38,12 @@
#endif // WITH_CRYPTOPP #endif // WITH_CRYPTOPP
size_t ethash_get_datasize(const uint32_t block_number) { size_t ethash_get_datasize(const uint32_t block_number) {
assert(block_number / EPOCH_LENGTH < 500); assert(block_number / EPOCH_LENGTH < 2048);
return dag_sizes[block_number / EPOCH_LENGTH]; return dag_sizes[block_number / EPOCH_LENGTH];
} }
size_t ethash_get_cachesize(const uint32_t block_number) { size_t ethash_get_cachesize(const uint32_t block_number) {
assert(block_number / EPOCH_LENGTH < 500); assert(block_number / EPOCH_LENGTH < 2048);
return cache_sizes[block_number / EPOCH_LENGTH]; return cache_sizes[block_number / EPOCH_LENGTH];
} }
@ -55,7 +55,7 @@ void static ethash_compute_cache_nodes(
ethash_params const *params, ethash_params const *params,
const uint8_t seed[32]) { const uint8_t seed[32]) {
assert((params->cache_size % sizeof(node)) == 0); assert((params->cache_size % sizeof(node)) == 0);
uint32_t const num_nodes = (uint32_t)(params->cache_size / sizeof(node)); uint32_t const num_nodes = (uint32_t) (params->cache_size / sizeof(node));
SHA3_512(nodes[0].bytes, seed, 32); SHA3_512(nodes[0].bytes, seed, 32);
@ -68,10 +68,9 @@ void static ethash_compute_cache_nodes(
uint32_t const idx = nodes[i].words[0] % num_nodes; uint32_t const idx = nodes[i].words[0] % num_nodes;
node data; node data;
data = nodes[(num_nodes - 1 + i) % num_nodes]; data = nodes[(num_nodes - 1 + i) % num_nodes];
for (unsigned w = 0; w != NODE_WORDS; ++w) for (unsigned w = 0; w != NODE_WORDS; ++w) {
{ data.words[w] ^= nodes[idx].words[w];
data.words[w] ^= nodes[idx].words[w]; }
}
SHA3_512(nodes[i].bytes, data.bytes, sizeof(data)); SHA3_512(nodes[i].bytes, data.bytes, sizeof(data));
} }
} }
@ -86,7 +85,7 @@ void static ethash_compute_cache_nodes(
} }
void ethash_mkcache( void ethash_mkcache(
ethash_cache *cache, ethash_cache *cache,
ethash_params const *params, ethash_params const *params,
const uint8_t seed[32]) { const uint8_t seed[32]) {
node *nodes = (node *) cache->mem; node *nodes = (node *) cache->mem;
@ -99,13 +98,13 @@ void ethash_calculate_dag_item(
const struct ethash_params *params, const struct ethash_params *params,
const struct ethash_cache *cache) { const struct ethash_cache *cache) {
uint32_t num_parent_nodes = (uint32_t)(params->cache_size / sizeof(node)); uint32_t num_parent_nodes = (uint32_t) (params->cache_size / sizeof(node));
node const *cache_nodes = (node const *) cache->mem; node const *cache_nodes = (node const *) cache->mem;
node const *init = &cache_nodes[node_index % num_parent_nodes]; node const *init = &cache_nodes[node_index % num_parent_nodes];
memcpy(ret, init, sizeof(node)); memcpy(ret, init, sizeof(node));
ret->words[0] ^= node_index; ret->words[0] ^= node_index;
SHA3_512(ret->bytes, ret->bytes, sizeof(node)); SHA3_512(ret->bytes, ret->bytes, sizeof(node));
#if defined(_M_X64) && ENABLE_SSE #if defined(_M_X64) && ENABLE_SSE
__m128i const fnv_prime = _mm_set1_epi32(FNV_PRIME); __m128i const fnv_prime = _mm_set1_epi32(FNV_PRIME);
@ -115,12 +114,11 @@ void ethash_calculate_dag_item(
__m128i xmm3 = ret->xmm[3]; __m128i xmm3 = ret->xmm[3];
#endif #endif
for (unsigned i = 0; i != DAG_PARENTS; ++i) for (unsigned i = 0; i != DATASET_PARENTS; ++i) {
{ uint32_t parent_index = ((node_index ^ i) * FNV_PRIME ^ ret->words[i % NODE_WORDS]) % num_parent_nodes;
uint32_t parent_index = ((node_index ^ i)*FNV_PRIME ^ ret->words[i % NODE_WORDS]) % num_parent_nodes;
node const *parent = &cache_nodes[parent_index]; node const *parent = &cache_nodes[parent_index];
#if defined(_M_X64) && ENABLE_SSE #if defined(_M_X64) && ENABLE_SSE
{ {
xmm0 = _mm_mullo_epi32(xmm0, fnv_prime); xmm0 = _mm_mullo_epi32(xmm0, fnv_prime);
xmm1 = _mm_mullo_epi32(xmm1, fnv_prime); xmm1 = _mm_mullo_epi32(xmm1, fnv_prime);
@ -143,10 +141,10 @@ void ethash_calculate_dag_item(
ret->words[w] = fnv_hash(ret->words[w], parent->words[w]); ret->words[w] = fnv_hash(ret->words[w], parent->words[w]);
} }
} }
#endif #endif
} }
SHA3_512(ret->bytes, ret->bytes, sizeof(node)); SHA3_512(ret->bytes, ret->bytes, sizeof(node));
} }
void ethash_compute_full_data( void ethash_compute_full_data(
@ -164,7 +162,7 @@ void ethash_compute_full_data(
} }
static void ethash_hash( static void ethash_hash(
ethash_return_value * ret, ethash_return_value *ret,
node const *full_nodes, node const *full_nodes,
ethash_cache const *cache, ethash_cache const *cache,
ethash_params const *params, ethash_params const *params,
@ -174,7 +172,7 @@ static void ethash_hash(
assert((params->full_size % MIX_WORDS) == 0); assert((params->full_size % MIX_WORDS) == 0);
// pack hash and nonce together into first 40 bytes of s_mix // pack hash and nonce together into first 40 bytes of s_mix
assert(sizeof(node)*8 == 512); assert(sizeof(node) * 8 == 512);
node s_mix[MIX_NODES + 1]; node s_mix[MIX_NODES + 1];
memcpy(s_mix[0].bytes, header_hash, 32); memcpy(s_mix[0].bytes, header_hash, 32);
@ -193,23 +191,21 @@ static void ethash_hash(
} }
#endif #endif
node* const mix = s_mix + 1; node *const mix = s_mix + 1;
for (unsigned w = 0; w != MIX_WORDS; ++w) { for (unsigned w = 0; w != MIX_WORDS; ++w) {
mix->words[w] = s_mix[0].words[w % NODE_WORDS]; mix->words[w] = s_mix[0].words[w % NODE_WORDS];
} }
unsigned const unsigned const
page_size = sizeof(uint32_t) * MIX_WORDS, page_size = sizeof(uint32_t) * MIX_WORDS,
num_full_pages = (unsigned)(params->full_size / page_size); num_full_pages = (unsigned) (params->full_size / page_size);
for (unsigned i = 0; i != ACCESSES; ++i) for (unsigned i = 0; i != ACCESSES; ++i) {
{ uint32_t const index = ((s_mix->words[0] ^ i) * FNV_PRIME ^ mix->words[i % MIX_WORDS]) % num_full_pages;
uint32_t const index = ((s_mix->words[0] ^ i)*FNV_PRIME ^ mix->words[i % MIX_WORDS]) % num_full_pages;
for (unsigned n = 0; n != MIX_NODES; ++n) for (unsigned n = 0; n != MIX_NODES; ++n) {
{ const node *dag_node = &full_nodes[MIX_NODES * index + n];
const node * dag_node = &full_nodes[MIX_NODES * index + n];
if (!full_nodes) { if (!full_nodes) {
node tmp_node; node tmp_node;
@ -217,7 +213,7 @@ static void ethash_hash(
dag_node = &tmp_node; dag_node = &tmp_node;
} }
#if defined(_M_X64) && ENABLE_SSE #if defined(_M_X64) && ENABLE_SSE
{ {
__m128i fnv_prime = _mm_set1_epi32(FNV_PRIME); __m128i fnv_prime = _mm_set1_epi32(FNV_PRIME);
__m128i xmm0 = _mm_mullo_epi32(fnv_prime, mix[n].xmm[0]); __m128i xmm0 = _mm_mullo_epi32(fnv_prime, mix[n].xmm[0]);
@ -235,20 +231,19 @@ static void ethash_hash(
mix[n].words[w] = fnv_hash(mix[n].words[w], dag_node->words[w]); mix[n].words[w] = fnv_hash(mix[n].words[w], dag_node->words[w]);
} }
} }
#endif #endif
} }
} }
// compress mix // compress mix
for (unsigned w = 0; w != MIX_WORDS; w += 4) for (unsigned w = 0; w != MIX_WORDS; w += 4) {
{ uint32_t reduction = mix->words[w + 0];
uint32_t reduction = mix->words[w+0]; reduction = reduction * FNV_PRIME ^ mix->words[w + 1];
reduction = reduction*FNV_PRIME ^ mix->words[w+1]; reduction = reduction * FNV_PRIME ^ mix->words[w + 2];
reduction = reduction*FNV_PRIME ^ mix->words[w+2]; reduction = reduction * FNV_PRIME ^ mix->words[w + 3];
reduction = reduction*FNV_PRIME ^ mix->words[w+3]; mix->words[w / 4] = reduction;
mix->words[w/4] = reduction; }
}
#if BYTE_ORDER != LITTLE_ENDIAN #if BYTE_ORDER != LITTLE_ENDIAN
for (unsigned w = 0; w != MIX_WORDS/4; ++w) { for (unsigned w = 0; w != MIX_WORDS/4; ++w) {
@ -258,7 +253,7 @@ static void ethash_hash(
memcpy(ret->mix_hash, mix->bytes, 32); memcpy(ret->mix_hash, mix->bytes, 32);
// final Keccak hash // final Keccak hash
SHA3_256(ret->result, s_mix->bytes, 64+32); // Keccak-256(s + compressed_mix) SHA3_256(ret->result, s_mix->bytes, 64 + 32); // Keccak-256(s + compressed_mix)
} }
void ethash_quick_hash( void ethash_quick_hash(
@ -267,7 +262,7 @@ void ethash_quick_hash(
const uint64_t nonce, const uint64_t nonce,
const uint8_t mix_hash[32]) { const uint8_t mix_hash[32]) {
uint8_t buf[64+32]; uint8_t buf[64 + 32];
memcpy(buf, header_hash, 32); memcpy(buf, header_hash, 32);
#if BYTE_ORDER != LITTLE_ENDIAN #if BYTE_ORDER != LITTLE_ENDIAN
nonce = fix_endian64(nonce); nonce = fix_endian64(nonce);
@ -275,7 +270,14 @@ void ethash_quick_hash(
memcpy(&(buf[32]), &nonce, 8); memcpy(&(buf[32]), &nonce, 8);
SHA3_512(buf, buf, 40); SHA3_512(buf, buf, 40);
memcpy(&(buf[64]), mix_hash, 32); memcpy(&(buf[64]), mix_hash, 32);
SHA3_256(return_hash, buf, 64+32); SHA3_256(return_hash, buf, 64 + 32);
}
void ethash_get_seedhash(uint8_t seedhash[32], const uint32_t block_number) {
memset(seedhash, 0, 32);
const uint32_t epochs = block_number / EPOCH_LENGTH;
for (uint32_t i = 0; i < epochs; ++i)
SHA3_256(seedhash, seedhash, 32);
} }
int ethash_quick_check_difficulty( int ethash_quick_check_difficulty(
@ -289,10 +291,10 @@ int ethash_quick_check_difficulty(
return ethash_check_difficulty(return_hash, difficulty); return ethash_check_difficulty(return_hash, difficulty);
} }
void ethash_full(ethash_return_value * ret, void const *full_mem, ethash_params const *params, const uint8_t previous_hash[32], const uint64_t nonce) { void ethash_full(ethash_return_value *ret, void const *full_mem, ethash_params const *params, const uint8_t previous_hash[32], const uint64_t nonce) {
ethash_hash(ret, (node const *) full_mem, NULL, params, previous_hash, nonce); ethash_hash(ret, (node const *) full_mem, NULL, params, previous_hash, nonce);
} }
void ethash_light(ethash_return_value * ret, ethash_cache const *cache, ethash_params const *params, const uint8_t previous_hash[32], const uint64_t nonce) { void ethash_light(ethash_return_value *ret, ethash_cache const *cache, ethash_params const *params, const uint8_t previous_hash[32], const uint64_t nonce) {
ethash_hash(ret, NULL, cache, params, previous_hash, nonce); ethash_hash(ret, NULL, cache, params, previous_hash, nonce);
} }

View File

@ -3,7 +3,7 @@
#include "endian.h" #include "endian.h"
#include "ethash.h" #include "ethash.h"
#define ENABLE_SSE 1 #define ENABLE_SSE 0
#if defined(_M_X64) && ENABLE_SSE #if defined(_M_X64) && ENABLE_SSE
#include <smmintrin.h> #include <smmintrin.h>

View File

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

View File

@ -1,18 +1,18 @@
/* /*
This file is part of cpp-ethereum. This file is part of ethash.
cpp-ethereum is free software: you can redistribute it and/or modify ethash is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or the Free Software Foundation, either version 3 of the License, or
(at your option) any later version. (at your option) any later version.
cpp-ethereum is distributed in the hope that it will be useful, ethash is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details. GNU General Public License for more details.
You should have received a copy of the GNU General Public License You should have received a copy of the GNU General Public License
along with cpp-ethereum. If not, see <http://www.gnu.org/licenses/>. along with ethash. If not, see <http://www.gnu.org/licenses/>.
*/ */
/** @file util.h /** @file util.h
* @author Tim Hughes <tim@twistedfury.com> * @author Tim Hughes <tim@twistedfury.com>

View File

@ -1,64 +1,305 @@
#include <Python.h> #include <Python.h>
#include <alloca.h> #include <alloca.h>
#include <stdint.h> #include <stdint.h>
#include <stdlib.h>
#include <time.h>
#include "../libethash/ethash.h" #include "../libethash/ethash.h"
static PyObject* #define MIX_WORDS (MIX_BYTES/4)
get_cache_size(PyObject* self, PyObject* args)
{ static PyObject *
get_cache_size(PyObject *self, PyObject *args) {
unsigned long block_number; unsigned long block_number;
if (!PyArg_ParseTuple(args, "k", &block_number)) if (!PyArg_ParseTuple(args, "k", &block_number))
return 0; return 0;
if (block_number >= EPOCH_LENGTH * 2048) {
char error_message[1024];
sprintf(error_message, "Block number must be less than %i (was %lu)", EPOCH_LENGTH * 2048, block_number);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
return Py_BuildValue("i", ethash_get_cachesize(block_number)); return Py_BuildValue("i", ethash_get_cachesize(block_number));
} }
static PyObject* static PyObject *
get_full_size(PyObject* self, PyObject* args) get_full_size(PyObject *self, PyObject *args) {
{
unsigned long block_number; unsigned long block_number;
if (!PyArg_ParseTuple(args, "k", &block_number)) if (!PyArg_ParseTuple(args, "k", &block_number))
return 0; return 0;
if (block_number >= EPOCH_LENGTH * 2048) {
char error_message[1024];
sprintf(error_message, "Block number must be less than %i (was %lu)", EPOCH_LENGTH * 2048, block_number);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
return Py_BuildValue("i", ethash_get_datasize(block_number)); return Py_BuildValue("i", ethash_get_datasize(block_number));
} }
static PyObject* static PyObject *
mkcache(PyObject* self, PyObject* args) mkcache_bytes(PyObject *self, PyObject *args) {
{ char *seed;
char * seed;
unsigned long cache_size; unsigned long cache_size;
int seed_len; int seed_len;
if (!PyArg_ParseTuple(args, "ks#", &cache_size, &seed, &seed_len)) if (!PyArg_ParseTuple(args, "ks#", &cache_size, &seed, &seed_len))
return 0; return 0;
if (seed_len != 32) if (seed_len != 32) {
{ char error_message[1024];
PyErr_SetString(PyExc_ValueError, sprintf(error_message, "Seed must be 32 bytes long (was %i)", seed_len);
"Seed must be 32 bytes long");
return 0; PyErr_SetString(PyExc_ValueError, error_message);
return 0;
} }
printf("cache size: %lu\n", cache_size);
ethash_params params; ethash_params params;
params.cache_size = (size_t) cache_size; params.cache_size = (size_t) cache_size;
ethash_cache cache; ethash_cache cache;
cache.mem = alloca(cache_size); cache.mem = malloc(cache_size);
ethash_mkcache(&cache, &params, (uint8_t *) seed); ethash_mkcache(&cache, &params, (uint8_t *) seed);
return PyString_FromStringAndSize(cache.mem, cache_size); PyObject * val = Py_BuildValue("s#", cache.mem, cache_size);
free(cache.mem);
return val;
} }
static PyMethodDef CoreMethods[] = static PyObject *
{ calc_dataset_bytes(PyObject *self, PyObject *args) {
{"get_cache_size", get_cache_size, METH_VARARGS, "Get the cache size for a given block number"}, char *cache_bytes;
{"get_full_size", get_full_size, METH_VARARGS, "Get the full size for a given block number"}, unsigned long full_size;
{"mkcache", mkcache, METH_VARARGS, "Makes the cache for given parameters and seed hash"}, int cache_size;
{NULL, NULL, 0, NULL}
}; if (!PyArg_ParseTuple(args, "ks#", &full_size, &cache_bytes, &cache_size))
return 0;
if (full_size % MIX_WORDS != 0) {
char error_message[1024];
sprintf(error_message, "The size of data set must be a multiple of %i bytes (was %lu)", MIX_WORDS, full_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
if (cache_size % HASH_BYTES != 0) {
char error_message[1024];
sprintf(error_message, "The size of the cache must be a multiple of %i bytes (was %i)", HASH_BYTES, cache_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
ethash_params params;
params.cache_size = (size_t) cache_size;
params.full_size = (size_t) full_size;
ethash_cache cache;
cache.mem = (void *) cache_bytes;
void *mem = malloc(params.full_size);
ethash_compute_full_data(mem, &params, &cache);
PyObject * val = Py_BuildValue("s#", (char *) mem, full_size);
free(mem);
return val;
}
// hashimoto_light(full_size, cache, header, nonce)
static PyObject *
hashimoto_light(PyObject *self, PyObject *args) {
char *cache_bytes, *header;
unsigned long full_size;
unsigned long long nonce;
int cache_size, header_size;
if (!PyArg_ParseTuple(args, "ks#s#K", &full_size, &cache_bytes, &cache_size, &header, &header_size, &nonce))
return 0;
if (full_size % MIX_WORDS != 0) {
char error_message[1024];
sprintf(error_message, "The size of data set must be a multiple of %i bytes (was %lu)", MIX_WORDS, full_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
if (cache_size % HASH_BYTES != 0) {
char error_message[1024];
sprintf(error_message, "The size of the cache must be a multiple of %i bytes (was %i)", HASH_BYTES, cache_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
if (header_size != 32) {
char error_message[1024];
sprintf(error_message, "Seed must be 32 bytes long (was %i)", header_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
ethash_return_value out;
ethash_params params;
params.cache_size = (size_t) cache_size;
params.full_size = (size_t) full_size;
ethash_cache cache;
cache.mem = (void *) cache_bytes;
ethash_light(&out, &cache, &params, (uint8_t *) header, nonce);
return Py_BuildValue("{s:s#,s:s#}",
"mix digest", out.mix_hash, 32,
"result", out.result, 32);
}
// hashimoto_full(dataset, header, nonce)
static PyObject *
hashimoto_full(PyObject *self, PyObject *args) {
char *full_bytes, *header;
unsigned long long nonce;
int full_size, header_size;
if (!PyArg_ParseTuple(args, "s#s#K", &full_bytes, &full_size, &header, &header_size, &nonce))
return 0;
if (full_size % MIX_WORDS != 0) {
char error_message[1024];
sprintf(error_message, "The size of data set must be a multiple of %i bytes (was %i)", MIX_WORDS, full_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
if (header_size != 32) {
char error_message[1024];
sprintf(error_message, "Header must be 32 bytes long (was %i)", header_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
ethash_return_value out;
ethash_params params;
params.full_size = (size_t) full_size;
ethash_full(&out, (void *) full_bytes, &params, (uint8_t *) header, nonce);
return Py_BuildValue("{s:s#, s:s#}",
"mix digest", out.mix_hash, 32,
"result", out.result, 32);
}
// mine(dataset_bytes, header, difficulty_bytes)
static PyObject *
mine(PyObject *self, PyObject *args) {
char *full_bytes, *header, *difficulty;
srand(time(0));
uint64_t nonce = ((uint64_t) rand()) << 32 | rand();
int full_size, header_size, difficulty_size;
if (!PyArg_ParseTuple(args, "s#s#s#", &full_bytes, &full_size, &header, &header_size, &difficulty, &difficulty_size))
return 0;
if (full_size % MIX_WORDS != 0) {
char error_message[1024];
sprintf(error_message, "The size of data set must be a multiple of %i bytes (was %i)", MIX_WORDS, full_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
if (header_size != 32) {
char error_message[1024];
sprintf(error_message, "Header must be 32 bytes long (was %i)", header_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
if (difficulty_size != 32) {
char error_message[1024];
sprintf(error_message, "Difficulty must be an array of 32 bytes (only had %i)", difficulty_size);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
ethash_return_value out;
ethash_params params;
params.full_size = (size_t) full_size;
// TODO: Multi threading?
do {
ethash_full(&out, (void *) full_bytes, &params, (const uint8_t *) header, nonce++);
// TODO: disagrees with the spec https://github.com/ethereum/wiki/wiki/Ethash#mining
} while (!ethash_check_difficulty(out.result, (const uint8_t *) difficulty));
return Py_BuildValue("{s:s#, s:s#, s:K}",
"mix digest", out.mix_hash, 32,
"result", out.result, 32,
"nonce", nonce);
}
//get_seedhash(block_number)
static PyObject *
get_seedhash(PyObject *self, PyObject *args) {
unsigned long block_number;
if (!PyArg_ParseTuple(args, "k", &block_number))
return 0;
if (block_number >= EPOCH_LENGTH * 2048) {
char error_message[1024];
sprintf(error_message, "Block number must be less than %i (was %lu)", EPOCH_LENGTH * 2048, block_number);
PyErr_SetString(PyExc_ValueError, error_message);
return 0;
}
uint8_t seedhash[32];
ethash_get_seedhash(seedhash, block_number);
return Py_BuildValue("s#", (char *) seedhash, 32);
}
static PyMethodDef PyethashMethods[] =
{
{"get_cache_size", get_cache_size, METH_VARARGS,
"get_cache_size(block_number)\n\n"
"Get the cache size for a given block number\n"
"\nExample:\n"
">>> get_cache_size(0)\n"
"1048384"},
{"get_full_size", get_full_size, METH_VARARGS,
"get_full_size(block_number)\n\n"
"Get the full size for a given block number\n"
"\nExample:\n"
">>> get_full_size(0)\n"
"1073739904"
},
{"get_seedhash", get_seedhash, METH_VARARGS,
"get_seedhash(block_number)\n\n"
"Gets the seedhash for a block."},
{"mkcache_bytes", mkcache_bytes, METH_VARARGS,
"mkcache_bytes(size, header)\n\n"
"Makes a byte array for the cache for given cache size and seed hash\n"
"\nExample:\n"
">>> pyethash.mkcache_bytes( 1024, \"~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~\").encode('hex')"
"\"2da2b506f21070e1143d908e867962486d6b0a02e31d468fd5e3a7143aafa76a14201f63374314e2a6aaf84ad2eb57105dea3378378965a1b3873453bb2b78f9a8620b2ebeca41fbc773bb837b5e724d6eb2de570d99858df0d7d97067fb8103b21757873b735097b35d3bea8fd1c359a9e8a63c1540c76c9784cf8d975e995ca8620b2ebeca41fbc773bb837b5e724d6eb2de570d99858df0d7d97067fb8103b21757873b735097b35d3bea8fd1c359a9e8a63c1540c76c9784cf8d975e995ca8620b2ebeca41fbc773bb837b5e724d6eb2de570d99858df0d7d97067fb8103b21757873b735097b35d3bea8fd1c359a9e8a63c1540c76c9784cf8d975e995c259440b89fa3481c2c33171477c305c8e1e421f8d8f6d59585449d0034f3e421808d8da6bbd0b6378f567647cc6c4ba6c434592b198ad444e7284905b7c6adaf70bf43ec2daa7bd5e8951aa609ab472c124cf9eba3d38cff5091dc3f58409edcc386c743c3bd66f92408796ee1e82dd149eaefbf52b00ce33014a6eb3e50625413b072a58bc01da28262f42cbe4f87d4abc2bf287d15618405a1fe4e386fcdafbb171064bd99901d8f81dd6789396ce5e364ac944bbbd75a7827291c70b42d26385910cd53ca535ab29433dd5c5714d26e0dce95514c5ef866329c12e958097e84462197c2b32087849dab33e88b11da61d52f9dbc0b92cc61f742c07dbbf751c49d7678624ee60dfbe62e5e8c47a03d8247643f3d16ad8c8e663953bcda1f59d7e2d4a9bf0768e789432212621967a8f41121ad1df6ae1fa78782530695414c6213942865b2730375019105cae91a4c17a558d4b63059661d9f108362143107babe0b848de412e4da59168cce82bfbff3c99e022dd6ac1e559db991f2e3f7bb910cefd173e65ed00a8d5d416534e2c8416ff23977dbf3eb7180b75c71580d08ce95efeb9b0afe904ea12285a392aff0c8561ff79fca67f694a62b9e52377485c57cc3598d84cac0a9d27960de0cc31ff9bbfe455acaa62c8aa5d2cce96f345da9afe843d258a99c4eaf3650fc62efd81c7b81cd0d534d2d71eeda7a6e315d540b4473c80f8730037dc2ae3e47b986240cfc65ccc565f0d8cde0bc68a57e39a271dda57440b3598bee19f799611d25731a96b5dbbbefdff6f4f656161462633030d62560ea4e9c161cf78fc96a2ca5aaa32453a6c5dea206f766244e8c9d9a8dc61185ce37f1fc804459c5f07434f8ecb34141b8dcae7eae704c950b55556c5f40140c3714b45eddb02637513268778cbf937a33e4e33183685f9deb31ef54e90161e76d969587dd782eaa94e289420e7c2ee908517f5893a26fdb5873d68f92d118d4bcf98d7a4916794d6ab290045e30f9ea00ca547c584b8482b0331ba1539a0f2714fddc3a0b06b0cfbb6a607b8339c39bcfd6640b1f653e9d70ef6c985b\""},
{"calc_dataset_bytes", calc_dataset_bytes, METH_VARARGS,
"calc_dataset_bytes(full_size, cache_bytes)\n\n"
"Makes a byte array for the dataset for a given size given cache bytes"},
{"hashimoto_light", hashimoto_light, METH_VARARGS,
"hashimoto_light(full_size, cache_bytes, header, nonce)\n\n"
"Runs the hashimoto hashing function just using cache bytes. Takes an int (full_size), byte array (cache_bytes), another byte array (header), and an int (nonce). Returns an object containing the mix digest, and hash result."},
{"hashimoto_full", hashimoto_full, METH_VARARGS,
"hashimoto_full(dataset_bytes, header, nonce)\n\n"
"Runs the hashimoto hashing function using the dataset bytes. Useful for testing. Returns an object containing the mix digest (byte array), and hash result (another byte array)."},
{"mine", mine, METH_VARARGS,
"mine(dataset_bytes, header, difficulty_bytes)\n\n"
"Mine for an adequate header. Returns an object containing the mix digest (byte array), hash result (another byte array) and nonce (an int)."},
{NULL, NULL, 0, NULL}
};
PyMODINIT_FUNC PyMODINIT_FUNC
initcore(void) initpyethash(void) {
{ PyObject *module = Py_InitModule("pyethash", PyethashMethods);
(void) Py_InitModule("core", CoreMethods); // Following Spec: https://github.com/ethereum/wiki/wiki/Ethash#definitions
PyModule_AddIntConstant(module, "REVISION", (long) REVISION);
PyModule_AddIntConstant(module, "DATASET_BYTES_INIT", (long) DATASET_BYTES_INIT);
PyModule_AddIntConstant(module, "DATASET_BYTES_GROWTH", (long) DATASET_BYTES_GROWTH);
PyModule_AddIntConstant(module, "CACHE_BYTES_INIT", (long) CACHE_BYTES_INIT);
PyModule_AddIntConstant(module, "CACHE_BYTES_GROWTH", (long) CACHE_BYTES_GROWTH);
PyModule_AddIntConstant(module, "EPOCH_LENGTH", (long) EPOCH_LENGTH);
PyModule_AddIntConstant(module, "MIX_BYTES", (long) MIX_BYTES);
PyModule_AddIntConstant(module, "HASH_BYTES", (long) HASH_BYTES);
PyModule_AddIntConstant(module, "DATASET_PARENTS", (long) DATASET_PARENTS);
PyModule_AddIntConstant(module, "CACHE_ROUNDS", (long) CACHE_ROUNDS);
PyModule_AddIntConstant(module, "ACCESSES", (long) ACCESSES);
} }

View File

@ -4,7 +4,9 @@
#include <libethash/internal.h> #include <libethash/internal.h>
#ifdef WITH_CRYPTOPP #ifdef WITH_CRYPTOPP
#include <libethash/sha3_cryptopp.h> #include <libethash/sha3_cryptopp.h>
#else #else
#include <libethash/sha3.h> #include <libethash/sha3.h>
#endif // WITH_CRYPTOPP #endif // WITH_CRYPTOPP
@ -28,7 +30,7 @@ BOOST_AUTO_TEST_CASE(fnv_hash_check) {
uint32_t x = 1235U; uint32_t x = 1235U;
const uint32_t const uint32_t
y = 9999999U, y = 9999999U,
expected = (FNV_PRIME * x) ^ y; expected = (FNV_PRIME * x) ^y;
x = fnv_hash(x, y); x = fnv_hash(x, y);
@ -65,43 +67,50 @@ BOOST_AUTO_TEST_CASE(SHA512_check) {
BOOST_AUTO_TEST_CASE(ethash_params_init_genesis_check) { BOOST_AUTO_TEST_CASE(ethash_params_init_genesis_check) {
ethash_params params; ethash_params params;
ethash_params_init(&params, 0); ethash_params_init(&params, 0);
BOOST_REQUIRE_MESSAGE(params.full_size < DAGSIZE_BYTES_INIT, BOOST_REQUIRE_MESSAGE(params.full_size < DATASET_BYTES_INIT,
"\nfull size: " << params.full_size << "\n" "\nfull size: " << params.full_size << "\n"
<< "should be less than or equal to: " << DAGSIZE_BYTES_INIT << "\n"); << "should be less than or equal to: " << DATASET_BYTES_INIT << "\n");
BOOST_REQUIRE_MESSAGE(params.full_size + 20*MIX_BYTES >= DAGSIZE_BYTES_INIT, BOOST_REQUIRE_MESSAGE(params.full_size + 20 * MIX_BYTES >= DATASET_BYTES_INIT,
"\nfull size + 20*MIX_BYTES: " << params.full_size + 20*MIX_BYTES << "\n" "\nfull size + 20*MIX_BYTES: " << params.full_size + 20 * MIX_BYTES << "\n"
<< "should be greater than or equal to: " << DAGSIZE_BYTES_INIT << "\n"); << "should be greater than or equal to: " << DATASET_BYTES_INIT << "\n");
BOOST_REQUIRE_MESSAGE(params.cache_size < DAGSIZE_BYTES_INIT / 32, BOOST_REQUIRE_MESSAGE(params.cache_size < DATASET_BYTES_INIT / 32,
"\ncache size: " << params.cache_size << "\n" "\ncache size: " << params.cache_size << "\n"
<< "should be less than or equal to: " << DAGSIZE_BYTES_INIT / 32 << "\n"); << "should be less than or equal to: " << DATASET_BYTES_INIT / 32 << "\n");
} }
BOOST_AUTO_TEST_CASE(ethash_params_init_genesis_calcifide_check) { BOOST_AUTO_TEST_CASE(ethash_params_init_genesis_calcifide_check) {
ethash_params params; ethash_params params;
ethash_params_init(&params, 0); ethash_params_init(&params, 0);
const uint32_t expected_full_size = 1073739904; const uint32_t expected_full_size = 1073739904;
const uint32_t expected_cache_size = 1048384; const uint32_t expected_cache_size = 16776896;
BOOST_REQUIRE_MESSAGE(params.full_size == expected_full_size, BOOST_REQUIRE_MESSAGE(params.full_size == expected_full_size,
"\nexpected: " << expected_cache_size << "\n" "\nexpected: " << expected_cache_size << "\n"
<< "actual: " << params.full_size << "\n"); << "actual: " << params.full_size << "\n");
BOOST_REQUIRE_MESSAGE(params.cache_size == expected_cache_size, BOOST_REQUIRE_MESSAGE(params.cache_size == expected_cache_size,
"\nexpected: " << expected_cache_size << "\n" "\nexpected: " << expected_cache_size << "\n"
<< "actual: " << params.cache_size << "\n"); << "actual: " << params.cache_size << "\n");
} }
BOOST_AUTO_TEST_CASE(light_and_full_client_checks) { BOOST_AUTO_TEST_CASE(light_and_full_client_checks) {
ethash_params params; ethash_params params;
uint8_t seed[32], hash[32]; uint8_t seed[32], hash[32], difficulty[32];
ethash_return_value light_out, full_out; ethash_return_value light_out, full_out;
memcpy(seed, "~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~", 32); memcpy(seed, "~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~", 32);
memcpy(hash, "~~~X~~~~~~~~~~~~~~~~~~~~~~~~~~~~", 32); memcpy(hash, "~~~X~~~~~~~~~~~~~~~~~~~~~~~~~~~~", 32);
// Set the difficulty
difficulty[0] = 197;
difficulty[1] = 90;
for (int i = 2; i < 32; i++)
difficulty[i] = (uint8_t) 255;
ethash_params_init(&params, 0); ethash_params_init(&params, 0);
params.cache_size = 1024; params.cache_size = 1024;
params.full_size = 1024 * 32; params.full_size = 1024 * 32;
ethash_cache cache; ethash_cache cache;
cache.mem = alloca(params.cache_size); cache.mem = alloca(params.cache_size);
ethash_mkcache(&cache, &params, seed); ethash_mkcache(&cache, &params, seed);
node * full_mem = (node *) alloca(params.full_size); node *full_mem = (node *) alloca(params.full_size);
ethash_compute_full_data(full_mem, &params, &cache); ethash_compute_full_data(full_mem, &params, &cache);
{ {
@ -115,7 +124,6 @@ BOOST_AUTO_TEST_CASE(light_and_full_client_checks) {
} }
{ {
node node; node node;
ethash_calculate_dag_item(&node, 0, &params, &cache); ethash_calculate_dag_item(&node, 0, &params, &cache);
@ -128,7 +136,7 @@ BOOST_AUTO_TEST_CASE(light_and_full_client_checks) {
} }
{ {
for (int i = 0 ; i < params.full_size / sizeof(node) ; ++i ) { for (int i = 0; i < params.full_size / sizeof(node); ++i) {
for (uint32_t j = 0; j < 32; ++j) { for (uint32_t j = 0; j < 32; ++j) {
node expected_node; node expected_node;
ethash_calculate_dag_item(&expected_node, j, &params, &cache); ethash_calculate_dag_item(&expected_node, j, &params, &cache);
@ -187,6 +195,12 @@ BOOST_AUTO_TEST_CASE(light_and_full_client_checks) {
BOOST_REQUIRE_MESSAGE(full_mix_hash_string == light_mix_hash_string, BOOST_REQUIRE_MESSAGE(full_mix_hash_string == light_mix_hash_string,
"\nlight mix hash: " << light_mix_hash_string.c_str() << "\n" "\nlight mix hash: " << light_mix_hash_string.c_str() << "\n"
<< "full mix hash: " << full_mix_hash_string.c_str() << "\n"); << "full mix hash: " << full_mix_hash_string.c_str() << "\n");
BOOST_REQUIRE_MESSAGE(ethash_check_difficulty(full_out.result, difficulty),
"ethash_check_difficulty failed"
);
BOOST_REQUIRE_MESSAGE(ethash_quick_check_difficulty(hash, 5U, full_out.mix_hash, difficulty),
"ethash_quick_check_difficulty failed"
);
} }
} }
@ -199,14 +213,14 @@ BOOST_AUTO_TEST_CASE(ethash_check_difficulty_check) {
memcpy(target, "22222222222222222222222222222222", 32); memcpy(target, "22222222222222222222222222222222", 32);
BOOST_REQUIRE_MESSAGE( BOOST_REQUIRE_MESSAGE(
ethash_check_difficulty(hash, target), ethash_check_difficulty(hash, target),
"\nexpected \"" << hash << "\" to have less difficulty than \"" << target << "\"\n"); "\nexpected \"" << std::string((char *) hash, 32).c_str() << "\" to have the same or less difficulty than \"" << std::string((char *) target, 32).c_str() << "\"\n");
BOOST_REQUIRE_MESSAGE( BOOST_REQUIRE_MESSAGE(
!ethash_check_difficulty(hash, hash), ethash_check_difficulty(hash, hash),
"\nexpected \"" << hash << "\" to have the same difficulty as \"" << hash << "\"\n"); "\nexpected \"" << hash << "\" to have the same or less difficulty than \"" << hash << "\"\n");
memcpy(target, "11111111111111111111111111111112", 32); memcpy(target, "11111111111111111111111111111112", 32);
BOOST_REQUIRE_MESSAGE( BOOST_REQUIRE_MESSAGE(
ethash_check_difficulty(hash, target), ethash_check_difficulty(hash, target),
"\nexpected \"" << hash << "\" to have less difficulty than \"" << target << "\"\n"); "\nexpected \"" << hash << "\" to have the same or less difficulty than \"" << target << "\"\n");
memcpy(target, "11111111111111111111111111111110", 32); memcpy(target, "11111111111111111111111111111110", 32);
BOOST_REQUIRE_MESSAGE( BOOST_REQUIRE_MESSAGE(
!ethash_check_difficulty(hash, target), !ethash_check_difficulty(hash, target),

View File

@ -0,0 +1,82 @@
package ethashTest
import (
"bytes"
"crypto/rand"
"encoding/hex"
"log"
"math/big"
"testing"
"github.com/ethereum/ethash"
"github.com/ethereum/go-ethereum/core"
"github.com/ethereum/go-ethereum/ethdb"
)
func TestEthash(t *testing.T) {
seedHash := make([]byte, 32)
_, err := rand.Read(seedHash)
if err != nil {
panic(err)
}
db, err := ethdb.NewMemDatabase()
if err != nil {
panic(err)
}
blockProcessor, err := core.NewCanonical(5, db)
if err != nil {
panic(err)
}
log.Println("Block Number: ", blockProcessor.ChainManager().CurrentBlock().Number())
e := ethash.New(blockProcessor.ChainManager())
miningHash := make([]byte, 32)
if _, err := rand.Read(miningHash); err != nil {
panic(err)
}
diff := big.NewInt(10000)
log.Println("difficulty", diff)
nonce := uint64(0)
ghash_full := e.FullHash(nonce, miningHash)
log.Printf("ethash full (on nonce): %x %x\n", ghash_full, nonce)
ghash_light := e.LightHash(nonce, miningHash)
log.Printf("ethash light (on nonce): %x %x\n", ghash_light, nonce)
if bytes.Compare(ghash_full, ghash_light) != 0 {
t.Errorf("full: %x, light: %x", ghash_full, ghash_light)
}
}
func TestGetSeedHash(t *testing.T) {
seed0, err := ethash.GetSeedHash(0)
if err != nil {
t.Errorf("Failed to get seedHash for block 0: %v", err)
}
if bytes.Compare(seed0, []byte{0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}) != 0 {
log.Printf("seedHash for block 0 should be 0s, was: %v\n", seed0)
}
seed1, err := ethash.GetSeedHash(30000)
if err != nil {
t.Error(err)
}
// From python:
// > from pyethash import get_seedhash
// > get_seedhash(30000)
expectedSeed1, err := hex.DecodeString("290decd9548b62a8d60345a988386fc84ba6bc95484008f6362f93160ef3e563")
if err != nil {
t.Error(err)
}
if bytes.Compare(seed1, expectedSeed1) != 0 {
log.Printf("seedHash for block 1 should be: %v,\nactual value: %v\n", expectedSeed1, seed1)
}
}

View File

@ -0,0 +1,20 @@
#!/bin/bash
# Strict mode
set -e
SOURCE="${BASH_SOURCE[0]}"
while [ -h "$SOURCE" ]; do
DIR="$( cd -P "$( dirname "$SOURCE" )" && pwd )"
SOURCE="$(readlink "$SOURCE")"
[[ $SOURCE != /* ]] && SOURCE="$DIR/$SOURCE"
done
TEST_DIR="$( cd -P "$( dirname "$SOURCE" )" && pwd )"
export GOPATH=${HOME}/.go
export PATH=$PATH:$GOPATH/bin
echo "# getting go dependencies (can take some time)..."
cd ${TEST_DIR}/../.. && go get
cd ${GOPATH}/src/github.com/ethereum/go-ethereum
git checkout poc-9
cd ${TEST_DIR} && go test

View File

@ -1,2 +1,3 @@
pyethereum==0.7.522 pyethereum==0.7.522
nose==1.3.4 nose==1.3.4
pysha3==0.3

View File

@ -14,6 +14,6 @@ TEST_DIR="$( cd -P "$( dirname "$SOURCE" )" && pwd )"
[ -d $TEST_DIR/python-virtual-env ] || virtualenv --system-site-packages $TEST_DIR/python-virtual-env [ -d $TEST_DIR/python-virtual-env ] || virtualenv --system-site-packages $TEST_DIR/python-virtual-env
source $TEST_DIR/python-virtual-env/bin/activate source $TEST_DIR/python-virtual-env/bin/activate
pip install -r $TEST_DIR/requirements.txt > /dev/null pip install -r $TEST_DIR/requirements.txt > /dev/null
pip install -e $TEST_DIR/../.. > /dev/null pip install --upgrade --no-deps --force-reinstall -e $TEST_DIR/../..
cd $TEST_DIR cd $TEST_DIR
nosetests --with-doctest -v nosetests --with-doctest -v --nocapture

View File

@ -4,42 +4,102 @@ from random import randint
def test_get_cache_size_not_None(): def test_get_cache_size_not_None():
for _ in range(100): for _ in range(100):
block_num = randint(0,12456789) block_num = randint(0,12456789)
out = pyethash.core.get_cache_size(block_num) out = pyethash.get_cache_size(block_num)
assert out != None assert out != None
def test_get_full_size_not_None(): def test_get_full_size_not_None():
for _ in range(100): for _ in range(100):
block_num = randint(0,12456789) block_num = randint(0,12456789)
out = pyethash.core.get_full_size(block_num) out = pyethash.get_full_size(block_num)
assert out != None assert out != None
def test_get_cache_size_based_on_EPOCH(): def test_get_cache_size_based_on_EPOCH():
for _ in range(100): for _ in range(100):
block_num = randint(0,12456789) block_num = randint(0,12456789)
out1 = pyethash.core.get_cache_size(block_num) out1 = pyethash.get_cache_size(block_num)
out2 = pyethash.core.get_cache_size((block_num // pyethash.EPOCH_LENGTH) * pyethash.EPOCH_LENGTH) out2 = pyethash.get_cache_size((block_num // pyethash.EPOCH_LENGTH) * pyethash.EPOCH_LENGTH)
assert out1 == out2 assert out1 == out2
def test_get_full_size_based_on_EPOCH(): def test_get_full_size_based_on_EPOCH():
for _ in range(100): for _ in range(100):
block_num = randint(0,12456789) block_num = randint(0,12456789)
out1 = pyethash.core.get_full_size(block_num) out1 = pyethash.get_full_size(block_num)
out2 = pyethash.core.get_full_size((block_num // pyethash.EPOCH_LENGTH) * pyethash.EPOCH_LENGTH) out2 = pyethash.get_full_size((block_num // pyethash.EPOCH_LENGTH) * pyethash.EPOCH_LENGTH)
assert out1 == out2 assert out1 == out2
#def test_get_params_based_on_EPOCH(): # See light_and_full_client_checks in test.cpp
# block_num = 123456 def test_mkcache_is_as_expected():
# out1 = pyethash.core.get_params(block_num) actual = pyethash.mkcache_bytes(
# out2 = pyethash.core.get_params((block_num // pyethash.EPOCH_LENGTH) * pyethash.EPOCH_LENGTH) 1024,
# assert out1["DAG Size"] == out2["DAG Size"] "~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~").encode('hex')
# assert out1["Cache Size"] == out2["Cache Size"] expected = "2da2b506f21070e1143d908e867962486d6b0a02e31d468fd5e3a7143aafa76a14201f63374314e2a6aaf84ad2eb57105dea3378378965a1b3873453bb2b78f9a8620b2ebeca41fbc773bb837b5e724d6eb2de570d99858df0d7d97067fb8103b21757873b735097b35d3bea8fd1c359a9e8a63c1540c76c9784cf8d975e995ca8620b2ebeca41fbc773bb837b5e724d6eb2de570d99858df0d7d97067fb8103b21757873b735097b35d3bea8fd1c359a9e8a63c1540c76c9784cf8d975e995ca8620b2ebeca41fbc773bb837b5e724d6eb2de570d99858df0d7d97067fb8103b21757873b735097b35d3bea8fd1c359a9e8a63c1540c76c9784cf8d975e995c259440b89fa3481c2c33171477c305c8e1e421f8d8f6d59585449d0034f3e421808d8da6bbd0b6378f567647cc6c4ba6c434592b198ad444e7284905b7c6adaf70bf43ec2daa7bd5e8951aa609ab472c124cf9eba3d38cff5091dc3f58409edcc386c743c3bd66f92408796ee1e82dd149eaefbf52b00ce33014a6eb3e50625413b072a58bc01da28262f42cbe4f87d4abc2bf287d15618405a1fe4e386fcdafbb171064bd99901d8f81dd6789396ce5e364ac944bbbd75a7827291c70b42d26385910cd53ca535ab29433dd5c5714d26e0dce95514c5ef866329c12e958097e84462197c2b32087849dab33e88b11da61d52f9dbc0b92cc61f742c07dbbf751c49d7678624ee60dfbe62e5e8c47a03d8247643f3d16ad8c8e663953bcda1f59d7e2d4a9bf0768e789432212621967a8f41121ad1df6ae1fa78782530695414c6213942865b2730375019105cae91a4c17a558d4b63059661d9f108362143107babe0b848de412e4da59168cce82bfbff3c99e022dd6ac1e559db991f2e3f7bb910cefd173e65ed00a8d5d416534e2c8416ff23977dbf3eb7180b75c71580d08ce95efeb9b0afe904ea12285a392aff0c8561ff79fca67f694a62b9e52377485c57cc3598d84cac0a9d27960de0cc31ff9bbfe455acaa62c8aa5d2cce96f345da9afe843d258a99c4eaf3650fc62efd81c7b81cd0d534d2d71eeda7a6e315d540b4473c80f8730037dc2ae3e47b986240cfc65ccc565f0d8cde0bc68a57e39a271dda57440b3598bee19f799611d25731a96b5dbbbefdff6f4f656161462633030d62560ea4e9c161cf78fc96a2ca5aaa32453a6c5dea206f766244e8c9d9a8dc61185ce37f1fc804459c5f07434f8ecb34141b8dcae7eae704c950b55556c5f40140c3714b45eddb02637513268778cbf937a33e4e33183685f9deb31ef54e90161e76d969587dd782eaa94e289420e7c2ee908517f5893a26fdb5873d68f92d118d4bcf98d7a4916794d6ab290045e30f9ea00ca547c584b8482b0331ba1539a0f2714fddc3a0b06b0cfbb6a607b8339c39bcfd6640b1f653e9d70ef6c985b"
# assert actual == expected
#def test_get_params_returns_different_values_based_on_different_block_input():
# out1 = pyethash.core.get_params(123456) def test_calc_dataset_is_not_None():
# out2 = pyethash.core.get_params(12345) cache = pyethash.mkcache_bytes(
# assert out1["DAG Size"] != out2["DAG Size"] 1024,
# assert out1["Cache Size"] != out2["Cache Size"] "~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~")
# assert pyethash.calc_dataset_bytes(1024 * 32, cache) != None
#def test_get_cache_smoke_test():
# params = pyethash.core.get_params(123456) def test_light_and_full_agree():
# assert pyethash.core.mkcache(params, "~~~~") != None cache = pyethash.mkcache_bytes(
1024,
"~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~")
full_size = 1024 * 32
header = "~~~~~X~~~~~~~~~~~~~~~~~~~~~~~~~~"
light_result = pyethash.hashimoto_light(full_size, cache, header, 0)
dataset = pyethash.calc_dataset_bytes(full_size, cache)
full_result = pyethash.hashimoto_full(dataset, header, 0)
assert light_result["mix digest"] != None
assert len(light_result["mix digest"]) == 32
assert light_result["mix digest"] == full_result["mix digest"]
assert light_result["result"] != None
assert len(light_result["result"]) == 32
assert light_result["result"] == full_result["result"]
def int_to_bytes(i):
b = []
for _ in range(32):
b.append(chr(i & 0xff))
i >>= 8
b.reverse()
return "".join(b)
def test_mining_basic():
easy_difficulty = int_to_bytes(2**256 - 1)
assert easy_difficulty.encode('hex') == 'f' * 64
cache = pyethash.mkcache_bytes(
1024,
"~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~")
full_size = 1024 * 32
header = "~~~~~X~~~~~~~~~~~~~~~~~~~~~~~~~~"
dataset = pyethash.calc_dataset_bytes(full_size, cache)
# Check type of outputs
assert type(pyethash.mine(dataset,header,easy_difficulty)) == dict
assert type(pyethash.mine(dataset,header,easy_difficulty)["nonce"]) == long
assert type(pyethash.mine(dataset,header,easy_difficulty)["mix digest"]) == str
assert type(pyethash.mine(dataset,header,easy_difficulty)["result"]) == str
def test_mining_doesnt_always_return_the_same_value():
easy_difficulty1 = int_to_bytes(int(2**256 * 0.999))
# 1 in 1000 difficulty
easy_difficulty2 = int_to_bytes(int(2**256 * 0.001))
assert easy_difficulty1 != easy_difficulty2
cache = pyethash.mkcache_bytes(
1024,
"~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~")
full_size = 1024 * 32
header = "~~~~~X~~~~~~~~~~~~~~~~~~~~~~~~~~"
dataset = pyethash.calc_dataset_bytes(full_size, cache)
# Check type of outputs
assert pyethash.mine(dataset, header, easy_difficulty1)['nonce'] != pyethash.mine(dataset, header, easy_difficulty2)['nonce']
def test_get_seedhash():
assert pyethash.get_seedhash(0).encode('hex') == '0' * 64
import hashlib, sha3
expected = pyethash.get_seedhash(0)
#print "checking seed hashes:",
for i in range(0, 30000*2048, 30000):
#print i // 30000,
assert pyethash.get_seedhash(i) == expected
expected = hashlib.sha3_256(expected).digest()

View File

@ -14,7 +14,12 @@ TEST_DIR="$( cd -P "$( dirname "$SOURCE" )" && pwd )"
echo -e "\n################# Testing JS ##################" echo -e "\n################# Testing JS ##################"
# TODO: Use mocha and real testing tools instead of rolling our own # TODO: Use mocha and real testing tools instead of rolling our own
cd $TEST_DIR/../js cd $TEST_DIR/../js
node test.js if [ -x "$(which nodejs)" ] ; then
nodejs test.js
fi
if [ -x "$(which node)" ] ; then
node test.js
fi
echo -e "\n################# Testing C ##################" echo -e "\n################# Testing C ##################"
$TEST_DIR/c/test.sh $TEST_DIR/c/test.sh