Skip to content

G
Geth-Modification
Commit a91bf014 authored by Felix Lange's avatar Felix Lange
Browse files
Options

Merge pull request #434 from fjl/poc-9 

Godeps: bump github.com/ethereum/ethash
parents 76e4e233 71e51054
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 3271 additions and 2522 deletions
Godeps/Godeps.json
View file @ a91bf014
...@@ -17,8 +17,8 @@ ...@@ -17,8 +17,8 @@
}, },
{ {
"ImportPath": "github.com/ethereum/ethash", "ImportPath": "github.com/ethereum/ethash",
"Comment": "v17-23-g2561e13", "Comment": "v17-63-gbca024b",
"Rev": "2561e1322a7e8e3d4a2cc903c44b1e96340bcb27" "Rev": "bca024b0b30d83ec6798a5d4fa8c5fc6f937009a"
}, },
{ {
"ImportPath": "github.com/ethereum/serpent-go", "ImportPath": "github.com/ethereum/serpent-go",
......
Godeps/_workspace/src/github.com/ethereum/ethash/.gitignore
View file @ a91bf014
...@@ -2,4 +2,9 @@ ...@@ -2,4 +2,9 @@
.DS_Store .DS_Store
*/**/*un~ */**/*un~
.vagrant/ .vagrant/
cpp-build/ *.pyc
build/
pyethash.egg-info/
*.so
*~
*.swp
Godeps/_workspace/src/github.com/ethereum/ethash/.travis.yml 0 → 100644
View file @ a91bf014
before_install:
- sudo apt-get update -qq
- sudo apt-get install -qq wget cmake gcc bash libboost-test-dev nodejs python-pip python-dev
- sudo pip install virtualenv -q
script: "./test/test.sh"
Godeps/_workspace/src/github.com/ethereum/ethash/CMakeLists.txt
View file @ a91bf014
cmake_minimum_required(VERSION 2.8.2) 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/")
...@@ -8,7 +8,8 @@ if (WIN32 AND WANT_CRYPTOPP) ...@@ -8,7 +8,8 @@ if (WIN32 AND WANT_CRYPTOPP)
add_subdirectory(cryptopp) add_subdirectory(cryptopp)
endif() endif()
add_subdirectory(libethash) add_subdirectory(src/libethash)
add_subdirectory(libethash-cl EXCLUDE_FROM_ALL) # bin2h.cmake doesn't work
add_subdirectory(benchmark EXCLUDE_FROM_ALL) #add_subdirectory(src/libethash-cl EXCLUDE_FROM_ALL)
add_subdirectory(test EXCLUDE_FROM_ALL) add_subdirectory(src/benchmark EXCLUDE_FROM_ALL)
add_subdirectory(test/c EXCLUDE_FROM_ALL)
Godeps/_workspace/src/github.com/ethereum/ethash/Makefile 0 → 100644
View file @ a91bf014
.PHONY: clean
clean:
rm -rf *.so pyethash.egg-info/ build/ test/python/python-virtual-env/ test/c/build/ pyethash/*.{so,pyc}
Godeps/_workspace/src/github.com/ethereum/ethash/README.md 0 → 100644
View file @ a91bf014
[![Build Status](https://travis-ci.org/ethereum/ethash.svg?branch=master)](https://travis-ci.org/ethereum/ethash)
# Ethash
For details on this project, please see the Ethereum wiki:
https://github.com/ethereum/wiki/wiki/Ethash
Godeps/_workspace/src/github.com/ethereum/ethash/Vagrantfile 0 → 100644
View file @ a91bf014
# -*- mode: ruby -*-
# vi: set ft=ruby :
Vagrant.configure(2) do |config|
config.vm.box = "Ubuntu 12.04"
config.vm.box_url = "https://cloud-images.ubuntu.com/vagrant/precise/current/precise-server-cloudimg-amd64-vagrant-disk1.box"
end
Godeps/_workspace/src/github.com/ethereum/ethash/cmake/modules/CMakeParseArguments.cmake 0 → 100644
View file @ a91bf014
#.rst:
# CMakeParseArguments
# -------------------
#
#
#
# CMAKE_PARSE_ARGUMENTS(<prefix> <options> <one_value_keywords>
# <multi_value_keywords> args...)
#
# CMAKE_PARSE_ARGUMENTS() is intended to be used in macros or functions
# for parsing the arguments given to that macro or function. It
# processes the arguments and defines a set of variables which hold the
# values of the respective options.
#
# The <options> argument contains all options for the respective macro,
# i.e. keywords which can be used when calling the macro without any
# value following, like e.g. the OPTIONAL keyword of the install()
# command.
#
# The <one_value_keywords> argument contains all keywords for this macro
# which are followed by one value, like e.g. DESTINATION keyword of the
# install() command.
#
# The <multi_value_keywords> argument contains all keywords for this
# macro which can be followed by more than one value, like e.g. the
# TARGETS or FILES keywords of the install() command.
#
# When done, CMAKE_PARSE_ARGUMENTS() will have defined for each of the
# keywords listed in <options>, <one_value_keywords> and
# <multi_value_keywords> a variable composed of the given <prefix>
# followed by "_" and the name of the respective keyword. These
# variables will then hold the respective value from the argument list.
# For the <options> keywords this will be TRUE or FALSE.
#
# All remaining arguments are collected in a variable
# <prefix>_UNPARSED_ARGUMENTS, this can be checked afterwards to see
# whether your macro was called with unrecognized parameters.
#
# As an example here a my_install() macro, which takes similar arguments
# as the real install() command:
#
# ::
#
# function(MY_INSTALL)
# set(options OPTIONAL FAST)
# set(oneValueArgs DESTINATION RENAME)
# set(multiValueArgs TARGETS CONFIGURATIONS)
# cmake_parse_arguments(MY_INSTALL "${options}" "${oneValueArgs}"
# "${multiValueArgs}" ${ARGN} )
# ...
#
#
#
# Assume my_install() has been called like this:
#
# ::
#
# my_install(TARGETS foo bar DESTINATION bin OPTIONAL blub)
#
#
#
# After the cmake_parse_arguments() call the macro will have set the
# following variables:
#
# ::
#
# MY_INSTALL_OPTIONAL = TRUE
# MY_INSTALL_FAST = FALSE (this option was not used when calling my_install()
# MY_INSTALL_DESTINATION = "bin"
# MY_INSTALL_RENAME = "" (was not used)
# MY_INSTALL_TARGETS = "foo;bar"
# MY_INSTALL_CONFIGURATIONS = "" (was not used)
# MY_INSTALL_UNPARSED_ARGUMENTS = "blub" (no value expected after "OPTIONAL"
#
#
#
# You can then continue and process these variables.
#
# Keywords terminate lists of values, e.g. if directly after a
# one_value_keyword another recognized keyword follows, this is
# interpreted as the beginning of the new option. E.g.
# my_install(TARGETS foo DESTINATION OPTIONAL) would result in
# MY_INSTALL_DESTINATION set to "OPTIONAL", but MY_INSTALL_DESTINATION
# would be empty and MY_INSTALL_OPTIONAL would be set to TRUE therefor.
#=============================================================================
# Copyright 2010 Alexander Neundorf <neundorf@kde.org>
#
# Distributed under the OSI-approved BSD License (the "License");
# see accompanying file Copyright.txt for details.
#
# This software is distributed WITHOUT ANY WARRANTY; without even the
# implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
# See the License for more information.
#=============================================================================
# (To distribute this file outside of CMake, substitute the full
# License text for the above reference.)
if(__CMAKE_PARSE_ARGUMENTS_INCLUDED)
return()
endif()
set(__CMAKE_PARSE_ARGUMENTS_INCLUDED TRUE)
function(CMAKE_PARSE_ARGUMENTS prefix _optionNames _singleArgNames _multiArgNames)
# first set all result variables to empty/FALSE
foreach(arg_name ${_singleArgNames} ${_multiArgNames})
set(${prefix}_${arg_name})
endforeach()
foreach(option ${_optionNames})
set(${prefix}_${option} FALSE)
endforeach()
set(${prefix}_UNPARSED_ARGUMENTS)
set(insideValues FALSE)
set(currentArgName)
# now iterate over all arguments and fill the result variables
foreach(currentArg ${ARGN})
list(FIND _optionNames "${currentArg}" optionIndex) # ... then this marks the end of the arguments belonging to this keyword
list(FIND _singleArgNames "${currentArg}" singleArgIndex) # ... then this marks the end of the arguments belonging to this keyword
list(FIND _multiArgNames "${currentArg}" multiArgIndex) # ... then this marks the end of the arguments belonging to this keyword
if(${optionIndex} EQUAL -1 AND ${singleArgIndex} EQUAL -1 AND ${multiArgIndex} EQUAL -1)
if(insideValues)
if("${insideValues}" STREQUAL "SINGLE")
set(${prefix}_${currentArgName} ${currentArg})
set(insideValues FALSE)
elseif("${insideValues}" STREQUAL "MULTI")
list(APPEND ${prefix}_${currentArgName} ${currentArg})
endif()
else()
list(APPEND ${prefix}_UNPARSED_ARGUMENTS ${currentArg})
endif()
else()
if(NOT ${optionIndex} EQUAL -1)
set(${prefix}_${currentArg} TRUE)
set(insideValues FALSE)
elseif(NOT ${singleArgIndex} EQUAL -1)
set(currentArgName ${currentArg})
set(${prefix}_${currentArgName})
set(insideValues "SINGLE")
elseif(NOT ${multiArgIndex} EQUAL -1)
set(currentArgName ${currentArg})
set(${prefix}_${currentArgName})
set(insideValues "MULTI")
endif()
endif()
endforeach()
# propagate the result variables to the caller:
foreach(arg_name ${_singleArgNames} ${_multiArgNames} ${_optionNames})
set(${prefix}_${arg_name} ${${prefix}_${arg_name}} PARENT_SCOPE)
endforeach()
set(${prefix}_UNPARSED_ARGUMENTS ${${prefix}_UNPARSED_ARGUMENTS} PARENT_SCOPE)
endfunction()
Godeps/_workspace/src/github.com/ethereum/ethash/cmake/modules/FindOpenCL.cmake
View file @ a91bf014
#.rst:
# FindOpenCL
# ----------
# #
# This file taken from FindOpenCL project @ http://gitorious.com/findopencl # Try to find OpenCL
# #
# - Try to find OpenCL # Once done this will define::
# This module tries to find an OpenCL implementation on your system. It supports
# AMD / ATI, Apple and NVIDIA implementations, but shoudl work, too.
# #
# Once done this will define # OpenCL_FOUND - True if OpenCL was found
# OPENCL_FOUND - system has OpenCL # OpenCL_INCLUDE_DIRS - include directories for OpenCL
# OPENCL_INCLUDE_DIRS - the OpenCL include directory # OpenCL_LIBRARIES - link against this library to use OpenCL
# OPENCL_LIBRARIES - link these to use OpenCL # OpenCL_VERSION_STRING - Highest supported OpenCL version (eg. 1.2)
# OpenCL_VERSION_MAJOR - The major version of the OpenCL implementation
# OpenCL_VERSION_MINOR - The minor version of the OpenCL implementation
#
# The module will also define two cache variables::
#
# OpenCL_INCLUDE_DIR - the OpenCL include directory
# OpenCL_LIBRARY - the path to the OpenCL library
# #
# WIN32 should work, but is untested
FIND_PACKAGE( PackageHandleStandardArgs )
SET (OPENCL_VERSION_STRING "0.1.0")
SET (OPENCL_VERSION_MAJOR 0)
SET (OPENCL_VERSION_MINOR 1)
SET (OPENCL_VERSION_PATCH 0)
IF (APPLE)
FIND_LIBRARY(OPENCL_LIBRARIES OpenCL DOC "OpenCL lib for OSX")
FIND_PATH(OPENCL_INCLUDE_DIRS OpenCL/cl.h DOC "Include for OpenCL on OSX")
FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS OpenCL/cl.hpp DOC "Include for OpenCL CPP bindings on OSX")
ELSE (APPLE)
IF (WIN32)
FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h)
FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp)
# The AMD SDK currently installs both x86 and x86_64 libraries
# This is only a hack to find out architecture
IF( ${CMAKE_SYSTEM_PROCESSOR} STREQUAL "AMD64" )
SET(OPENCL_LIB_DIR "$ENV{ATISTREAMSDKROOT}/lib/x86_64")
SET(OPENCL_LIB_DIR "$ENV{ATIINTERNALSTREAMSDKROOT}/lib/x86_64")
ELSE (${CMAKE_SYSTEM_PROCESSOR} STREQUAL "AMD64")
SET(OPENCL_LIB_DIR "$ENV{ATISTREAMSDKROOT}/lib/x86")
SET(OPENCL_LIB_DIR "$ENV{ATIINTERNALSTREAMSDKROOT}/lib/x86")
ENDIF( ${CMAKE_SYSTEM_PROCESSOR} STREQUAL "AMD64" )
# find out if the user asked for a 64-bit build, and use the corresponding
# 64 or 32 bit NVIDIA library paths to the search:
STRING(REGEX MATCH "Win64" ISWIN64 ${CMAKE_GENERATOR})
IF("${ISWIN64}" STREQUAL "Win64")
FIND_LIBRARY(OPENCL_LIBRARIES OpenCL.lib ${OPENCL_LIB_DIR} $ENV{CUDA_LIB_PATH} $ENV{CUDA_PATH}/lib/x64)
ELSE("${ISWIN64}" STREQUAL "Win64")
FIND_LIBRARY(OPENCL_LIBRARIES OpenCL.lib ${OPENCL_LIB_DIR} $ENV{CUDA_LIB_PATH} $ENV{CUDA_PATH}/lib/Win32)
ENDIF("${ISWIN64}" STREQUAL "Win64")
GET_FILENAME_COMPONENT(_OPENCL_INC_CAND ${OPENCL_LIB_DIR}/../../include ABSOLUTE)
# On Win32 search relative to the library
FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h PATHS "${_OPENCL_INC_CAND}" $ENV{CUDA_INC_PATH} $ENV{CUDA_PATH}/include)
FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp PATHS "${_OPENCL_INC_CAND}" $ENV{CUDA_INC_PATH} $ENV{CUDA_PATH}/include)
ELSE (WIN32)
# Unix style platforms
FIND_LIBRARY(OPENCL_LIBRARIES OpenCL
ENV LD_LIBRARY_PATH
)
GET_FILENAME_COMPONENT(OPENCL_LIB_DIR ${OPENCL_LIBRARIES} PATH)
GET_FILENAME_COMPONENT(_OPENCL_INC_CAND ${OPENCL_LIB_DIR}/../../include ABSOLUTE)
# The AMD SDK currently does not place its headers
# in /usr/include, therefore also search relative
# to the library
FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h PATHS ${_OPENCL_INC_CAND} "/usr/local/cuda/include")
FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp PATHS ${_OPENCL_INC_CAND} "/usr/local/cuda/include")
ENDIF (WIN32)
ENDIF (APPLE)
FIND_PACKAGE_HANDLE_STANDARD_ARGS( OpenCL DEFAULT_MSG OPENCL_LIBRARIES OPENCL_INCLUDE_DIRS )
IF( _OPENCL_CPP_INCLUDE_DIRS )
SET( OPENCL_HAS_CPP_BINDINGS TRUE )
LIST( APPEND OPENCL_INCLUDE_DIRS ${_OPENCL_CPP_INCLUDE_DIRS} )
# This is often the same, so clean up
LIST( REMOVE_DUPLICATES OPENCL_INCLUDE_DIRS )
ENDIF( _OPENCL_CPP_INCLUDE_DIRS )
MARK_AS_ADVANCED( #=============================================================================
OPENCL_INCLUDE_DIRS # Copyright 2014 Matthaeus G. Chajdas
) #
\ No newline at end of file # Distributed under the OSI-approved BSD License (the "License");
# see accompanying file Copyright.txt for details.
#
# This software is distributed WITHOUT ANY WARRANTY; without even the
# implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
# See the License for more information.
#=============================================================================
# (To distribute this file outside of CMake, substitute the full
# License text for the above reference.)
function(_FIND_OPENCL_VERSION)
include(CheckSymbolExists)
include(CMakePushCheckState)
set(CMAKE_REQUIRED_QUIET ${OpenCL_FIND_QUIETLY})
CMAKE_PUSH_CHECK_STATE()
foreach(VERSION "2_0" "1_2" "1_1" "1_0")
set(CMAKE_REQUIRED_INCLUDES "${OpenCL_INCLUDE_DIR}")
if(APPLE)
CHECK_SYMBOL_EXISTS(
CL_VERSION_${VERSION}
"${OpenCL_INCLUDE_DIR}/OpenCL/cl.h"
OPENCL_VERSION_${VERSION})
else()
CHECK_SYMBOL_EXISTS(
CL_VERSION_${VERSION}
"${OpenCL_INCLUDE_DIR}/CL/cl.h"
OPENCL_VERSION_${VERSION})
endif()
if(OPENCL_VERSION_${VERSION})
string(REPLACE "_" "." VERSION "${VERSION}")
set(OpenCL_VERSION_STRING ${VERSION} PARENT_SCOPE)
string(REGEX MATCHALL "[0-9]+" version_components "${VERSION}")
list(GET version_components 0 major_version)
list(GET version_components 1 minor_version)
set(OpenCL_VERSION_MAJOR ${major_version} PARENT_SCOPE)
set(OpenCL_VERSION_MINOR ${minor_version} PARENT_SCOPE)
break()
endif()
endforeach()
CMAKE_POP_CHECK_STATE()
endfunction()
find_path(OpenCL_INCLUDE_DIR
NAMES
CL/cl.h OpenCL/cl.h
PATHS
ENV "PROGRAMFILES(X86)"
ENV AMDAPPSDKROOT
ENV INTELOCLSDKROOT
ENV NVSDKCOMPUTE_ROOT
ENV CUDA_PATH
ENV ATISTREAMSDKROOT
PATH_SUFFIXES
include
OpenCL/common/inc
"AMD APP/include")
_FIND_OPENCL_VERSION()
if(WIN32)
if(CMAKE_SIZEOF_VOID_P EQUAL 4)
find_library(OpenCL_LIBRARY
NAMES OpenCL
PATHS
ENV "PROGRAMFILES(X86)"
ENV AMDAPPSDKROOT
ENV INTELOCLSDKROOT
ENV CUDA_PATH
ENV NVSDKCOMPUTE_ROOT
ENV ATISTREAMSDKROOT
PATH_SUFFIXES
"AMD APP/lib/x86"
lib/x86
lib/Win32
OpenCL/common/lib/Win32)
elseif(CMAKE_SIZEOF_VOID_P EQUAL 8)
find_library(OpenCL_LIBRARY
NAMES OpenCL
PATHS
ENV "PROGRAMFILES(X86)"
ENV AMDAPPSDKROOT
ENV INTELOCLSDKROOT
ENV CUDA_PATH
ENV NVSDKCOMPUTE_ROOT
ENV ATISTREAMSDKROOT
PATH_SUFFIXES
"AMD APP/lib/x86_64"
lib/x86_64
lib/x64
OpenCL/common/lib/x64)
endif()
else()
find_library(OpenCL_LIBRARY
NAMES OpenCL)
endif()
set(OpenCL_LIBRARIES ${OpenCL_LIBRARY})
set(OpenCL_INCLUDE_DIRS ${OpenCL_INCLUDE_DIR})
include(${CMAKE_CURRENT_LIST_DIR}/FindPackageHandleStandardArgs.cmake)
find_package_handle_standard_args(
OpenCL
FOUND_VAR OpenCL_FOUND
REQUIRED_VARS OpenCL_LIBRARY OpenCL_INCLUDE_DIR
VERSION_VAR OpenCL_VERSION_STRING)
mark_as_advanced(
OpenCL_INCLUDE_DIR
OpenCL_LIBRARY)
Godeps/_workspace/src/github.com/ethereum/ethash/cmake/modules/FindPackageMessage.cmake 0 → 100644
View file @ a91bf014
#.rst:
# FindPackageMessage
# ------------------
#
#
#
# FIND_PACKAGE_MESSAGE(<name> "message for user" "find result details")
#
# This macro is intended to be used in FindXXX.cmake modules files. It
# will print a message once for each unique find result. This is useful
# for telling the user where a package was found. The first argument
# specifies the name (XXX) of the package. The second argument
# specifies the message to display. The third argument lists details
# about the find result so that if they change the message will be
# displayed again. The macro also obeys the QUIET argument to the
# find_package command.
#
# Example:
#
# ::
#
# if(X11_FOUND)
# FIND_PACKAGE_MESSAGE(X11 "Found X11: ${X11_X11_LIB}"
# "[${X11_X11_LIB}][${X11_INCLUDE_DIR}]")
# else()
# ...
# endif()
#=============================================================================
# Copyright 2008-2009 Kitware, Inc.
#
# Distributed under the OSI-approved BSD License (the "License");
# see accompanying file Copyright.txt for details.
#
# This software is distributed WITHOUT ANY WARRANTY; without even the
# implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
# See the License for more information.
#=============================================================================
# (To distribute this file outside of CMake, substitute the full
# License text for the above reference.)
function(FIND_PACKAGE_MESSAGE pkg msg details)
# Avoid printing a message repeatedly for the same find result.
if(NOT ${pkg}_FIND_QUIETLY)
string(REPLACE "\n" "" details "${details}")
set(DETAILS_VAR FIND_PACKAGE_MESSAGE_DETAILS_${pkg})
if(NOT "${details}" STREQUAL "${${DETAILS_VAR}}")
# The message has not yet been printed.
message(STATUS "${msg}")
# Save the find details in the cache to avoid printing the same
# message again.
set("${DETAILS_VAR}" "${details}"
CACHE INTERNAL "Details about finding ${pkg}")
endif()
endif()
endfunction()
Godeps/_workspace/src/github.com/ethereum/ethash/js/LICENSE 0 → 100644
View file @ a91bf014
The MIT License (MIT)
Copyright (c) 2015 Tim Hughes
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.
Godeps/_workspace/src/github.com/ethereum/ethash/js/ethash.js 0 → 100644
View file @ a91bf014
// ethash.js
// Tim Hughes <tim@twistedfury.com>
// Revision 19
/*jslint node: true, shadow:true */
"use strict";
var Keccak = require('./keccak');
var util = require('./util');
// 32-bit unsigned modulo
function mod32(x, n)
{
return (x>>>0) % (n>>>0);
}
function fnv(x, y)
{
// js integer multiply by 0x01000193 will lose precision
return ((x*0x01000000 | 0) + (x*0x193 | 0)) ^ y;
}
function computeCache(params, seedWords)
{
var cache = new Uint32Array(params.cacheSize >> 2);
var cacheNodeCount = params.cacheSize >> 6;
// Initialize cache
var keccak = new Keccak();
keccak.digestWords(cache, 0, 16, seedWords, 0, seedWords.length);
for (var n = 1; n < cacheNodeCount; ++n)
{
keccak.digestWords(cache, n<<4, 16, cache, (n-1)<<4, 16);
}
var tmp = new Uint32Array(16);
// Do randmemohash passes
for (var r = 0; r < params.cacheRounds; ++r)
{
for (var n = 0; n < cacheNodeCount; ++n)
{
var p0 = mod32(n + cacheNodeCount - 1, cacheNodeCount) << 4;
var p1 = mod32(cache[n<<4|0], cacheNodeCount) << 4;
for (var w = 0; w < 16; w=(w+1)|0)
{
tmp[w] = cache[p0 | w] ^ cache[p1 | w];
}
keccak.digestWords(cache, n<<4, 16, tmp, 0, tmp.length);
}
}
return cache;
}
function computeDagNode(o_node, params, cache, keccak, nodeIndex)
{
var cacheNodeCount = params.cacheSize >> 6;
var dagParents = params.dagParents;
var c = (nodeIndex % cacheNodeCount) << 4;
var mix = o_node;
for (var w = 0; w < 16; ++w)
{
mix[w] = cache[c|w];
}
mix[0] ^= nodeIndex;
keccak.digestWords(mix, 0, 16, mix, 0, 16);
for (var p = 0; p < dagParents; ++p)
{
// compute cache node (word) index
c = mod32(fnv(nodeIndex ^ p, mix[p&15]), cacheNodeCount) << 4;
for (var w = 0; w < 16; ++w)
{
mix[w] = fnv(mix[w], cache[c|w]);
}
}
keccak.digestWords(mix, 0, 16, mix, 0, 16);
}
function computeHashInner(mix, params, cache, keccak, tempNode)
{
var mixParents = params.mixParents|0;
var mixWordCount = params.mixSize >> 2;
var mixNodeCount = mixWordCount >> 4;
var dagPageCount = (params.dagSize / params.mixSize) >> 0;
// grab initial first word
var s0 = mix[0];
// initialise mix from initial 64 bytes
for (var w = 16; w < mixWordCount; ++w)
{
mix[w] = mix[w & 15];
}
for (var a = 0; a < mixParents; ++a)
{
var p = mod32(fnv(s0 ^ a, mix[a & (mixWordCount-1)]), dagPageCount);
var d = (p * mixNodeCount)|0;
for (var n = 0, w = 0; n < mixNodeCount; ++n, w += 16)
{
computeDagNode(tempNode, params, cache, keccak, (d + n)|0);
for (var v = 0; v < 16; ++v)
{
mix[w|v] = fnv(mix[w|v], tempNode[v]);
}
}
}
}
function convertSeed(seed)
{
// todo, reconcile with spec, byte ordering?
// todo, big-endian conversion
var newSeed = util.toWords(seed);
if (newSeed === null)
throw Error("Invalid seed '" + seed + "'");
return newSeed;
}
exports.defaultParams = function()
{
return {
cacheSize: 1048384,
cacheRounds: 3,
dagSize: 1073739904,
dagParents: 256,
mixSize: 128,
mixParents: 64,
};
};
exports.Ethash = function(params, seed)
{
// precompute cache and related values
seed = convertSeed(seed);
var cache = computeCache(params, seed);
// preallocate buffers/etc
var initBuf = new ArrayBuffer(96);
var initBytes = new Uint8Array(initBuf);
var initWords = new Uint32Array(initBuf);
var mixWords = new Uint32Array(params.mixSize / 4);
var tempNode = new Uint32Array(16);
var keccak = new Keccak();
var retWords = new Uint32Array(8);
var retBytes = new Uint8Array(retWords.buffer); // supposedly read-only
this.hash = function(header, nonce)
{
// compute initial hash
initBytes.set(header, 0);
initBytes.set(nonce, 32);
keccak.digestWords(initWords, 0, 16, initWords, 0, 8 + nonce.length/4);
// compute mix
for (var i = 0; i != 16; ++i)
{
mixWords[i] = initWords[i];
}
computeHashInner(mixWords, params, cache, keccak, tempNode);
// compress mix and append to initWords
for (var i = 0; i != mixWords.length; i += 4)
{
initWords[16 + i/4] = fnv(fnv(fnv(mixWords[i], mixWords[i+1]), mixWords[i+2]), mixWords[i+3]);
}
// final Keccak hashes
keccak.digestWords(retWords, 0, 8, initWords, 0, 24); // Keccak-256(s + cmix)
return retBytes;
};
this.cacheDigest = function()
{
return keccak.digest(32, new Uint8Array(cache.buffer));
};
};
Godeps/_workspace/src/github.com/ethereum/ethash/js/keccak.js 0 → 100644
View file @ a91bf014
This diff is collapsed.
Godeps/_workspace/src/github.com/ethereum/ethash/js/makekeccak.js 0 → 100644
View file @ a91bf014
#!/usr/bin/env node
// makekeccak.js
// Tim Hughes <tim@twistedfury.com>
/*jslint node: true, shadow:true */
"use strict";
var Keccak_f1600_Rho = [
1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14,
27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44
];
var Keccak_f1600_Pi= [
10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4,
15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1
];
var Keccak_f1600_RC = [
0x00000001, 0x00000000,
0x00008082, 0x00000000,
0x0000808a, 0x80000000,
0x80008000, 0x80000000,
0x0000808b, 0x00000000,
0x80000001, 0x00000000,
0x80008081, 0x80000000,
0x00008009, 0x80000000,
0x0000008a, 0x00000000,
0x00000088, 0x00000000,
0x80008009, 0x00000000,
0x8000000a, 0x00000000,
0x8000808b, 0x00000000,
0x0000008b, 0x80000000,
0x00008089, 0x80000000,
0x00008003, 0x80000000,
0x00008002, 0x80000000,
0x00000080, 0x80000000,
0x0000800a, 0x00000000,
0x8000000a, 0x80000000,
0x80008081, 0x80000000,
0x00008080, 0x80000000,
0x80000001, 0x00000000,
0x80008008, 0x80000000,
];
function makeRotLow(lo, hi, n)
{
if (n === 0 || n === 32) throw Error("unsupported");
if ((n & 0x20) !== 0)
{
n &= ~0x20;
var t = hi;
hi = lo;
lo = t;
}
var hir = hi + " >>> " + (32 - n);
var los = lo + " << " + n;
return los + " | " + hir;
}
function makeRotHigh(lo, hi, n)
{
if (n === 0 || n === 32) throw Error("unsupported");
if ((n & 0x20) !== 0)
{
n &= ~0x20;
var t = hi;
hi = lo;
lo = t;
}
var his = hi + " << " + n;
var lor = lo + " >>> " + (32 - n);
return his + " | " + lor;
}
function makeKeccak_f1600()
{
var format = function(n)
{
return n < 10 ? "0"+n : ""+n;
};
var a = function(n, w)
{
return "a" + format(n) + (w !== 0?'h':'l');
};
var b = function(n, w)
{
return "b" + format(n) + (w !== 0?'h':'l');
};
var str = "";
str += "function keccak_f1600(outState, outOffset, outSize, inState)\n";
str += "{\n";
for (var i = 0; i < 25; ++i)
{
for (var w = 0; w <= 1; ++w)
{
str += "\tvar " + a(i,w) + " = inState["+(i<<1|w)+"]|0;\n";
}
}
for (var j = 0; j < 5; ++j)
{
str += "\tvar ";
for (var i = 0; i < 5; ++i)
{
if (i !== 0)
str += ", ";
str += b(j*5+i,0) + ", " + b(j*5+i,1);
}
str += ";\n";
}
str += "\tvar tl, th;\n";
str += "\n";
str += "\tfor (var r = 0; r < 48; r = (r+2)|0)\n";
str += "\t{\n";
// Theta
str += "\t\t// Theta\n";
for (var i = 0; i < 5; ++i)
{
for (var w = 0; w <= 1; ++w)
{
str += "\t\t" + b(i,w) + " = " + a(i,w) + " ^ " + a(i+5,w) + " ^ " + a(i+10,w) + " ^ " + a(i+15,w) + " ^ " + a(i+20,w) + ";\n";
}
}
for (var i = 0; i < 5; ++i)
{
var i4 = (i + 4) % 5;
var i1 = (i + 1) % 5;
str += "\t\ttl = " + b(i4,0) + " ^ (" + b(i1,0) + " << 1 | " + b(i1,1) + " >>> 31);\n";
str += "\t\tth = " + b(i4,1) + " ^ (" + b(i1,1) + " << 1 | " + b(i1,0) + " >>> 31);\n";
for (var j = 0; j < 25; j = (j+5)|0)
{
str += "\t\t" + a((j+i),0) + " ^= tl;\n";
str += "\t\t" + a((j+i),1) + " ^= th;\n";
}
}
// Rho Pi
str += "\n\t\t// Rho Pi\n";
for (var w = 0; w <= 1; ++w)
{
str += "\t\t" + b(0,w) + " = " + a(0,w) + ";\n";
}
var opi = 1;
for (var i = 0; i < 24; ++i)
{
var pi = Keccak_f1600_Pi[i];
str += "\t\t" + b(pi,0) + " = " + makeRotLow(a(opi,0), a(opi,1), Keccak_f1600_Rho[i]) + ";\n";
str += "\t\t" + b(pi,1) + " = " + makeRotHigh(a(opi,0), a(opi,1), Keccak_f1600_Rho[i]) + ";\n";
opi = pi;
}
// Chi
str += "\n\t\t// Chi\n";
for (var j = 0; j < 25; j += 5)
{
for (var i = 0; i < 5; ++i)
{
for (var w = 0; w <= 1; ++w)
{
str += "\t\t" + a(j+i,w) + " = " + b(j+i,w) + " ^ ~" + b(j+(i+1)%5,w) + " & " + b(j+(i+2)%5,w) + ";\n";
}
}
}
// Iota
str += "\n\t\t// Iota\n";
for (var w = 0; w <= 1; ++w)
{
str += "\t\t" + a(0,w) + " ^= Keccak_f1600_RC[r|" + w + "];\n";
}
str += "\t}\n";
for (var i = 0; i < 25; ++i)
{
if (i == 4 || i == 8)
{
str += "\tif (outSize == " + i*2 + ")\n\t\treturn;\n";
}
for (var w = 0; w <= 1; ++w)
{
str += "\toutState[outOffset|"+(i<<1|w)+"] = " + a(i,w) + ";\n";
}
}
str += "}\n";
return str;
}
console.log(makeKeccak_f1600());
Godeps/_workspace/src/github.com/ethereum/ethash/js/test.js 0 → 100644
View file @ a91bf014
// test.js
// Tim Hughes <tim@twistedfury.com>
/*jslint node: true, shadow:true */
"use strict";
var ethash = require('./ethash');
var util = require('./util');
var Keccak = require('./keccak');
// sanity check hash functions
var src = util.stringToBytes("");
if (util.bytesToHexString(new Keccak().digest(32, src)) != "c5d2460186f7233c927e7db2dcc703c0e500b653ca82273b7bfad8045d85a470") throw Error("Keccak-256 failed");
if (util.bytesToHexString(new Keccak().digest(64, src)) != "0eab42de4c3ceb9235fc91acffe746b29c29a8c366b7c60e4e67c466f36a4304c00fa9caf9d87976ba469bcbe06713b435f091ef2769fb160cdab33d3670680e") throw Error("Keccak-512 failed");
src = new Uint32Array(src.buffer);
var dst = new Uint32Array(8);
new Keccak().digestWords(dst, 0, dst.length, src, 0, src.length);
if (util.wordsToHexString(dst) != "c5d2460186f7233c927e7db2dcc703c0e500b653ca82273b7bfad8045d85a470") throw Error("Keccak-256 Fast failed");
var dst = new Uint32Array(16);
new Keccak().digestWords(dst, 0, dst.length, src, 0, src.length);
if (util.wordsToHexString(dst) != "0eab42de4c3ceb9235fc91acffe746b29c29a8c366b7c60e4e67c466f36a4304c00fa9caf9d87976ba469bcbe06713b435f091ef2769fb160cdab33d3670680e") throw Error("Keccak-512 Fast failed");
// init params
var ethashParams = ethash.defaultParams();
//ethashParams.cacheRounds = 0;
// create hasher
var seed = util.hexStringToBytes("9410b944535a83d9adf6bbdcc80e051f30676173c16ca0d32d6f1263fc246466")
var startTime = new Date().getTime();
var hasher = new ethash.Ethash(ethashParams, seed);
console.log('Ethash startup took: '+(new Date().getTime() - startTime) + "ms");
console.log('Ethash cache hash: ' + util.bytesToHexString(hasher.cacheDigest()));
var testHexString = "c5d2460186f7233c927e7db2dcc703c0e500b653ca82273b7bfad8045d85a470";
if (testHexString != util.bytesToHexString(util.hexStringToBytes(testHexString)))
throw Error("bytesToHexString or hexStringToBytes broken");
var header = util.hexStringToBytes("c5d2460186f7233c927e7db2dcc703c0e500b653ca82273b7bfad8045d85a470");
var nonce = util.hexStringToBytes("0000000000000000");
var hash;
startTime = new Date().getTime();
var trials = 10;
for (var i = 0; i < trials; ++i)
{
hash = hasher.hash(header, nonce);
}
console.log("Light client hashes averaged: " + (new Date().getTime() - startTime)/trials + "ms");
console.log("Hash = " + util.bytesToHexString(hash));
Godeps/_workspace/src/github.com/ethereum/ethash/js/util.js 0 → 100644
View file @ a91bf014
// util.js
// Tim Hughes <tim@twistedfury.com>
/*jslint node: true, shadow:true */
"use strict";
function nibbleToChar(nibble)
{
return String.fromCharCode((nibble < 10 ? 48 : 87) + nibble);
}
function charToNibble(chr)
{
if (chr >= 48 && chr <= 57)
{
return chr - 48;
}
if (chr >= 65 && chr <= 70)
{
return chr - 65 + 10;
}
if (chr >= 97 && chr <= 102)
{
return chr - 97 + 10;
}
return 0;
}
function stringToBytes(str)
{
var bytes = new Uint8Array(str.length);
for (var i = 0; i != str.length; ++i)
{
bytes[i] = str.charCodeAt(i);
}
return bytes;
}
function hexStringToBytes(str)
{
var bytes = new Uint8Array(str.length>>>1);
for (var i = 0; i != bytes.length; ++i)
{
bytes[i] = charToNibble(str.charCodeAt(i<<1 | 0)) << 4;
bytes[i] |= charToNibble(str.charCodeAt(i<<1 | 1));
}
return bytes;
}
function bytesToHexString(bytes)
{
var str = "";
for (var i = 0; i != bytes.length; ++i)
{
str += nibbleToChar(bytes[i] >>> 4);
str += nibbleToChar(bytes[i] & 0xf);
}
return str;
}
function wordsToHexString(words)
{
return bytesToHexString(new Uint8Array(words.buffer));
}
function uint32ToHexString(num)
{
var buf = new Uint8Array(4);
buf[0] = (num >> 24) & 0xff;
buf[1] = (num >> 16) & 0xff;
buf[2] = (num >> 8) & 0xff;
buf[3] = (num >> 0) & 0xff;
return bytesToHexString(buf);
}
function toWords(input)
{
if (input instanceof Uint32Array)
{
return input;
}
else if (input instanceof Uint8Array)
{
var tmp = new Uint8Array((input.length + 3) & ~3);
tmp.set(input);
return new Uint32Array(tmp.buffer);
}
else if (typeof input === typeof "")
{
return toWords(stringToBytes(input));
}
return null;
}
exports.stringToBytes = stringToBytes;
exports.hexStringToBytes = hexStringToBytes;
exports.bytesToHexString = bytesToHexString;
exports.wordsToHexString = wordsToHexString;
exports.uint32ToHexString = uint32ToHexString;
exports.toWords = toWords;
\ No newline at end of file
Godeps/_workspace/src/github.com/ethereum/ethash/libethash-cuda/CMakeLists.txt deleted 100644 → 0
View file @ 76e4e233
find_package(CUDA)
# Pass options to NVCC
if (CUDA_FOUND)
set(CUDA_NVCC_FLAGS " -gencode;arch=compute_30,code=sm_30;
-gencode;arch=compute_20,code=sm_20;
-gencode;arch=compute_11,code=sm_11;
-gencode;arch=compute_12,code=sm_12;
-gencode;arch=compute_13,code=sm_13;")
cuda_add_executable(
ethash-cuda
libethash.cu)
endif()
\ No newline at end of file
Godeps/_workspace/src/github.com/ethereum/ethash/libethash-cuda/cuPrintf.cuh deleted 100644 → 0
View file @ 76e4e233
/*
Copyright 2009 NVIDIA Corporation. All rights reserved.
NOTICE TO LICENSEE:
This source code and/or documentation ("Licensed Deliverables") are subject
to NVIDIA intellectual property rights under U.S. and international Copyright
laws.
These Licensed Deliverables contained herein is PROPRIETARY and CONFIDENTIAL
to NVIDIA and is being provided under the terms and conditions of a form of
NVIDIA software license agreement by and between NVIDIA and Licensee ("License
Agreement") or electronically accepted by Licensee. Notwithstanding any terms
or conditions to the contrary in the License Agreement, reproduction or
disclosure of the Licensed Deliverables to any third party without the express
written consent of NVIDIA is prohibited.
NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE LICENSE AGREEMENT,
NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THESE LICENSED
DELIVERABLES FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED
WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE
LICENSED DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. NOTWITHSTANDING ANY
TERMS OR CONDITIONS TO THE CONTRARY IN THE LICENSE AGREEMENT, IN NO EVENT SHALL
NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES,
OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS, WHETHER
IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS ACTION, ARISING OUT OF
OR IN CONNECTION WITH THE USE OR PERFORMANCE OF THESE LICENSED DELIVERABLES.
U.S. Government End Users. These Licensed Deliverables are a "commercial item"
as that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of
"commercial computer software" and "commercial computer software documentation"
as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) and is provided to the
U.S. Government only as a commercial end item. Consistent with 48 C.F.R.12.212
and 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all U.S. Government
End Users acquire the Licensed Deliverables with only those rights set forth
herein.
Any use of the Licensed Deliverables in individual and commercial software must
include, in the user documentation and internal comments to the code, the above
Disclaimer and U.S. Government End Users Notice.
*/
#ifndef CUPRINTF_H
#define CUPRINTF_H
/*
* This is the header file supporting cuPrintf.cu and defining both
* the host and device-side interfaces. See that file for some more
* explanation and sample use code. See also below for details of the
* host-side interfaces.
*
* Quick sample code:
*
#include "cuPrintf.cu"
__global__ void testKernel(int val)
{
cuPrintf("Value is: %d\n", val);
}
int main()
{
cudaPrintfInit();
testKernel<<< 2, 3 >>>(10);
cudaPrintfDisplay(stdout, true);
cudaPrintfEnd();
return 0;
}
*/
///////////////////////////////////////////////////////////////////////////////
// DEVICE SIDE
// External function definitions for device-side code
// Abuse of templates to simulate varargs
__device__ int cuPrintf(const char *fmt);
template <typename T1> __device__ int cuPrintf(const char *fmt, T1 arg1);
template <typename T1, typename T2> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2);
template <typename T1, typename T2, typename T3> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3);
template <typename T1, typename T2, typename T3, typename T4> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4);
template <typename T1, typename T2, typename T3, typename T4, typename T5> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9);
template <typename T1, typename T2, typename T3, typename T4, typename T5, typename T6, typename T7, typename T8, typename T9, typename T10> __device__ int cuPrintf(const char *fmt, T1 arg1, T2 arg2, T3 arg3, T4 arg4, T5 arg5, T6 arg6, T7 arg7, T8 arg8, T9 arg9, T10 arg10);
//
// cuPrintfRestrict
//
// Called to restrict output to a given thread/block. Pass
// the constant CUPRINTF_UNRESTRICTED to unrestrict output
// for thread/block IDs. Note you can therefore allow
// "all printfs from block 3" or "printfs from thread 2
// on all blocks", or "printfs only from block 1, thread 5".
//
// Arguments:
// threadid - Thread ID to allow printfs from
// blockid - Block ID to allow printfs from
//
// NOTE: Restrictions last between invocations of
// kernels unless cudaPrintfInit() is called again.
//
#define CUPRINTF_UNRESTRICTED -1
__device__ void cuPrintfRestrict(int threadid, int blockid);
///////////////////////////////////////////////////////////////////////////////
// HOST SIDE
// External function definitions for host-side code
//
// cudaPrintfInit
//
// Call this once to initialise the printf system. If the output
// file or buffer size needs to be changed, call cudaPrintfEnd()
// before re-calling cudaPrintfInit().
//
// The default size for the buffer is 1 megabyte. For CUDA
// architecture 1.1 and above, the buffer is filled linearly and
// is completely used; however for architecture 1.0, the buffer
// is divided into as many segments are there are threads, even
// if some threads do not call cuPrintf().
//
// Arguments:
// bufferLen - Length, in bytes, of total space to reserve
// (in device global memory) for output.
//
// Returns:
// cudaSuccess if all is well.
//
extern "C" cudaError_t cudaPrintfInit(size_t bufferLen=1048576); // 1-meg - that's enough for 4096 printfs by all threads put together
//
// cudaPrintfEnd
//
// Cleans up all memories allocated by cudaPrintfInit().
// Call this at exit, or before calling cudaPrintfInit() again.
//
extern "C" void cudaPrintfEnd();
//
// cudaPrintfDisplay
//
// Dumps the contents of the output buffer to the specified
// file pointer. If the output pointer is not specified,
// the default "stdout" is used.
//
// Arguments:
// outputFP - A file pointer to an output stream.
// showThreadID - If "true", output strings are prefixed
// by "[blockid, threadid] " at output.
//
// Returns:
// cudaSuccess if all is well.
//
extern "C" cudaError_t cudaPrintfDisplay(void *outputFP=NULL, bool showThreadID=false);
#endif // CUPRINTF_H
Godeps/_workspace/src/github.com/ethereum/ethash/libethash-cuda/libethash.cu deleted 100644 → 0
View file @ 76e4e233
#include "cuPrintf.cu"
#include <stdio.h>
__global__ void device_greetings(void)
{
cuPrintf("Hello, world from the device!\n");
}
int main(void)
{
// greet from the host
printf("Hello, world from the host!\n");
// initialize cuPrintf
cudaPrintfInit();
// launch a kernel with a single thread to greet from the device
device_greetings<<<1,1>>>();
// display the device's greeting
cudaPrintfDisplay();
// clean up after cuPrintf
cudaPrintfEnd();
return 0;
}
Godeps/_workspace/src/github.com/ethereum/ethash/node-ethash/binding.gyp deleted 100644 → 0
View file @ 76e4e233
{
"targets":
[{
"target_name": "ethash",
"sources": [
'./ethash.cc',
'../libethash/ethash.h',
'../libethash/util.c',
'../libethash/util.h',
'../libethash/blum_blum_shub.h',
'../libethash/blum_blum_shub.c',
'../libethash/sha3.h',
'../libethash/sha3.c',
'../libethash/internal.h',
'../libethash/internal.c'
],
"include_dirs": [
"../",
"<!(node -e \"require('nan')\")"
],
"cflags": [
"-Wall",
"-Wno-maybe-uninitialized",
"-Wno-uninitialized",
"-Wno-unused-function",
"-Wextra"
]
}]
}
Godeps/_workspace/src/github.com/ethereum/ethash/node-ethash/package.json deleted 100644 → 0
View file @ 76e4e233
{
"name": "node-ethash",
"version": "1.0.0",
"description": "",
"main": "index.js",
"scripts": {
"test": "echo \"Error: no test specified\" && exit 1",
"install": "node-gyp rebuild"
},
"author": "",
"license": "ISC",
"gypfile": true
}
Godeps/_workspace/src/github.com/ethereum/ethash/node-ethash/readme.md deleted 100644 → 0
View file @ 76e4e233
# To Develop
`npm install -g node-gyp`
`npm install .`
# To rebuild
`node-gyp rebuild`
# notes
nan is good https://github.com/rvagg/nan
Godeps/_workspace/src/github.com/ethereum/ethash/pyethash/__init__.py 0 → 100644
View file @ a91bf014
import pyethash.core
core = pyethash.core
EPOCH_LENGTH = 30000
Godeps/_workspace/src/github.com/ethereum/ethash/setup.py 0 → 100644
View file @ a91bf014
#!/usr/bin/env python
from distutils.core import setup, Extension
pyethash_core = Extension('pyethash.core',
sources = [
'src/python/core.c',
'src/libethash/util.c',
'src/libethash/internal.c',
'src/libethash/sha3.c'
],
extra_compile_args = ["-std=gnu99"])
setup (
name = 'pyethash',
author = "Matthew Wampler-Doty",
author_email = "matthew.wampler.doty@gmail.com",
license = 'GPL',
version = '1.0',
description = 'Python wrappers for ethash, the ethereum proof of work hashing function',
ext_modules = [pyethash_core],
)
Godeps/_workspace/src/github.com/ethereum/ethash/libethash-cl/CMakeLists.txt → Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/CMakeLists.txt
View file @ a91bf014
set(LIBRARY ethash-cl) set(LIBRARY ethash-cl)
set(CMAKE_BUILD_TYPE Release) set(CMAKE_BUILD_TYPE Release)
if (NOT OPENCL_FOUND) 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)
if (NOT OpenCL_FOUND)
find_package(OpenCL) find_package(OpenCL)
endif() endif()
if (OPENCL_FOUND) if (OpenCL_FOUND)
include_directories(${OPENCL_INCLUDE_DIRS}) 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)
TARGET_LINK_LIBRARIES(${LIBRARY} ${OPENCL_LIBRARIES} ethash) TARGET_LINK_LIBRARIES(${LIBRARY} ${OpenCL_LIBRARIES} ethash)
endif() endif()
\ No newline at end of file
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/bin2h.cmake 0 → 100644
View file @ a91bf014
# https://gist.github.com/sivachandran/3a0de157dccef822a230
include(CMakeParseArguments)
# Function to wrap a given string into multiple lines at the given column position.
# Parameters:
# VARIABLE - The name of the CMake variable holding the string.
# AT_COLUMN - The column position at which string will be wrapped.
function(WRAP_STRING)
set(oneValueArgs VARIABLE AT_COLUMN)
cmake_parse_arguments(WRAP_STRING "${options}" "${oneValueArgs}" "" ${ARGN})
string(LENGTH ${${WRAP_STRING_VARIABLE}} stringLength)
math(EXPR offset "0")
while(stringLength GREATER 0)
if(stringLength GREATER ${WRAP_STRING_AT_COLUMN})
math(EXPR length "${WRAP_STRING_AT_COLUMN}")
else()
math(EXPR length "${stringLength}")
endif()
string(SUBSTRING ${${WRAP_STRING_VARIABLE}} ${offset} ${length} line)
set(lines "${lines}\n${line}")
math(EXPR stringLength "${stringLength} - ${length}")
math(EXPR offset "${offset} + ${length}")
endwhile()
set(${WRAP_STRING_VARIABLE} "${lines}" PARENT_SCOPE)
endfunction()
# Function to embed contents of a file as byte array in C/C++ header file(.h). The header file
# will contain a byte array and integer variable holding the size of the array.
# Parameters
# SOURCE_FILE - The path of source file whose contents will be embedded in the header file.
# VARIABLE_NAME - The name of the variable for the byte array. The string "_SIZE" will be append
# to this name and will be used a variable name for size variable.
# HEADER_FILE - The path of header file.
# APPEND - If specified appends to the header file instead of overwriting it
# NULL_TERMINATE - If specified a null byte(zero) will be append to the byte array. This will be
# useful if the source file is a text file and we want to use the file contents
# as string. But the size variable holds size of the byte array without this
# null byte.
# Usage:
# bin2h(SOURCE_FILE "Logo.png" HEADER_FILE "Logo.h" VARIABLE_NAME "LOGO_PNG")
function(BIN2H)
set(options APPEND NULL_TERMINATE)
set(oneValueArgs SOURCE_FILE VARIABLE_NAME HEADER_FILE)
cmake_parse_arguments(BIN2H "${options}" "${oneValueArgs}" "" ${ARGN})
# reads source file contents as hex string
file(READ ${BIN2H_SOURCE_FILE} hexString HEX)
string(LENGTH ${hexString} hexStringLength)
# appends null byte if asked
if(BIN2H_NULL_TERMINATE)
set(hexString "${hexString}00")
endif()
# wraps the hex string into multiple lines at column 32(i.e. 16 bytes per line)
wrap_string(VARIABLE hexString AT_COLUMN 32)
math(EXPR arraySize "${hexStringLength} / 2")
# adds '0x' prefix and comma suffix before and after every byte respectively
string(REGEX REPLACE "([0-9a-f][0-9a-f])" "0x\\1, " arrayValues ${hexString})
# removes trailing comma
string(REGEX REPLACE ", $" "" arrayValues ${arrayValues})
# converts the variable name into proper C identifier
# TODO: fix for legacy cmake
IF (${CMAKE_VERSION} GREATER 2.8.10)
string(MAKE_C_IDENTIFIER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME)
ENDIF()
string(TOUPPER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME)
# declares byte array and the length variables
set(arrayDefinition "const unsigned char ${BIN2H_VARIABLE_NAME}[] = { ${arrayValues} };")
set(arraySizeDefinition "const size_t ${BIN2H_VARIABLE_NAME}_SIZE = ${arraySize};")
set(declarations "${arrayDefinition}\n\n${arraySizeDefinition}\n\n")
if(BIN2H_APPEND)
file(APPEND ${BIN2H_HEADER_FILE} "${declarations}")
else()
file(WRITE ${BIN2H_HEADER_FILE} "${declarations}")
endif()
endfunction()
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner.cpp 0 → 100644
View file @ a91bf014
/*
This file is part of c-ethash.
c-ethash is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or
(at your option) any later version.
c-ethash is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with cpp-ethereum. If not, see <http://www.gnu.org/licenses/>.
*/
/** @file ethash_cl_miner.cpp
* @author Tim Hughes <tim@twistedfury.com>
* @date 2015
*/
#define _CRT_SECURE_NO_WARNINGS
#include <assert.h>
#include <queue>
#include "ethash_cl_miner.h"
#include "ethash_cl_miner_kernel.h"
#include <libethash/util.h>
#undef min
#undef max
#define HASH_BYTES 32
static void add_definition(std::string& source, char const* id, unsigned value)
{
char buf[256];
sprintf(buf, "#define %s %uu\n", id, value);
source.insert(source.begin(), buf, buf + strlen(buf));
}
ethash_cl_miner::ethash_cl_miner()
{
}
bool ethash_cl_miner::init(ethash_params const& params, const uint8_t seed[32], unsigned workgroup_size)
{
// store params
m_params = params;
// get all platforms
std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty())
{
debugf("No OpenCL platforms found.\n");
return false;
}
// use default platform
debugf("Using platform: %s\n", platforms[0].getInfo<CL_PLATFORM_NAME>().c_str());
// get GPU device of the default platform
std::vector<cl::Device> devices;
platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices);
if (devices.empty())
{
debugf("No OpenCL devices found.\n");
return false;
}
// use default device
cl::Device& device = devices[0];
debugf("Using device: %s\n", device.getInfo<CL_DEVICE_NAME>().c_str());
// create context
m_context = cl::Context({device});
m_queue = cl::CommandQueue(m_context, device);
// use requested workgroup size, but we require multiple of 8
m_workgroup_size = ((workgroup_size + 7) / 8) * 8;
// patch source code
std::string code(ETHASH_CL_MINER_KERNEL, ETHASH_CL_MINER_KERNEL + ETHASH_CL_MINER_KERNEL_SIZE);
add_definition(code, "GROUP_SIZE", m_workgroup_size);
add_definition(code, "DAG_SIZE", (unsigned)(params.full_size / MIX_BYTES));
add_definition(code, "ACCESSES", ACCESSES);
add_definition(code, "MAX_OUTPUTS", c_max_search_results);
//debugf("%s", code.c_str());
// create miner OpenCL program
cl::Program::Sources sources;
sources.push_back({code.c_str(), code.size()});
cl::Program program(m_context, sources);
try
{
program.build({device});
}
catch (cl::Error err)
{
debugf("%s\n", program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device).c_str());
return false;
}
m_hash_kernel = cl::Kernel(program, "ethash_hash");
m_search_kernel = cl::Kernel(program, "ethash_search");
// create buffer for dag
m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, params.full_size);
// create buffer for header
m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32);
// compute dag on CPU
{
void* cache_mem = malloc(params.cache_size + 63);
ethash_cache cache;
cache.mem = (void*)(((uintptr_t)cache_mem + 63) & ~63);
ethash_mkcache(&cache, &params, seed);
// 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);
ethash_compute_full_data(dag_ptr, &params, &cache);
m_queue.enqueueUnmapMemObject(m_dag, dag_ptr);
free(cache_mem);
}
// create mining buffers
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_search_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_max_search_results + 1) * sizeof(uint32_t));
}
return true;
}
void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count)
{
struct pending_batch
{
unsigned base;
unsigned count;
unsigned buf;
};
std::queue<pending_batch> pending;
// update header constant buffer
m_queue.enqueueWriteBuffer(m_header, true, 0, 32, header);
/*
__kernel void ethash_combined_hash(
__global hash32_t* g_hashes,
__constant hash32_t const* g_header,
__global hash128_t const* g_dag,
ulong start_nonce,
uint isolate
)
*/
m_hash_kernel.setArg(1, m_header);
m_hash_kernel.setArg(2, m_dag);
m_hash_kernel.setArg(3, nonce);
m_hash_kernel.setArg(4, ~0u); // have to pass this to stop the compile unrolling the loop
unsigned buf = 0;
for (unsigned i = 0; i < count || !pending.empty(); )
{
// how many this batch
if (i < count)
{
unsigned const this_count = std::min(count - i, c_hash_batch_size);
unsigned const batch_count = std::max(this_count, m_workgroup_size);
// supply output hash buffer to kernel
m_hash_kernel.setArg(0, m_hash_buf[buf]);
// execute it!
clock_t start_time = clock();
m_queue.enqueueNDRangeKernel(
m_hash_kernel,
cl::NullRange,
cl::NDRange(batch_count),
cl::NDRange(m_workgroup_size)
);
m_queue.flush();
pending.push({i, this_count, buf});
i += this_count;
buf = (buf + 1) % c_num_buffers;
}
// read results
if (i == count || pending.size() == c_num_buffers)
{
pending_batch const& batch = pending.front();
// 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);
memcpy(ret + batch.base*HASH_BYTES, hashes, batch.count*HASH_BYTES);
m_queue.enqueueUnmapMemObject(m_hash_buf[batch.buf], hashes);
pending.pop();
}
}
}
void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook)
{
struct pending_batch
{
uint64_t start_nonce;
unsigned buf;
};
std::queue<pending_batch> pending;
static uint32_t const c_zero = 0;
// update header constant buffer
m_queue.enqueueWriteBuffer(m_header, false, 0, 32, header);
for (unsigned i = 0; i != c_num_buffers; ++i)
{
m_queue.enqueueWriteBuffer(m_search_buf[i], false, 0, 4, &c_zero);
}
cl::Event pre_return_event;
m_queue.enqueueBarrierWithWaitList(NULL, &pre_return_event);
/*
__kernel void ethash_combined_search(
__global hash32_t* g_hashes, // 0
__constant hash32_t const* g_header, // 1
__global hash128_t const* g_dag, // 2
ulong start_nonce, // 3
ulong target, // 4
uint isolate // 5
)
*/
m_search_kernel.setArg(1, m_header);
m_search_kernel.setArg(2, m_dag);
// pass these to stop the compiler unrolling the loops
m_search_kernel.setArg(4, target);
m_search_kernel.setArg(5, ~0u);
unsigned buf = 0;
for (uint64_t start_nonce = 0; ; start_nonce += c_search_batch_size)
{
// supply output buffer to kernel
m_search_kernel.setArg(0, m_search_buf[buf]);
m_search_kernel.setArg(3, start_nonce);
// execute it!
m_queue.enqueueNDRangeKernel(m_search_kernel, cl::NullRange, c_search_batch_size, m_workgroup_size);
pending.push({start_nonce, buf});
buf = (buf + 1) % c_num_buffers;
// read results
if (pending.size() == c_num_buffers)
{
pending_batch const& batch = pending.front();
// could use pinned host pointer instead
uint32_t* results = (uint32_t*)m_queue.enqueueMapBuffer(m_search_buf[batch.buf], true, CL_MAP_READ, 0, (1+c_max_search_results) * sizeof(uint32_t));
unsigned num_found = std::min(results[0], c_max_search_results);
uint64_t nonces[c_max_search_results];
for (unsigned i = 0; i != num_found; ++i)
{
nonces[i] = batch.start_nonce + results[i+1];
}
m_queue.enqueueUnmapMemObject(m_search_buf[batch.buf], results);
bool exit = num_found && hook.found(nonces, num_found);
exit |= hook.searched(batch.start_nonce, c_search_batch_size); // always report searched before exit
if (exit)
break;
pending.pop();
}
}
// not safe to return until this is ready
pre_return_event.wait();
}
Godeps/_workspace/src/github.com/ethereum/ethash/libethash-cl/ethash_cl_miner.cpp → Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner_kernel.cl
View file @ a91bf014
/*
This file is part of c-ethash.
c-ethash is free software: you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation, either version 3 of the License, or
(at your option) any later version.
c-ethash is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with cpp-ethereum. If not, see <http://www.gnu.org/licenses/>.
*/
/** @file ethash_cl_miner.cpp
* @author Tim Hughes <tim@twistedfury.com>
* @date 2015
*/
#define _CRT_SECURE_NO_WARNINGS
#include <assert.h>
#include <queue>
#include "ethash_cl_miner.h"
#include <libethash/util.h>
#undef min
#undef max
#define HASH_BYTES 32
static char const ethash_inner_code[] = R"(
// author Tim Hughes <tim@twistedfury.com> // author Tim Hughes <tim@twistedfury.com>
// Tested on Radeon HD 7850 // Tested on Radeon HD 7850
// Hashrate: 15940347 hashes/s // Hashrate: 15940347 hashes/s
...@@ -495,260 +459,3 @@ __kernel void ethash_search( ...@@ -495,260 +459,3 @@ __kernel void ethash_search(
g_output[slot] = gid; g_output[slot] = gid;
} }
} }
)";
static void add_definition(std::string& source, char const* id, unsigned value)
{
char buf[256];
sprintf(buf, "#define %s %uu\n", id, value);
source.insert(source.begin(), buf, buf + strlen(buf));
}
ethash_cl_miner::ethash_cl_miner()
{
}
bool ethash_cl_miner::init(ethash_params const& params, const uint8_t seed[32], unsigned workgroup_size)
{
// store params
m_params = params;
// get all platforms
std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty())
{
debugf("No OpenCL platforms found.\n");
return false;
}
// use default platform
debugf("Using platform: %s\n", platforms[0].getInfo<CL_PLATFORM_NAME>().c_str());
// get GPU device of the default platform
std::vector<cl::Device> devices;
platforms[0].getDevices(CL_DEVICE_TYPE_ALL, &devices);
if (devices.empty())
{
debugf("No OpenCL devices found.\n");
return false;
}
// use default device
cl::Device& device = devices[0];
debugf("Using device: %s\n", device.getInfo<CL_DEVICE_NAME>().c_str());
// create context
m_context = cl::Context({device});
m_queue = cl::CommandQueue(m_context, device);
// use requested workgroup size, but we require multiple of 8
m_workgroup_size = ((workgroup_size + 7) / 8) * 8;
// patch source code
std::string code = ethash_inner_code;
add_definition(code, "GROUP_SIZE", m_workgroup_size);
add_definition(code, "DAG_SIZE", (unsigned)(params.full_size / MIX_BYTES));
add_definition(code, "ACCESSES", ACCESSES);
add_definition(code, "MAX_OUTPUTS", c_max_search_results);
//debugf("%s", code.c_str());
// create miner OpenCL program
cl::Program::Sources sources;
sources.push_back({code.c_str(), code.size()});
cl::Program program(m_context, sources);
try
{
program.build({device});
}
catch (cl::Error err)
{
debugf("%s\n", program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device).c_str());
return false;
}
m_hash_kernel = cl::Kernel(program, "ethash_hash");
m_search_kernel = cl::Kernel(program, "ethash_search");
// create buffer for dag
m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, params.full_size);
// create buffer for header
m_header = cl::Buffer(m_context, CL_MEM_READ_ONLY, 32);
// compute dag on CPU
{
void* cache_mem = malloc(params.cache_size + 63);
ethash_cache cache;
cache.mem = (void*)(((uintptr_t)cache_mem + 63) & ~63);
ethash_mkcache(&cache, &params, seed);
// 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);
ethash_compute_full_data(dag_ptr, &params, &cache);
m_queue.enqueueUnmapMemObject(m_dag, dag_ptr);
free(cache_mem);
}
// create mining buffers
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_search_buf[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_max_search_results + 1) * sizeof(uint32_t));
}
return true;
}
void ethash_cl_miner::hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count)
{
struct pending_batch
{
unsigned base;
unsigned count;
unsigned buf;
};
std::queue<pending_batch> pending;
// update header constant buffer
m_queue.enqueueWriteBuffer(m_header, true, 0, 32, header);
/*
__kernel void ethash_combined_hash(
__global hash32_t* g_hashes,
__constant hash32_t const* g_header,
__global hash128_t const* g_dag,
ulong start_nonce,
uint isolate
)
*/
m_hash_kernel.setArg(1, m_header);
m_hash_kernel.setArg(2, m_dag);
m_hash_kernel.setArg(3, nonce);
m_hash_kernel.setArg(4, ~0u); // have to pass this to stop the compile unrolling the loop
unsigned buf = 0;
for (unsigned i = 0; i < count || !pending.empty(); )
{
// how many this batch
if (i < count)
{
unsigned const this_count = std::min(count - i, c_hash_batch_size);
unsigned const batch_count = std::max(this_count, m_workgroup_size);
// supply output hash buffer to kernel
m_hash_kernel.setArg(0, m_hash_buf[buf]);
// execute it!
clock_t start_time = clock();
m_queue.enqueueNDRangeKernel(
m_hash_kernel,
cl::NullRange,
cl::NDRange(batch_count),
cl::NDRange(m_workgroup_size)
);
m_queue.flush();
pending.push({i, this_count, buf});
i += this_count;
buf = (buf + 1) % c_num_buffers;
}
// read results
if (i == count || pending.size() == c_num_buffers)
{
pending_batch const& batch = pending.front();
// 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);
memcpy(ret + batch.base*HASH_BYTES, hashes, batch.count*HASH_BYTES);
m_queue.enqueueUnmapMemObject(m_hash_buf[batch.buf], hashes);
pending.pop();
}
}
}
void ethash_cl_miner::search(uint8_t const* header, uint64_t target, search_hook& hook)
{
struct pending_batch
{
uint64_t start_nonce;
unsigned buf;
};
std::queue<pending_batch> pending;
static uint32_t const c_zero = 0;
// update header constant buffer
m_queue.enqueueWriteBuffer(m_header, false, 0, 32, header);
for (unsigned i = 0; i != c_num_buffers; ++i)
{
m_queue.enqueueWriteBuffer(m_search_buf[i], false, 0, 4, &c_zero);
}
cl::Event pre_return_event;
m_queue.enqueueBarrierWithWaitList(NULL, &pre_return_event);
/*
__kernel void ethash_combined_search(
__global hash32_t* g_hashes, // 0
__constant hash32_t const* g_header, // 1
__global hash128_t const* g_dag, // 2
ulong start_nonce, // 3
ulong target, // 4
uint isolate // 5
)
*/
m_search_kernel.setArg(1, m_header);
m_search_kernel.setArg(2, m_dag);
// pass these to stop the compiler unrolling the loops
m_search_kernel.setArg(4, target);
m_search_kernel.setArg(5, ~0u);
unsigned buf = 0;
for (uint64_t start_nonce = 0; ; start_nonce += c_search_batch_size)
{
// supply output buffer to kernel
m_search_kernel.setArg(0, m_search_buf[buf]);
m_search_kernel.setArg(3, start_nonce);
// execute it!
m_queue.enqueueNDRangeKernel(m_search_kernel, cl::NullRange, c_search_batch_size, m_workgroup_size);
pending.push({start_nonce, buf});
buf = (buf + 1) % c_num_buffers;
// read results
if (pending.size() == c_num_buffers)
{
pending_batch const& batch = pending.front();
// could use pinned host pointer instead
uint32_t* results = (uint32_t*)m_queue.enqueueMapBuffer(m_search_buf[batch.buf], true, CL_MAP_READ, 0, (1+c_max_search_results) * sizeof(uint32_t));
unsigned num_found = std::min(results[0], c_max_search_results);
uint64_t nonces[c_max_search_results];
for (unsigned i = 0; i != num_found; ++i)
{
nonces[i] = batch.start_nonce + results[i+1];
}
m_queue.enqueueUnmapMemObject(m_search_buf[batch.buf], results);
bool exit = num_found && hook.found(nonces, num_found);
exit |= hook.searched(batch.start_nonce, c_search_batch_size); // always report searched before exit
if (exit)
break;
pending.pop();
}
}
// not safe to return until this is ready
pre_return_event.wait();
}
Godeps/_workspace/src/github.com/ethereum/ethash/libethash/CMakeLists.txt → Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/CMakeLists.txt
View file @ a91bf014
set(LIBRARY ethash) set(LIBRARY ethash)
if (CPPETHEREUM)
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -fPIC")
#else ()
endif ()
set(CMAKE_BUILD_TYPE Release) set(CMAKE_BUILD_TYPE Release)
if (NOT MSVC) if (NOT MSVC)
...@@ -30,4 +36,4 @@ add_library(${LIBRARY} ${FILES}) ...@@ -30,4 +36,4 @@ add_library(${LIBRARY} ${FILES})
if (CRYPTOPP_FOUND) if (CRYPTOPP_FOUND)
TARGET_LINK_LIBRARIES(${LIBRARY} ${CRYPTOPP_LIBRARIES}) TARGET_LINK_LIBRARIES(${LIBRARY} ${CRYPTOPP_LIBRARIES})
endif() endif()
\ No newline at end of file
Godeps/_workspace/src/github.com/ethereum/ethash/libethash/ethash.h → Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/ethash.h
View file @ a91bf014
...@@ -14,6 +14,7 @@ ...@@ -14,6 +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 ethash.h /** @file ethash.h
* @date 2015 * @date 2015
*/ */
...@@ -25,9 +26,10 @@ ...@@ -25,9 +26,10 @@
#include <stddef.h> #include <stddef.h>
#include "compiler.h" #include "compiler.h"
#define REVISION 18 #define REVISION 20
#define DAGSIZE_BYTES_INIT 1073741824U #define DAGSIZE_BYTES_INIT 1073741824U // 2**30
#define DAG_GROWTH 113000000U #define DAG_GROWTH 8388608U // 2**23
#define CACHE_MULTIPLIER 1024
#define EPOCH_LENGTH 30000U #define EPOCH_LENGTH 30000U
#define MIX_BYTES 128 #define MIX_BYTES 128
#define DAG_PARENTS 256 #define DAG_PARENTS 256
...@@ -48,8 +50,8 @@ typedef struct ethash_return_value { ...@@ -48,8 +50,8 @@ typedef struct ethash_return_value {
uint8_t mix_hash[32]; uint8_t mix_hash[32];
} ethash_return_value; } ethash_return_value;
size_t const ethash_get_datasize(const uint32_t block_number); size_t ethash_get_datasize(const uint32_t block_number);
size_t const ethash_get_cachesize(const uint32_t block_number); size_t ethash_get_cachesize(const uint32_t block_number);
// initialize the parameters // initialize the parameters
static inline void ethash_params_init(ethash_params *params, const uint32_t block_number) { static inline void ethash_params_init(ethash_params *params, const uint32_t block_number) {
...@@ -58,7 +60,7 @@ static inline void ethash_params_init(ethash_params *params, const uint32_t bloc ...@@ -58,7 +60,7 @@ 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]);
...@@ -66,6 +68,11 @@ void ethash_compute_full_data(void *mem, ethash_params const *params, ethash_cac ...@@ -66,6 +68,11 @@ void ethash_compute_full_data(void *mem, ethash_params const *params, ethash_cac
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);
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_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); }
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]) {
......
Godeps/_workspace/src/github.com/ethereum/ethash/test/test.sh 0 → 100644
View file @ a91bf014
This diff is collapsed.
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment