subtly
10 years ago
99 changed files with 8604 additions and 1968 deletions
@ -0,0 +1,161 @@ |
|||||
|
#.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() |
||||
@ -0,0 +1,33 @@ |
|||||
|
# Find libcpuid |
||||
|
# |
||||
|
# Find the libcpuid includes and library |
||||
|
# |
||||
|
# if you nee to add a custom library search path, do it via via CMAKE_PREFIX_PATH |
||||
|
# |
||||
|
# This module defines |
||||
|
# CPUID_INCLUDE_DIRS, where to find header, etc. |
||||
|
# CPUID_LIBRARIES, the libraries needed to use cpuid. |
||||
|
# CPUID_FOUND, If false, do not try to use cpuid. |
||||
|
|
||||
|
# only look in default directories |
||||
|
find_path( |
||||
|
CPUID_INCLUDE_DIR |
||||
|
NAMES libcpuid/libcpuid.h |
||||
|
DOC "libcpuid include dir" |
||||
|
) |
||||
|
|
||||
|
find_library( |
||||
|
CPUID_LIBRARY |
||||
|
NAMES cpuid |
||||
|
DOC "libcpuid library" |
||||
|
) |
||||
|
|
||||
|
set(CPUID_INCLUDE_DIRS ${CPUID_INCLUDE_DIR}) |
||||
|
set(CPUID_LIBRARIES ${CPUID_LIBRARY}) |
||||
|
|
||||
|
# handle the QUIETLY and REQUIRED arguments and set CPUID_FOUND to TRUE |
||||
|
# if all listed variables are TRUE, hide their existence from configuration view |
||||
|
include(FindPackageHandleStandardArgs) |
||||
|
find_package_handle_standard_args(cpuid DEFAULT_MSG CPUID_INCLUDE_DIR CPUID_LIBRARY) |
||||
|
mark_as_advanced (CPUID_INCLUDE_DIR CPUID_LIBRARY) |
||||
|
|
||||
@ -0,0 +1,136 @@ |
|||||
|
#.rst: |
||||
|
# FindOpenCL |
||||
|
# ---------- |
||||
|
# |
||||
|
# Try to find OpenCL |
||||
|
# |
||||
|
# Once done this will define:: |
||||
|
# |
||||
|
# OpenCL_FOUND - True if OpenCL was found |
||||
|
# OpenCL_INCLUDE_DIRS - include directories for OpenCL |
||||
|
# OpenCL_LIBRARIES - link against this library 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 |
||||
|
# |
||||
|
|
||||
|
#============================================================================= |
||||
|
# Copyright 2014 Matthaeus G. Chajdas |
||||
|
# |
||||
|
# 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) |
||||
@ -0,0 +1,382 @@ |
|||||
|
#.rst: |
||||
|
# FindPackageHandleStandardArgs |
||||
|
# ----------------------------- |
||||
|
# |
||||
|
# |
||||
|
# |
||||
|
# FIND_PACKAGE_HANDLE_STANDARD_ARGS(<name> ... ) |
||||
|
# |
||||
|
# This function is intended to be used in FindXXX.cmake modules files. |
||||
|
# It handles the REQUIRED, QUIET and version-related arguments to |
||||
|
# find_package(). It also sets the <packagename>_FOUND variable. The |
||||
|
# package is considered found if all variables <var1>... listed contain |
||||
|
# valid results, e.g. valid filepaths. |
||||
|
# |
||||
|
# There are two modes of this function. The first argument in both |
||||
|
# modes is the name of the Find-module where it is called (in original |
||||
|
# casing). |
||||
|
# |
||||
|
# The first simple mode looks like this: |
||||
|
# |
||||
|
# :: |
||||
|
# |
||||
|
# FIND_PACKAGE_HANDLE_STANDARD_ARGS(<name> |
||||
|
# (DEFAULT_MSG|"Custom failure message") <var1>...<varN> ) |
||||
|
# |
||||
|
# If the variables <var1> to <varN> are all valid, then |
||||
|
# <UPPERCASED_NAME>_FOUND will be set to TRUE. If DEFAULT_MSG is given |
||||
|
# as second argument, then the function will generate itself useful |
||||
|
# success and error messages. You can also supply a custom error |
||||
|
# message for the failure case. This is not recommended. |
||||
|
# |
||||
|
# The second mode is more powerful and also supports version checking: |
||||
|
# |
||||
|
# :: |
||||
|
# |
||||
|
# FIND_PACKAGE_HANDLE_STANDARD_ARGS(NAME |
||||
|
# [FOUND_VAR <resultVar>] |
||||
|
# [REQUIRED_VARS <var1>...<varN>] |
||||
|
# [VERSION_VAR <versionvar>] |
||||
|
# [HANDLE_COMPONENTS] |
||||
|
# [CONFIG_MODE] |
||||
|
# [FAIL_MESSAGE "Custom failure message"] ) |
||||
|
# |
||||
|
# In this mode, the name of the result-variable can be set either to |
||||
|
# either <UPPERCASED_NAME>_FOUND or <OriginalCase_Name>_FOUND using the |
||||
|
# FOUND_VAR option. Other names for the result-variable are not |
||||
|
# allowed. So for a Find-module named FindFooBar.cmake, the two |
||||
|
# possible names are FooBar_FOUND and FOOBAR_FOUND. It is recommended |
||||
|
# to use the original case version. If the FOUND_VAR option is not |
||||
|
# used, the default is <UPPERCASED_NAME>_FOUND. |
||||
|
# |
||||
|
# As in the simple mode, if <var1> through <varN> are all valid, |
||||
|
# <packagename>_FOUND will be set to TRUE. After REQUIRED_VARS the |
||||
|
# variables which are required for this package are listed. Following |
||||
|
# VERSION_VAR the name of the variable can be specified which holds the |
||||
|
# version of the package which has been found. If this is done, this |
||||
|
# version will be checked against the (potentially) specified required |
||||
|
# version used in the find_package() call. The EXACT keyword is also |
||||
|
# handled. The default messages include information about the required |
||||
|
# version and the version which has been actually found, both if the |
||||
|
# version is ok or not. If the package supports components, use the |
||||
|
# HANDLE_COMPONENTS option to enable handling them. In this case, |
||||
|
# find_package_handle_standard_args() will report which components have |
||||
|
# been found and which are missing, and the <packagename>_FOUND variable |
||||
|
# will be set to FALSE if any of the required components (i.e. not the |
||||
|
# ones listed after OPTIONAL_COMPONENTS) are missing. Use the option |
||||
|
# CONFIG_MODE if your FindXXX.cmake module is a wrapper for a |
||||
|
# find_package(... NO_MODULE) call. In this case VERSION_VAR will be |
||||
|
# set to <NAME>_VERSION and the macro will automatically check whether |
||||
|
# the Config module was found. Via FAIL_MESSAGE a custom failure |
||||
|
# message can be specified, if this is not used, the default message |
||||
|
# will be displayed. |
||||
|
# |
||||
|
# Example for mode 1: |
||||
|
# |
||||
|
# :: |
||||
|
# |
||||
|
# find_package_handle_standard_args(LibXml2 DEFAULT_MSG |
||||
|
# LIBXML2_LIBRARY LIBXML2_INCLUDE_DIR) |
||||
|
# |
||||
|
# |
||||
|
# |
||||
|
# LibXml2 is considered to be found, if both LIBXML2_LIBRARY and |
||||
|
# LIBXML2_INCLUDE_DIR are valid. Then also LIBXML2_FOUND is set to |
||||
|
# TRUE. If it is not found and REQUIRED was used, it fails with |
||||
|
# FATAL_ERROR, independent whether QUIET was used or not. If it is |
||||
|
# found, success will be reported, including the content of <var1>. On |
||||
|
# repeated Cmake runs, the same message won't be printed again. |
||||
|
# |
||||
|
# Example for mode 2: |
||||
|
# |
||||
|
# :: |
||||
|
# |
||||
|
# find_package_handle_standard_args(LibXslt |
||||
|
# FOUND_VAR LibXslt_FOUND |
||||
|
# REQUIRED_VARS LibXslt_LIBRARIES LibXslt_INCLUDE_DIRS |
||||
|
# VERSION_VAR LibXslt_VERSION_STRING) |
||||
|
# |
||||
|
# In this case, LibXslt is considered to be found if the variable(s) |
||||
|
# listed after REQUIRED_VAR are all valid, i.e. LibXslt_LIBRARIES and |
||||
|
# LibXslt_INCLUDE_DIRS in this case. The result will then be stored in |
||||
|
# LibXslt_FOUND . Also the version of LibXslt will be checked by using |
||||
|
# the version contained in LibXslt_VERSION_STRING. Since no |
||||
|
# FAIL_MESSAGE is given, the default messages will be printed. |
||||
|
# |
||||
|
# Another example for mode 2: |
||||
|
# |
||||
|
# :: |
||||
|
# |
||||
|
# find_package(Automoc4 QUIET NO_MODULE HINTS /opt/automoc4) |
||||
|
# find_package_handle_standard_args(Automoc4 CONFIG_MODE) |
||||
|
# |
||||
|
# In this case, FindAutmoc4.cmake wraps a call to find_package(Automoc4 |
||||
|
# NO_MODULE) and adds an additional search directory for automoc4. Here |
||||
|
# the result will be stored in AUTOMOC4_FOUND. The following |
||||
|
# FIND_PACKAGE_HANDLE_STANDARD_ARGS() call produces a proper |
||||
|
# success/error message. |
||||
|
|
||||
|
#============================================================================= |
||||
|
# Copyright 2007-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.) |
||||
|
|
||||
|
include(${CMAKE_CURRENT_LIST_DIR}/FindPackageMessage.cmake) |
||||
|
include(${CMAKE_CURRENT_LIST_DIR}/CMakeParseArguments.cmake) |
||||
|
|
||||
|
# internal helper macro |
||||
|
macro(_FPHSA_FAILURE_MESSAGE _msg) |
||||
|
if (${_NAME}_FIND_REQUIRED) |
||||
|
message(FATAL_ERROR "${_msg}") |
||||
|
else () |
||||
|
if (NOT ${_NAME}_FIND_QUIETLY) |
||||
|
message(STATUS "${_msg}") |
||||
|
endif () |
||||
|
endif () |
||||
|
endmacro() |
||||
|
|
||||
|
|
||||
|
# internal helper macro to generate the failure message when used in CONFIG_MODE: |
||||
|
macro(_FPHSA_HANDLE_FAILURE_CONFIG_MODE) |
||||
|
# <name>_CONFIG is set, but FOUND is false, this means that some other of the REQUIRED_VARS was not found: |
||||
|
if(${_NAME}_CONFIG) |
||||
|
_FPHSA_FAILURE_MESSAGE("${FPHSA_FAIL_MESSAGE}: missing: ${MISSING_VARS} (found ${${_NAME}_CONFIG} ${VERSION_MSG})") |
||||
|
else() |
||||
|
# If _CONSIDERED_CONFIGS is set, the config-file has been found, but no suitable version. |
||||
|
# List them all in the error message: |
||||
|
if(${_NAME}_CONSIDERED_CONFIGS) |
||||
|
set(configsText "") |
||||
|
list(LENGTH ${_NAME}_CONSIDERED_CONFIGS configsCount) |
||||
|
math(EXPR configsCount "${configsCount} - 1") |
||||
|
foreach(currentConfigIndex RANGE ${configsCount}) |
||||
|
list(GET ${_NAME}_CONSIDERED_CONFIGS ${currentConfigIndex} filename) |
||||
|
list(GET ${_NAME}_CONSIDERED_VERSIONS ${currentConfigIndex} version) |
||||
|
set(configsText "${configsText} ${filename} (version ${version})\n") |
||||
|
endforeach() |
||||
|
if (${_NAME}_NOT_FOUND_MESSAGE) |
||||
|
set(configsText "${configsText} Reason given by package: ${${_NAME}_NOT_FOUND_MESSAGE}\n") |
||||
|
endif() |
||||
|
_FPHSA_FAILURE_MESSAGE("${FPHSA_FAIL_MESSAGE} ${VERSION_MSG}, checked the following files:\n${configsText}") |
||||
|
|
||||
|
else() |
||||
|
# Simple case: No Config-file was found at all: |
||||
|
_FPHSA_FAILURE_MESSAGE("${FPHSA_FAIL_MESSAGE}: found neither ${_NAME}Config.cmake nor ${_NAME_LOWER}-config.cmake ${VERSION_MSG}") |
||||
|
endif() |
||||
|
endif() |
||||
|
endmacro() |
||||
|
|
||||
|
|
||||
|
function(FIND_PACKAGE_HANDLE_STANDARD_ARGS _NAME _FIRST_ARG) |
||||
|
|
||||
|
# set up the arguments for CMAKE_PARSE_ARGUMENTS and check whether we are in |
||||
|
# new extended or in the "old" mode: |
||||
|
set(options CONFIG_MODE HANDLE_COMPONENTS) |
||||
|
set(oneValueArgs FAIL_MESSAGE VERSION_VAR FOUND_VAR) |
||||
|
set(multiValueArgs REQUIRED_VARS) |
||||
|
set(_KEYWORDS_FOR_EXTENDED_MODE ${options} ${oneValueArgs} ${multiValueArgs} ) |
||||
|
list(FIND _KEYWORDS_FOR_EXTENDED_MODE "${_FIRST_ARG}" INDEX) |
||||
|
|
||||
|
if(${INDEX} EQUAL -1) |
||||
|
set(FPHSA_FAIL_MESSAGE ${_FIRST_ARG}) |
||||
|
set(FPHSA_REQUIRED_VARS ${ARGN}) |
||||
|
set(FPHSA_VERSION_VAR) |
||||
|
else() |
||||
|
|
||||
|
CMAKE_PARSE_ARGUMENTS(FPHSA "${options}" "${oneValueArgs}" "${multiValueArgs}" ${_FIRST_ARG} ${ARGN}) |
||||
|
|
||||
|
if(FPHSA_UNPARSED_ARGUMENTS) |
||||
|
message(FATAL_ERROR "Unknown keywords given to FIND_PACKAGE_HANDLE_STANDARD_ARGS(): \"${FPHSA_UNPARSED_ARGUMENTS}\"") |
||||
|
endif() |
||||
|
|
||||
|
if(NOT FPHSA_FAIL_MESSAGE) |
||||
|
set(FPHSA_FAIL_MESSAGE "DEFAULT_MSG") |
||||
|
endif() |
||||
|
endif() |
||||
|
|
||||
|
# now that we collected all arguments, process them |
||||
|
|
||||
|
if("x${FPHSA_FAIL_MESSAGE}" STREQUAL "xDEFAULT_MSG") |
||||
|
set(FPHSA_FAIL_MESSAGE "Could NOT find ${_NAME}") |
||||
|
endif() |
||||
|
|
||||
|
# In config-mode, we rely on the variable <package>_CONFIG, which is set by find_package() |
||||
|
# when it successfully found the config-file, including version checking: |
||||
|
if(FPHSA_CONFIG_MODE) |
||||
|
list(INSERT FPHSA_REQUIRED_VARS 0 ${_NAME}_CONFIG) |
||||
|
list(REMOVE_DUPLICATES FPHSA_REQUIRED_VARS) |
||||
|
set(FPHSA_VERSION_VAR ${_NAME}_VERSION) |
||||
|
endif() |
||||
|
|
||||
|
if(NOT FPHSA_REQUIRED_VARS) |
||||
|
message(FATAL_ERROR "No REQUIRED_VARS specified for FIND_PACKAGE_HANDLE_STANDARD_ARGS()") |
||||
|
endif() |
||||
|
|
||||
|
list(GET FPHSA_REQUIRED_VARS 0 _FIRST_REQUIRED_VAR) |
||||
|
|
||||
|
string(TOUPPER ${_NAME} _NAME_UPPER) |
||||
|
string(TOLOWER ${_NAME} _NAME_LOWER) |
||||
|
|
||||
|
if(FPHSA_FOUND_VAR) |
||||
|
if(FPHSA_FOUND_VAR MATCHES "^${_NAME}_FOUND$" OR FPHSA_FOUND_VAR MATCHES "^${_NAME_UPPER}_FOUND$") |
||||
|
set(_FOUND_VAR ${FPHSA_FOUND_VAR}) |
||||
|
else() |
||||
|
message(FATAL_ERROR "The argument for FOUND_VAR is \"${FPHSA_FOUND_VAR}\", but only \"${_NAME}_FOUND\" and \"${_NAME_UPPER}_FOUND\" are valid names.") |
||||
|
endif() |
||||
|
else() |
||||
|
set(_FOUND_VAR ${_NAME_UPPER}_FOUND) |
||||
|
endif() |
||||
|
|
||||
|
# collect all variables which were not found, so they can be printed, so the |
||||
|
# user knows better what went wrong (#6375) |
||||
|
set(MISSING_VARS "") |
||||
|
set(DETAILS "") |
||||
|
# check if all passed variables are valid |
||||
|
unset(${_FOUND_VAR}) |
||||
|
foreach(_CURRENT_VAR ${FPHSA_REQUIRED_VARS}) |
||||
|
if(NOT ${_CURRENT_VAR}) |
||||
|
set(${_FOUND_VAR} FALSE) |
||||
|
set(MISSING_VARS "${MISSING_VARS} ${_CURRENT_VAR}") |
||||
|
else() |
||||
|
set(DETAILS "${DETAILS}[${${_CURRENT_VAR}}]") |
||||
|
endif() |
||||
|
endforeach() |
||||
|
if(NOT "${${_FOUND_VAR}}" STREQUAL "FALSE") |
||||
|
set(${_FOUND_VAR} TRUE) |
||||
|
endif() |
||||
|
|
||||
|
# component handling |
||||
|
unset(FOUND_COMPONENTS_MSG) |
||||
|
unset(MISSING_COMPONENTS_MSG) |
||||
|
|
||||
|
if(FPHSA_HANDLE_COMPONENTS) |
||||
|
foreach(comp ${${_NAME}_FIND_COMPONENTS}) |
||||
|
if(${_NAME}_${comp}_FOUND) |
||||
|
|
||||
|
if(NOT DEFINED FOUND_COMPONENTS_MSG) |
||||
|
set(FOUND_COMPONENTS_MSG "found components: ") |
||||
|
endif() |
||||
|
set(FOUND_COMPONENTS_MSG "${FOUND_COMPONENTS_MSG} ${comp}") |
||||
|
|
||||
|
else() |
||||
|
|
||||
|
if(NOT DEFINED MISSING_COMPONENTS_MSG) |
||||
|
set(MISSING_COMPONENTS_MSG "missing components: ") |
||||
|
endif() |
||||
|
set(MISSING_COMPONENTS_MSG "${MISSING_COMPONENTS_MSG} ${comp}") |
||||
|
|
||||
|
if(${_NAME}_FIND_REQUIRED_${comp}) |
||||
|
set(${_FOUND_VAR} FALSE) |
||||
|
set(MISSING_VARS "${MISSING_VARS} ${comp}") |
||||
|
endif() |
||||
|
|
||||
|
endif() |
||||
|
endforeach() |
||||
|
set(COMPONENT_MSG "${FOUND_COMPONENTS_MSG} ${MISSING_COMPONENTS_MSG}") |
||||
|
set(DETAILS "${DETAILS}[c${COMPONENT_MSG}]") |
||||
|
endif() |
||||
|
|
||||
|
# version handling: |
||||
|
set(VERSION_MSG "") |
||||
|
set(VERSION_OK TRUE) |
||||
|
set(VERSION ${${FPHSA_VERSION_VAR}}) |
||||
|
|
||||
|
# check with DEFINED here as the requested or found version may be "0" |
||||
|
if (DEFINED ${_NAME}_FIND_VERSION) |
||||
|
if(DEFINED ${FPHSA_VERSION_VAR}) |
||||
|
|
||||
|
if(${_NAME}_FIND_VERSION_EXACT) # exact version required |
||||
|
# count the dots in the version string |
||||
|
string(REGEX REPLACE "[^.]" "" _VERSION_DOTS "${VERSION}") |
||||
|
# add one dot because there is one dot more than there are components |
||||
|
string(LENGTH "${_VERSION_DOTS}." _VERSION_DOTS) |
||||
|
if (_VERSION_DOTS GREATER ${_NAME}_FIND_VERSION_COUNT) |
||||
|
# Because of the C++ implementation of find_package() ${_NAME}_FIND_VERSION_COUNT |
||||
|
# is at most 4 here. Therefore a simple lookup table is used. |
||||
|
if (${_NAME}_FIND_VERSION_COUNT EQUAL 1) |
||||
|
set(_VERSION_REGEX "[^.]*") |
||||
|
elseif (${_NAME}_FIND_VERSION_COUNT EQUAL 2) |
||||
|
set(_VERSION_REGEX "[^.]*\\.[^.]*") |
||||
|
elseif (${_NAME}_FIND_VERSION_COUNT EQUAL 3) |
||||
|
set(_VERSION_REGEX "[^.]*\\.[^.]*\\.[^.]*") |
||||
|
else () |
||||
|
set(_VERSION_REGEX "[^.]*\\.[^.]*\\.[^.]*\\.[^.]*") |
||||
|
endif () |
||||
|
string(REGEX REPLACE "^(${_VERSION_REGEX})\\..*" "\\1" _VERSION_HEAD "${VERSION}") |
||||
|
unset(_VERSION_REGEX) |
||||
|
if (NOT ${_NAME}_FIND_VERSION VERSION_EQUAL _VERSION_HEAD) |
||||
|
set(VERSION_MSG "Found unsuitable version \"${VERSION}\", but required is exact version \"${${_NAME}_FIND_VERSION}\"") |
||||
|
set(VERSION_OK FALSE) |
||||
|
else () |
||||
|
set(VERSION_MSG "(found suitable exact version \"${VERSION}\")") |
||||
|
endif () |
||||
|
unset(_VERSION_HEAD) |
||||
|
else () |
||||
|
if (NOT "${${_NAME}_FIND_VERSION}" VERSION_EQUAL "${VERSION}") |
||||
|
set(VERSION_MSG "Found unsuitable version \"${VERSION}\", but required is exact version \"${${_NAME}_FIND_VERSION}\"") |
||||
|
set(VERSION_OK FALSE) |
||||
|
else () |
||||
|
set(VERSION_MSG "(found suitable exact version \"${VERSION}\")") |
||||
|
endif () |
||||
|
endif () |
||||
|
unset(_VERSION_DOTS) |
||||
|
|
||||
|
else() # minimum version specified: |
||||
|
if ("${${_NAME}_FIND_VERSION}" VERSION_GREATER "${VERSION}") |
||||
|
set(VERSION_MSG "Found unsuitable version \"${VERSION}\", but required is at least \"${${_NAME}_FIND_VERSION}\"") |
||||
|
set(VERSION_OK FALSE) |
||||
|
else () |
||||
|
set(VERSION_MSG "(found suitable version \"${VERSION}\", minimum required is \"${${_NAME}_FIND_VERSION}\")") |
||||
|
endif () |
||||
|
endif() |
||||
|
|
||||
|
else() |
||||
|
|
||||
|
# if the package was not found, but a version was given, add that to the output: |
||||
|
if(${_NAME}_FIND_VERSION_EXACT) |
||||
|
set(VERSION_MSG "(Required is exact version \"${${_NAME}_FIND_VERSION}\")") |
||||
|
else() |
||||
|
set(VERSION_MSG "(Required is at least version \"${${_NAME}_FIND_VERSION}\")") |
||||
|
endif() |
||||
|
|
||||
|
endif() |
||||
|
else () |
||||
|
if(VERSION) |
||||
|
set(VERSION_MSG "(found version \"${VERSION}\")") |
||||
|
endif() |
||||
|
endif () |
||||
|
|
||||
|
if(VERSION_OK) |
||||
|
set(DETAILS "${DETAILS}[v${VERSION}(${${_NAME}_FIND_VERSION})]") |
||||
|
else() |
||||
|
set(${_FOUND_VAR} FALSE) |
||||
|
endif() |
||||
|
|
||||
|
|
||||
|
# print the result: |
||||
|
if (${_FOUND_VAR}) |
||||
|
FIND_PACKAGE_MESSAGE(${_NAME} "Found ${_NAME}: ${${_FIRST_REQUIRED_VAR}} ${VERSION_MSG} ${COMPONENT_MSG}" "${DETAILS}") |
||||
|
else () |
||||
|
|
||||
|
if(FPHSA_CONFIG_MODE) |
||||
|
_FPHSA_HANDLE_FAILURE_CONFIG_MODE() |
||||
|
else() |
||||
|
if(NOT VERSION_OK) |
||||
|
_FPHSA_FAILURE_MESSAGE("${FPHSA_FAIL_MESSAGE}: ${VERSION_MSG} (found ${${_FIRST_REQUIRED_VAR}})") |
||||
|
else() |
||||
|
_FPHSA_FAILURE_MESSAGE("${FPHSA_FAIL_MESSAGE} (missing: ${MISSING_VARS}) ${VERSION_MSG}") |
||||
|
endif() |
||||
|
endif() |
||||
|
|
||||
|
endif () |
||||
|
|
||||
|
set(${_FOUND_VAR} ${${_FOUND_VAR}} PARENT_SCOPE) |
||||
|
|
||||
|
endfunction() |
||||
@ -0,0 +1,57 @@ |
|||||
|
#.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() |
||||
@ -0,0 +1,19 @@ |
|||||
|
set(EXECUTABLE ethash-cl) |
||||
|
|
||||
|
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) |
||||
|
|
||||
|
aux_source_directory(. SRC_LIST) |
||||
|
file(GLOB HEADERS "*.h") |
||||
|
|
||||
|
include_directories(${CMAKE_CURRENT_BINARY_DIR}) |
||||
|
include_directories(${OpenCL_INCLUDE_DIRS}}) |
||||
|
include_directories(..) |
||||
|
add_library(${EXECUTABLE} ${SRC_LIST} ${HEADERS}) |
||||
|
TARGET_LINK_LIBRARIES(${EXECUTABLE} ${OpenCL_LIBRARIES} ethash) |
||||
|
|
||||
|
install( TARGETS ${EXECUTABLE} RUNTIME DESTINATION bin ARCHIVE DESTINATION lib LIBRARY DESTINATION lib ) |
||||
|
install( FILES ${HEADERS} DESTINATION include/${EXECUTABLE} ) |
||||
|
|
||||
@ -0,0 +1,86 @@ |
|||||
|
# 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 |
||||
|
IF (${CMAKE_VERSION} GREATER 2.8.10) # fix for legacy cmake |
||||
|
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() |
||||
File diff suppressed because it is too large
@ -0,0 +1,366 @@ |
|||||
|
/*
|
||||
|
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 <cstdio> |
||||
|
#include <cstdlib> |
||||
|
#include <assert.h> |
||||
|
#include <queue> |
||||
|
#include <vector> |
||||
|
#include <libethash/util.h> |
||||
|
#include <libethash/ethash.h> |
||||
|
#include "ethash_cl_miner.h" |
||||
|
#include "ethash_cl_miner_kernel.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 max |
||||
|
|
||||
|
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::search_hook::~search_hook() {} |
||||
|
|
||||
|
ethash_cl_miner::ethash_cl_miner() |
||||
|
: m_opencl_1_1() |
||||
|
{ |
||||
|
} |
||||
|
|
||||
|
std::string ethash_cl_miner::platform_info() |
||||
|
{ |
||||
|
std::vector<cl::Platform> platforms; |
||||
|
cl::Platform::get(&platforms); |
||||
|
if (platforms.empty()) |
||||
|
{ |
||||
|
debugf("No OpenCL platforms found.\n"); |
||||
|
return std::string(); |
||||
|
} |
||||
|
|
||||
|
// 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 std::string(); |
||||
|
} |
||||
|
|
||||
|
// use default device
|
||||
|
unsigned device_num = 0; |
||||
|
cl::Device& device = devices[device_num]; |
||||
|
std::string device_version = device.getInfo<CL_DEVICE_VERSION>(); |
||||
|
|
||||
|
return "{ platform: '" + platforms[0].getInfo<CL_PLATFORM_NAME>() + "', device: '" + device.getInfo<CL_DEVICE_NAME>() + "', version: '" + device_version + "' }"; |
||||
|
} |
||||
|
|
||||
|
void ethash_cl_miner::finish() |
||||
|
{ |
||||
|
if (m_queue()) |
||||
|
{ |
||||
|
m_queue.finish(); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
bool ethash_cl_miner::init(ethash_params const& params, std::function<void(void*)> _fillDAG, unsigned workgroup_size, unsigned _deviceId) |
||||
|
{ |
||||
|
// 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
|
||||
|
fprintf(stderr, "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[std::min<unsigned>(_deviceId, devices.size() - 1)]; |
||||
|
for (unsigned n = 0; n < devices.size(); ++n) |
||||
|
{ |
||||
|
auto version = devices[n].getInfo<CL_DEVICE_VERSION>(); |
||||
|
auto name = devices[n].getInfo<CL_DEVICE_NAME>(); |
||||
|
fprintf(stderr, "%s %d: %s (%s)\n", n == _deviceId ? "USING " : " ", n, name.c_str(), version.c_str()); |
||||
|
} |
||||
|
std::string device_version = device.getInfo<CL_DEVICE_VERSION>(); |
||||
|
fprintf(stderr, "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
|
||||
|
m_context = cl::Context(std::vector<cl::Device>(&device, &device + 1)); |
||||
|
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 / ETHASH_MIX_BYTES)); |
||||
|
add_definition(code, "ACCESSES", ETHASH_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
|
||||
|
{ |
||||
|
// 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, m_opencl_1_1 ? CL_MAP_WRITE : CL_MAP_WRITE_INVALIDATE_REGION, 0, params.full_size); |
||||
|
// memcpying 1GB: horrible... really. horrible. but necessary since we can't mmap *and* gpumap.
|
||||
|
_fillDAG(dag_ptr); |
||||
|
m_queue.enqueueUnmapMemObject(m_dag, dag_ptr); |
||||
|
} |
||||
|
|
||||
|
// create mining buffers
|
||||
|
for (unsigned i = 0; i != c_num_buffers; ++i) |
||||
|
{ |
||||
|
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)); |
||||
|
} |
||||
|
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<unsigned>(count - i, c_hash_batch_size); |
||||
|
unsigned const batch_count = std::max<unsigned>(this_count, m_workgroup_size); |
||||
|
|
||||
|
// supply output hash buffer to kernel
|
||||
|
m_hash_kernel.setArg(0, m_hash_buf[buf]); |
||||
|
|
||||
|
// execute it!
|
||||
|
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 * ETHASH_BYTES); |
||||
|
memcpy(ret + batch.base*ETHASH_BYTES, hashes, batch.count*ETHASH_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); |
||||
|
} |
||||
|
|
||||
|
#if CL_VERSION_1_2 && 0 |
||||
|
cl::Event pre_return_event; |
||||
|
if (!m_opencl_1_1) |
||||
|
{ |
||||
|
m_queue.enqueueBarrierWithWaitList(NULL, &pre_return_event); |
||||
|
} |
||||
|
else |
||||
|
#endif |
||||
|
{ |
||||
|
m_queue.finish(); |
||||
|
} |
||||
|
|
||||
|
/*
|
||||
|
__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<unsigned>(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; |
||||
|
|
||||
|
// reset search buffer if we're still going
|
||||
|
if (num_found) |
||||
|
m_queue.enqueueWriteBuffer(m_search_buf[batch.buf], true, 0, 4, &c_zero); |
||||
|
|
||||
|
pending.pop(); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
// not safe to return until this is ready
|
||||
|
#if CL_VERSION_1_2 && 0 |
||||
|
if (!m_opencl_1_1) |
||||
|
{ |
||||
|
pre_return_event.wait(); |
||||
|
} |
||||
|
#endif |
||||
|
} |
||||
|
|
||||
@ -0,0 +1,55 @@ |
|||||
|
#pragma once |
||||
|
|
||||
|
#define __CL_ENABLE_EXCEPTIONS |
||||
|
#define CL_USE_DEPRECATED_OPENCL_2_0_APIS |
||||
|
|
||||
|
#if defined(__clang__) |
||||
|
#pragma clang diagnostic push |
||||
|
#pragma clang diagnostic ignored "-Wunused-parameter" |
||||
|
#include "cl.hpp" |
||||
|
#pragma clang diagnostic pop |
||||
|
#else |
||||
|
#include "cl.hpp" |
||||
|
#endif |
||||
|
|
||||
|
#include <time.h> |
||||
|
#include <functional> |
||||
|
#include <libethash/ethash.h> |
||||
|
|
||||
|
class ethash_cl_miner |
||||
|
{ |
||||
|
public: |
||||
|
struct search_hook |
||||
|
{ |
||||
|
virtual ~search_hook(); // always a virtual destructor for a class with virtuals.
|
||||
|
|
||||
|
// reports progress, return true to abort
|
||||
|
virtual bool found(uint64_t const* nonces, uint32_t count) = 0; |
||||
|
virtual bool searched(uint64_t start_nonce, uint32_t count) = 0; |
||||
|
}; |
||||
|
|
||||
|
public: |
||||
|
ethash_cl_miner(); |
||||
|
|
||||
|
bool init(ethash_params const& params, std::function<void(void*)> _fillDAG, unsigned workgroup_size = 64, unsigned _deviceId = 0); |
||||
|
static std::string platform_info(); |
||||
|
|
||||
|
void finish(); |
||||
|
void hash(uint8_t* ret, uint8_t const* header, uint64_t nonce, unsigned count); |
||||
|
void search(uint8_t const* header, uint64_t target, search_hook& hook); |
||||
|
|
||||
|
private: |
||||
|
enum { c_max_search_results = 63, c_num_buffers = 2, c_hash_batch_size = 1024, c_search_batch_size = 1024*256 }; |
||||
|
|
||||
|
ethash_params m_params; |
||||
|
cl::Context m_context; |
||||
|
cl::CommandQueue m_queue; |
||||
|
cl::Kernel m_hash_kernel; |
||||
|
cl::Kernel m_search_kernel; |
||||
|
cl::Buffer m_dag; |
||||
|
cl::Buffer m_header; |
||||
|
cl::Buffer m_hash_buf[c_num_buffers]; |
||||
|
cl::Buffer m_search_buf[c_num_buffers]; |
||||
|
unsigned m_workgroup_size; |
||||
|
bool m_opencl_1_1; |
||||
|
}; |
||||
@ -0,0 +1,460 @@ |
|||||
|
// author Tim Hughes <tim@twistedfury.com> |
||||
|
// Tested on Radeon HD 7850 |
||||
|
// Hashrate: 15940347 hashes/s |
||||
|
// Bandwidth: 124533 MB/s |
||||
|
// search kernel should fit in <= 84 VGPRS (3 wavefronts) |
||||
|
|
||||
|
#define THREADS_PER_HASH (128 / 16) |
||||
|
#define HASHES_PER_LOOP (GROUP_SIZE / THREADS_PER_HASH) |
||||
|
|
||||
|
#define FNV_PRIME 0x01000193 |
||||
|
|
||||
|
__constant uint2 const Keccak_f1600_RC[24] = { |
||||
|
(uint2)(0x00000001, 0x00000000), |
||||
|
(uint2)(0x00008082, 0x00000000), |
||||
|
(uint2)(0x0000808a, 0x80000000), |
||||
|
(uint2)(0x80008000, 0x80000000), |
||||
|
(uint2)(0x0000808b, 0x00000000), |
||||
|
(uint2)(0x80000001, 0x00000000), |
||||
|
(uint2)(0x80008081, 0x80000000), |
||||
|
(uint2)(0x00008009, 0x80000000), |
||||
|
(uint2)(0x0000008a, 0x00000000), |
||||
|
(uint2)(0x00000088, 0x00000000), |
||||
|
(uint2)(0x80008009, 0x00000000), |
||||
|
(uint2)(0x8000000a, 0x00000000), |
||||
|
(uint2)(0x8000808b, 0x00000000), |
||||
|
(uint2)(0x0000008b, 0x80000000), |
||||
|
(uint2)(0x00008089, 0x80000000), |
||||
|
(uint2)(0x00008003, 0x80000000), |
||||
|
(uint2)(0x00008002, 0x80000000), |
||||
|
(uint2)(0x00000080, 0x80000000), |
||||
|
(uint2)(0x0000800a, 0x00000000), |
||||
|
(uint2)(0x8000000a, 0x80000000), |
||||
|
(uint2)(0x80008081, 0x80000000), |
||||
|
(uint2)(0x00008080, 0x80000000), |
||||
|
(uint2)(0x80000001, 0x00000000), |
||||
|
(uint2)(0x80008008, 0x80000000), |
||||
|
}; |
||||
|
|
||||
|
void keccak_f1600_round(uint2* a, uint r, uint out_size) |
||||
|
{ |
||||
|
#if !__ENDIAN_LITTLE__ |
||||
|
for (uint i = 0; i != 25; ++i) |
||||
|
a[i] = a[i].yx; |
||||
|
#endif |
||||
|
|
||||
|
uint2 b[25]; |
||||
|
uint2 t; |
||||
|
|
||||
|
// Theta |
||||
|
b[0] = a[0] ^ a[5] ^ a[10] ^ a[15] ^ a[20]; |
||||
|
b[1] = a[1] ^ a[6] ^ a[11] ^ a[16] ^ a[21]; |
||||
|
b[2] = a[2] ^ a[7] ^ a[12] ^ a[17] ^ a[22]; |
||||
|
b[3] = a[3] ^ a[8] ^ a[13] ^ a[18] ^ a[23]; |
||||
|
b[4] = a[4] ^ a[9] ^ a[14] ^ a[19] ^ a[24]; |
||||
|
t = b[4] ^ (uint2)(b[1].x << 1 | b[1].y >> 31, b[1].y << 1 | b[1].x >> 31); |
||||
|
a[0] ^= t; |
||||
|
a[5] ^= t; |
||||
|
a[10] ^= t; |
||||
|
a[15] ^= t; |
||||
|
a[20] ^= t; |
||||
|
t = b[0] ^ (uint2)(b[2].x << 1 | b[2].y >> 31, b[2].y << 1 | b[2].x >> 31); |
||||
|
a[1] ^= t; |
||||
|
a[6] ^= t; |
||||
|
a[11] ^= t; |
||||
|
a[16] ^= t; |
||||
|
a[21] ^= t; |
||||
|
t = b[1] ^ (uint2)(b[3].x << 1 | b[3].y >> 31, b[3].y << 1 | b[3].x >> 31); |
||||
|
a[2] ^= t; |
||||
|
a[7] ^= t; |
||||
|
a[12] ^= t; |
||||
|
a[17] ^= t; |
||||
|
a[22] ^= t; |
||||
|
t = b[2] ^ (uint2)(b[4].x << 1 | b[4].y >> 31, b[4].y << 1 | b[4].x >> 31); |
||||
|
a[3] ^= t; |
||||
|
a[8] ^= t; |
||||
|
a[13] ^= t; |
||||
|
a[18] ^= t; |
||||
|
a[23] ^= t; |
||||
|
t = b[3] ^ (uint2)(b[0].x << 1 | b[0].y >> 31, b[0].y << 1 | b[0].x >> 31); |
||||
|
a[4] ^= t; |
||||
|
a[9] ^= t; |
||||
|
a[14] ^= t; |
||||
|
a[19] ^= t; |
||||
|
a[24] ^= t; |
||||
|
|
||||
|
// Rho Pi |
||||
|
b[0] = a[0]; |
||||
|
b[10] = (uint2)(a[1].x << 1 | a[1].y >> 31, a[1].y << 1 | a[1].x >> 31); |
||||
|
b[7] = (uint2)(a[10].x << 3 | a[10].y >> 29, a[10].y << 3 | a[10].x >> 29); |
||||
|
b[11] = (uint2)(a[7].x << 6 | a[7].y >> 26, a[7].y << 6 | a[7].x >> 26); |
||||
|
b[17] = (uint2)(a[11].x << 10 | a[11].y >> 22, a[11].y << 10 | a[11].x >> 22); |
||||
|
b[18] = (uint2)(a[17].x << 15 | a[17].y >> 17, a[17].y << 15 | a[17].x >> 17); |
||||
|
b[3] = (uint2)(a[18].x << 21 | a[18].y >> 11, a[18].y << 21 | a[18].x >> 11); |
||||
|
b[5] = (uint2)(a[3].x << 28 | a[3].y >> 4, a[3].y << 28 | a[3].x >> 4); |
||||
|
b[16] = (uint2)(a[5].y << 4 | a[5].x >> 28, a[5].x << 4 | a[5].y >> 28); |
||||
|
b[8] = (uint2)(a[16].y << 13 | a[16].x >> 19, a[16].x << 13 | a[16].y >> 19); |
||||
|
b[21] = (uint2)(a[8].y << 23 | a[8].x >> 9, a[8].x << 23 | a[8].y >> 9); |
||||
|
b[24] = (uint2)(a[21].x << 2 | a[21].y >> 30, a[21].y << 2 | a[21].x >> 30); |
||||
|
b[4] = (uint2)(a[24].x << 14 | a[24].y >> 18, a[24].y << 14 | a[24].x >> 18); |
||||
|
b[15] = (uint2)(a[4].x << 27 | a[4].y >> 5, a[4].y << 27 | a[4].x >> 5); |
||||
|
b[23] = (uint2)(a[15].y << 9 | a[15].x >> 23, a[15].x << 9 | a[15].y >> 23); |
||||
|
b[19] = (uint2)(a[23].y << 24 | a[23].x >> 8, a[23].x << 24 | a[23].y >> 8); |
||||
|
b[13] = (uint2)(a[19].x << 8 | a[19].y >> 24, a[19].y << 8 | a[19].x >> 24); |
||||
|
b[12] = (uint2)(a[13].x << 25 | a[13].y >> 7, a[13].y << 25 | a[13].x >> 7); |
||||
|
b[2] = (uint2)(a[12].y << 11 | a[12].x >> 21, a[12].x << 11 | a[12].y >> 21); |
||||
|
b[20] = (uint2)(a[2].y << 30 | a[2].x >> 2, a[2].x << 30 | a[2].y >> 2); |
||||
|
b[14] = (uint2)(a[20].x << 18 | a[20].y >> 14, a[20].y << 18 | a[20].x >> 14); |
||||
|
b[22] = (uint2)(a[14].y << 7 | a[14].x >> 25, a[14].x << 7 | a[14].y >> 25); |
||||
|
b[9] = (uint2)(a[22].y << 29 | a[22].x >> 3, a[22].x << 29 | a[22].y >> 3); |
||||
|
b[6] = (uint2)(a[9].x << 20 | a[9].y >> 12, a[9].y << 20 | a[9].x >> 12); |
||||
|
b[1] = (uint2)(a[6].y << 12 | a[6].x >> 20, a[6].x << 12 | a[6].y >> 20); |
||||
|
|
||||
|
// Chi |
||||
|
a[0] = bitselect(b[0] ^ b[2], b[0], b[1]); |
||||
|
a[1] = bitselect(b[1] ^ b[3], b[1], b[2]); |
||||
|
a[2] = bitselect(b[2] ^ b[4], b[2], b[3]); |
||||
|
a[3] = bitselect(b[3] ^ b[0], b[3], b[4]); |
||||
|
if (out_size >= 4) |
||||
|
{ |
||||
|
a[4] = bitselect(b[4] ^ b[1], b[4], b[0]); |
||||
|
a[5] = bitselect(b[5] ^ b[7], b[5], b[6]); |
||||
|
a[6] = bitselect(b[6] ^ b[8], b[6], b[7]); |
||||
|
a[7] = bitselect(b[7] ^ b[9], b[7], b[8]); |
||||
|
a[8] = bitselect(b[8] ^ b[5], b[8], b[9]); |
||||
|
if (out_size >= 8) |
||||
|
{ |
||||
|
a[9] = bitselect(b[9] ^ b[6], b[9], b[5]); |
||||
|
a[10] = bitselect(b[10] ^ b[12], b[10], b[11]); |
||||
|
a[11] = bitselect(b[11] ^ b[13], b[11], b[12]); |
||||
|
a[12] = bitselect(b[12] ^ b[14], b[12], b[13]); |
||||
|
a[13] = bitselect(b[13] ^ b[10], b[13], b[14]); |
||||
|
a[14] = bitselect(b[14] ^ b[11], b[14], b[10]); |
||||
|
a[15] = bitselect(b[15] ^ b[17], b[15], b[16]); |
||||
|
a[16] = bitselect(b[16] ^ b[18], b[16], b[17]); |
||||
|
a[17] = bitselect(b[17] ^ b[19], b[17], b[18]); |
||||
|
a[18] = bitselect(b[18] ^ b[15], b[18], b[19]); |
||||
|
a[19] = bitselect(b[19] ^ b[16], b[19], b[15]); |
||||
|
a[20] = bitselect(b[20] ^ b[22], b[20], b[21]); |
||||
|
a[21] = bitselect(b[21] ^ b[23], b[21], b[22]); |
||||
|
a[22] = bitselect(b[22] ^ b[24], b[22], b[23]); |
||||
|
a[23] = bitselect(b[23] ^ b[20], b[23], b[24]); |
||||
|
a[24] = bitselect(b[24] ^ b[21], b[24], b[20]); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
// Iota |
||||
|
a[0] ^= Keccak_f1600_RC[r]; |
||||
|
|
||||
|
#if !__ENDIAN_LITTLE__ |
||||
|
for (uint i = 0; i != 25; ++i) |
||||
|
a[i] = a[i].yx; |
||||
|
#endif |
||||
|
} |
||||
|
|
||||
|
void keccak_f1600_no_absorb(ulong* a, uint in_size, uint out_size, uint isolate) |
||||
|
{ |
||||
|
for (uint i = in_size; i != 25; ++i) |
||||
|
{ |
||||
|
a[i] = 0; |
||||
|
} |
||||
|
#if __ENDIAN_LITTLE__ |
||||
|
a[in_size] ^= 0x0000000000000001; |
||||
|
a[24-out_size*2] ^= 0x8000000000000000; |
||||
|
#else |
||||
|
a[in_size] ^= 0x0100000000000000; |
||||
|
a[24-out_size*2] ^= 0x0000000000000080; |
||||
|
#endif |
||||
|
|
||||
|
// Originally I unrolled the first and last rounds to interface |
||||
|
// better with surrounding code, however I haven't done this |
||||
|
// without causing the AMD compiler to blow up the VGPR usage. |
||||
|
uint r = 0; |
||||
|
do |
||||
|
{ |
||||
|
// This dynamic branch stops the AMD compiler unrolling the loop |
||||
|
// and additionally saves about 33% of the VGPRs, enough to gain another |
||||
|
// wavefront. Ideally we'd get 4 in flight, but 3 is the best I can |
||||
|
// massage out of the compiler. It doesn't really seem to matter how |
||||
|
// much we try and help the compiler save VGPRs because it seems to throw |
||||
|
// that information away, hence the implementation of keccak here |
||||
|
// doesn't bother. |
||||
|
if (isolate) |
||||
|
{ |
||||
|
keccak_f1600_round((uint2*)a, r++, 25); |
||||
|
} |
||||
|
} |
||||
|
while (r < 23); |
||||
|
|
||||
|
// final round optimised for digest size |
||||
|
keccak_f1600_round((uint2*)a, r++, out_size); |
||||
|
} |
||||
|
|
||||
|
#define copy(dst, src, count) for (uint i = 0; i != count; ++i) { (dst)[i] = (src)[i]; } |
||||
|
|
||||
|
#define countof(x) (sizeof(x) / sizeof(x[0])) |
||||
|
|
||||
|
uint fnv(uint x, uint y) |
||||
|
{ |
||||
|
return x * FNV_PRIME ^ y; |
||||
|
} |
||||
|
|
||||
|
uint4 fnv4(uint4 x, uint4 y) |
||||
|
{ |
||||
|
return x * FNV_PRIME ^ y; |
||||
|
} |
||||
|
|
||||
|
uint fnv_reduce(uint4 v) |
||||
|
{ |
||||
|
return fnv(fnv(fnv(v.x, v.y), v.z), v.w); |
||||
|
} |
||||
|
|
||||
|
typedef union |
||||
|
{ |
||||
|
ulong ulongs[32 / sizeof(ulong)]; |
||||
|
uint uints[32 / sizeof(uint)]; |
||||
|
} hash32_t; |
||||
|
|
||||
|
typedef union |
||||
|
{ |
||||
|
ulong ulongs[64 / sizeof(ulong)]; |
||||
|
uint4 uint4s[64 / sizeof(uint4)]; |
||||
|
} hash64_t; |
||||
|
|
||||
|
typedef union |
||||
|
{ |
||||
|
uint uints[128 / sizeof(uint)]; |
||||
|
uint4 uint4s[128 / sizeof(uint4)]; |
||||
|
} hash128_t; |
||||
|
|
||||
|
hash64_t init_hash(__constant hash32_t const* header, ulong nonce, uint isolate) |
||||
|
{ |
||||
|
hash64_t init; |
||||
|
uint const init_size = countof(init.ulongs); |
||||
|
uint const hash_size = countof(header->ulongs); |
||||
|
|
||||
|
// sha3_512(header .. nonce) |
||||
|
ulong state[25]; |
||||
|
copy(state, header->ulongs, hash_size); |
||||
|
state[hash_size] = nonce; |
||||
|
keccak_f1600_no_absorb(state, hash_size + 1, init_size, isolate); |
||||
|
|
||||
|
copy(init.ulongs, state, init_size); |
||||
|
return init; |
||||
|
} |
||||
|
|
||||
|
uint inner_loop(uint4 init, uint thread_id, __local uint* share, __global hash128_t const* g_dag, uint isolate) |
||||
|
{ |
||||
|
uint4 mix = init; |
||||
|
|
||||
|
// share init0 |
||||
|
if (thread_id == 0) |
||||
|
*share = mix.x; |
||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
uint init0 = *share; |
||||
|
|
||||
|
uint a = 0; |
||||
|
do |
||||
|
{ |
||||
|
bool update_share = thread_id == (a/4) % THREADS_PER_HASH; |
||||
|
|
||||
|
#pragma unroll |
||||
|
for (uint i = 0; i != 4; ++i) |
||||
|
{ |
||||
|
if (update_share) |
||||
|
{ |
||||
|
uint m[4] = { mix.x, mix.y, mix.z, mix.w }; |
||||
|
*share = fnv(init0 ^ (a+i), m[i]) % DAG_SIZE; |
||||
|
} |
||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
|
||||
|
mix = fnv4(mix, g_dag[*share].uint4s[thread_id]); |
||||
|
} |
||||
|
} |
||||
|
while ((a += 4) != (ACCESSES & isolate)); |
||||
|
|
||||
|
return fnv_reduce(mix); |
||||
|
} |
||||
|
|
||||
|
hash32_t final_hash(hash64_t const* init, hash32_t const* mix, uint isolate) |
||||
|
{ |
||||
|
ulong state[25]; |
||||
|
|
||||
|
hash32_t hash; |
||||
|
uint const hash_size = countof(hash.ulongs); |
||||
|
uint const init_size = countof(init->ulongs); |
||||
|
uint const mix_size = countof(mix->ulongs); |
||||
|
|
||||
|
// keccak_256(keccak_512(header..nonce) .. mix); |
||||
|
copy(state, init->ulongs, init_size); |
||||
|
copy(state + init_size, mix->ulongs, mix_size); |
||||
|
keccak_f1600_no_absorb(state, init_size+mix_size, hash_size, isolate); |
||||
|
|
||||
|
// copy out |
||||
|
copy(hash.ulongs, state, hash_size); |
||||
|
return hash; |
||||
|
} |
||||
|
|
||||
|
hash32_t compute_hash_simple( |
||||
|
__constant hash32_t const* g_header, |
||||
|
__global hash128_t const* g_dag, |
||||
|
ulong nonce, |
||||
|
uint isolate |
||||
|
) |
||||
|
{ |
||||
|
hash64_t init = init_hash(g_header, nonce, isolate); |
||||
|
|
||||
|
hash128_t mix; |
||||
|
for (uint i = 0; i != countof(mix.uint4s); ++i) |
||||
|
{ |
||||
|
mix.uint4s[i] = init.uint4s[i % countof(init.uint4s)]; |
||||
|
} |
||||
|
|
||||
|
uint mix_val = mix.uints[0]; |
||||
|
uint init0 = mix.uints[0]; |
||||
|
uint a = 0; |
||||
|
do |
||||
|
{ |
||||
|
uint pi = fnv(init0 ^ a, mix_val) % DAG_SIZE; |
||||
|
uint n = (a+1) % countof(mix.uints); |
||||
|
|
||||
|
#pragma unroll |
||||
|
for (uint i = 0; i != countof(mix.uints); ++i) |
||||
|
{ |
||||
|
mix.uints[i] = fnv(mix.uints[i], g_dag[pi].uints[i]); |
||||
|
mix_val = i == n ? mix.uints[i] : mix_val; |
||||
|
} |
||||
|
} |
||||
|
while (++a != (ACCESSES & isolate)); |
||||
|
|
||||
|
// reduce to output |
||||
|
hash32_t fnv_mix; |
||||
|
for (uint i = 0; i != countof(fnv_mix.uints); ++i) |
||||
|
{ |
||||
|
fnv_mix.uints[i] = fnv_reduce(mix.uint4s[i]); |
||||
|
} |
||||
|
|
||||
|
return final_hash(&init, &fnv_mix, isolate); |
||||
|
} |
||||
|
|
||||
|
typedef union |
||||
|
{ |
||||
|
struct |
||||
|
{ |
||||
|
hash64_t init; |
||||
|
uint pad; // avoid lds bank conflicts |
||||
|
}; |
||||
|
hash32_t mix; |
||||
|
} compute_hash_share; |
||||
|
|
||||
|
hash32_t compute_hash( |
||||
|
__local compute_hash_share* share, |
||||
|
__constant hash32_t const* g_header, |
||||
|
__global hash128_t const* g_dag, |
||||
|
ulong nonce, |
||||
|
uint isolate |
||||
|
) |
||||
|
{ |
||||
|
uint const gid = get_global_id(0); |
||||
|
|
||||
|
// Compute one init hash per work item. |
||||
|
hash64_t init = init_hash(g_header, nonce, isolate); |
||||
|
|
||||
|
// Threads work together in this phase in groups of 8. |
||||
|
uint const thread_id = gid % THREADS_PER_HASH; |
||||
|
uint const hash_id = (gid % GROUP_SIZE) / THREADS_PER_HASH; |
||||
|
|
||||
|
hash32_t mix; |
||||
|
uint i = 0; |
||||
|
do |
||||
|
{ |
||||
|
// share init with other threads |
||||
|
if (i == thread_id) |
||||
|
share[hash_id].init = init; |
||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
|
||||
|
uint4 thread_init = share[hash_id].init.uint4s[thread_id % (64 / sizeof(uint4))]; |
||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
|
||||
|
uint thread_mix = inner_loop(thread_init, thread_id, share[hash_id].mix.uints, g_dag, isolate); |
||||
|
|
||||
|
share[hash_id].mix.uints[thread_id] = thread_mix; |
||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
|
||||
|
if (i == thread_id) |
||||
|
mix = share[hash_id].mix; |
||||
|
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
} |
||||
|
while (++i != (THREADS_PER_HASH & isolate)); |
||||
|
|
||||
|
return final_hash(&init, &mix, isolate); |
||||
|
} |
||||
|
|
||||
|
__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) |
||||
|
__kernel void ethash_hash_simple( |
||||
|
__global hash32_t* g_hashes, |
||||
|
__constant hash32_t const* g_header, |
||||
|
__global hash128_t const* g_dag, |
||||
|
ulong start_nonce, |
||||
|
uint isolate |
||||
|
) |
||||
|
{ |
||||
|
uint const gid = get_global_id(0); |
||||
|
g_hashes[gid] = compute_hash_simple(g_header, g_dag, start_nonce + gid, isolate); |
||||
|
} |
||||
|
|
||||
|
__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) |
||||
|
__kernel void ethash_search_simple( |
||||
|
__global volatile uint* restrict g_output, |
||||
|
__constant hash32_t const* g_header, |
||||
|
__global hash128_t const* g_dag, |
||||
|
ulong start_nonce, |
||||
|
ulong target, |
||||
|
uint isolate |
||||
|
) |
||||
|
{ |
||||
|
uint const gid = get_global_id(0); |
||||
|
hash32_t hash = compute_hash_simple(g_header, g_dag, start_nonce + gid, isolate); |
||||
|
if (as_ulong(as_uchar8(hash.ulongs[0]).s76543210) < target) |
||||
|
{ |
||||
|
uint slot = min(MAX_OUTPUTS, atomic_inc(&g_output[0]) + 1); |
||||
|
g_output[slot] = gid; |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) |
||||
|
__kernel void ethash_hash( |
||||
|
__global hash32_t* g_hashes, |
||||
|
__constant hash32_t const* g_header, |
||||
|
__global hash128_t const* g_dag, |
||||
|
ulong start_nonce, |
||||
|
uint isolate |
||||
|
) |
||||
|
{ |
||||
|
__local compute_hash_share share[HASHES_PER_LOOP]; |
||||
|
|
||||
|
uint const gid = get_global_id(0); |
||||
|
g_hashes[gid] = compute_hash(share, g_header, g_dag, start_nonce + gid, isolate); |
||||
|
} |
||||
|
|
||||
|
__attribute__((reqd_work_group_size(GROUP_SIZE, 1, 1))) |
||||
|
__kernel void ethash_search( |
||||
|
__global volatile uint* restrict g_output, |
||||
|
__constant hash32_t const* g_header, |
||||
|
__global hash128_t const* g_dag, |
||||
|
ulong start_nonce, |
||||
|
ulong target, |
||||
|
uint isolate |
||||
|
) |
||||
|
{ |
||||
|
__local compute_hash_share share[HASHES_PER_LOOP]; |
||||
|
|
||||
|
uint const gid = get_global_id(0); |
||||
|
hash32_t hash = compute_hash(share, g_header, g_dag, start_nonce + gid, isolate); |
||||
|
|
||||
|
if (as_ulong(as_uchar8(hash.ulongs[0]).s76543210) < target) |
||||
|
{ |
||||
|
uint slot = min(MAX_OUTPUTS, atomic_inc(&g_output[0]) + 1); |
||||
|
g_output[slot] = gid; |
||||
|
} |
||||
|
} |
||||
@ -0,0 +1,328 @@ |
|||||
|
/*
|
||||
|
This file is part of cpp-ethereum. |
||||
|
|
||||
|
cpp-ethereum 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. |
||||
|
|
||||
|
cpp-ethereum 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.cpp
|
||||
|
* @author Gav Wood <i@gavwood.com> |
||||
|
* @date 2014 |
||||
|
*/ |
||||
|
|
||||
|
#include "Ethash.h" |
||||
|
|
||||
|
#include <boost/detail/endian.hpp> |
||||
|
#include <boost/filesystem.hpp> |
||||
|
#include <boost/algorithm/string.hpp> |
||||
|
#include <chrono> |
||||
|
#include <array> |
||||
|
#include <thread> |
||||
|
#include <random> |
||||
|
#include <thread> |
||||
|
#include <libdevcore/Guards.h> |
||||
|
#include <libdevcore/Log.h> |
||||
|
#include <libdevcore/Common.h> |
||||
|
#include <libdevcore/CommonIO.h> |
||||
|
#include <libdevcrypto/CryptoPP.h> |
||||
|
#include <libdevcrypto/FileSystem.h> |
||||
|
#include <libethash/ethash.h> |
||||
|
#if ETH_ETHASHCL || !ETH_TRUE |
||||
|
#include <libethash-cl/ethash_cl_miner.h> |
||||
|
#endif |
||||
|
#if ETH_CPUID || !ETH_TRUE |
||||
|
#define HAVE_STDINT_H |
||||
|
#include <libcpuid/libcpuid.h> |
||||
|
#endif |
||||
|
#include "BlockInfo.h" |
||||
|
#include "EthashAux.h" |
||||
|
using namespace std; |
||||
|
using namespace std::chrono; |
||||
|
|
||||
|
namespace dev |
||||
|
{ |
||||
|
namespace eth |
||||
|
{ |
||||
|
|
||||
|
const Ethash::WorkPackage Ethash::NullWorkPackage = Ethash::WorkPackage(); |
||||
|
|
||||
|
std::string Ethash::name() |
||||
|
{ |
||||
|
return "Ethash"; |
||||
|
} |
||||
|
|
||||
|
unsigned Ethash::revision() |
||||
|
{ |
||||
|
return ETHASH_REVISION; |
||||
|
} |
||||
|
|
||||
|
Ethash::WorkPackage Ethash::package(BlockInfo const& _bi) |
||||
|
{ |
||||
|
WorkPackage ret; |
||||
|
ret.boundary = _bi.boundary(); |
||||
|
ret.headerHash = _bi.headerHash(WithoutNonce); |
||||
|
ret.seedHash = _bi.seedHash(); |
||||
|
return ret; |
||||
|
} |
||||
|
|
||||
|
void Ethash::prep(BlockInfo const& _header) |
||||
|
{ |
||||
|
EthashAux::full(_header); |
||||
|
} |
||||
|
|
||||
|
bool Ethash::preVerify(BlockInfo const& _header) |
||||
|
{ |
||||
|
if (_header.number >= ETHASH_EPOCH_LENGTH * 2048) |
||||
|
return false; |
||||
|
|
||||
|
h256 boundary = u256((bigint(1) << 256) / _header.difficulty); |
||||
|
|
||||
|
return !!ethash_quick_check_difficulty( |
||||
|
_header.headerHash(WithoutNonce).data(), |
||||
|
(uint64_t)(u64)_header.nonce, |
||||
|
_header.mixHash.data(), |
||||
|
boundary.data()); |
||||
|
} |
||||
|
|
||||
|
bool Ethash::verify(BlockInfo const& _header) |
||||
|
{ |
||||
|
bool pre = preVerify(_header); |
||||
|
#if !ETH_DEBUG |
||||
|
if (!pre) |
||||
|
return false; |
||||
|
#endif |
||||
|
|
||||
|
auto result = EthashAux::eval(_header); |
||||
|
bool slow = result.value <= _header.boundary() && result.mixHash == _header.mixHash; |
||||
|
|
||||
|
#if ETH_DEBUG || !ETH_TRUE |
||||
|
if (!pre && slow) |
||||
|
{ |
||||
|
cwarn << "WARNING: evaluated result gives true whereas ethash_quick_check_difficulty gives false."; |
||||
|
cwarn << "headerHash:" << _header.headerHash(WithoutNonce); |
||||
|
cwarn << "nonce:" << _header.nonce; |
||||
|
cwarn << "mixHash:" << _header.mixHash; |
||||
|
cwarn << "difficulty:" << _header.difficulty; |
||||
|
cwarn << "boundary:" << _header.boundary(); |
||||
|
cwarn << "result.value:" << result.value; |
||||
|
cwarn << "result.mixHash:" << result.mixHash; |
||||
|
} |
||||
|
#endif |
||||
|
|
||||
|
return slow; |
||||
|
} |
||||
|
|
||||
|
void Ethash::CPUMiner::workLoop() |
||||
|
{ |
||||
|
auto tid = std::this_thread::get_id(); |
||||
|
static std::mt19937_64 s_eng((time(0) + std::hash<decltype(tid)>()(tid))); |
||||
|
|
||||
|
uint64_t tryNonce = (uint64_t)(u64)Nonce::random(s_eng); |
||||
|
ethash_return_value ethashReturn; |
||||
|
|
||||
|
WorkPackage w = work(); |
||||
|
|
||||
|
auto p = EthashAux::params(w.seedHash); |
||||
|
void const* dagPointer = EthashAux::full(w.seedHash).data(); |
||||
|
uint8_t const* headerHashPointer = w.headerHash.data(); |
||||
|
h256 boundary = w.boundary; |
||||
|
unsigned hashCount = 1; |
||||
|
for (; !shouldStop(); tryNonce++, hashCount++) |
||||
|
{ |
||||
|
ethash_compute_full(ðashReturn, dagPointer, &p, headerHashPointer, tryNonce); |
||||
|
h256 value = h256(ethashReturn.result, h256::ConstructFromPointer); |
||||
|
if (value <= boundary && submitProof(Solution{(Nonce)(u64)tryNonce, h256(ethashReturn.mix_hash, h256::ConstructFromPointer)})) |
||||
|
break; |
||||
|
if (!(hashCount % 1000)) |
||||
|
accumulateHashes(1000); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
static string jsonEncode(map<string, string> const& _m) |
||||
|
{ |
||||
|
string ret = "{"; |
||||
|
|
||||
|
for (auto const& i: _m) |
||||
|
{ |
||||
|
string k = boost::replace_all_copy(boost::replace_all_copy(i.first, "\\", "\\\\"), "'", "\\'"); |
||||
|
string v = boost::replace_all_copy(boost::replace_all_copy(i.second, "\\", "\\\\"), "'", "\\'"); |
||||
|
if (ret.size() > 1) |
||||
|
ret += ", "; |
||||
|
ret += "\"" + k + "\":\"" + v + "\""; |
||||
|
} |
||||
|
|
||||
|
return ret + "}"; |
||||
|
} |
||||
|
|
||||
|
std::string Ethash::CPUMiner::platformInfo() |
||||
|
{ |
||||
|
string baseline = toString(std::thread::hardware_concurrency()) + "-thread CPU"; |
||||
|
#if ETH_CPUID || !ETH_TRUE |
||||
|
if (!cpuid_present()) |
||||
|
return baseline; |
||||
|
struct cpu_raw_data_t raw; |
||||
|
struct cpu_id_t data; |
||||
|
if (cpuid_get_raw_data(&raw) < 0) |
||||
|
return baseline; |
||||
|
if (cpu_identify(&raw, &data) < 0) |
||||
|
return baseline; |
||||
|
map<string, string> m; |
||||
|
m["vendor"] = data.vendor_str; |
||||
|
m["codename"] = data.cpu_codename; |
||||
|
m["brand"] = data.brand_str; |
||||
|
m["L1 cache"] = toString(data.l1_data_cache); |
||||
|
m["L2 cache"] = toString(data.l2_cache); |
||||
|
m["L3 cache"] = toString(data.l3_cache); |
||||
|
m["cores"] = toString(data.num_cores); |
||||
|
m["threads"] = toString(data.num_logical_cpus); |
||||
|
m["clocknominal"] = toString(cpu_clock_by_os()); |
||||
|
m["clocktested"] = toString(cpu_clock_measure(200, 0)); |
||||
|
/*
|
||||
|
printf(" MMX : %s\n", data.flags[CPU_FEATURE_MMX] ? "present" : "absent"); |
||||
|
printf(" MMX-extended: %s\n", data.flags[CPU_FEATURE_MMXEXT] ? "present" : "absent"); |
||||
|
printf(" SSE : %s\n", data.flags[CPU_FEATURE_SSE] ? "present" : "absent"); |
||||
|
printf(" SSE2 : %s\n", data.flags[CPU_FEATURE_SSE2] ? "present" : "absent"); |
||||
|
printf(" 3DNow! : %s\n", data.flags[CPU_FEATURE_3DNOW] ? "present" : "absent"); |
||||
|
*/ |
||||
|
return jsonEncode(m); |
||||
|
#else |
||||
|
return baseline; |
||||
|
#endif |
||||
|
} |
||||
|
|
||||
|
#if ETH_ETHASHCL || !ETH_TRUE |
||||
|
|
||||
|
class EthashCLHook: public ethash_cl_miner::search_hook |
||||
|
{ |
||||
|
public: |
||||
|
EthashCLHook(Ethash::GPUMiner* _owner): m_owner(_owner) {} |
||||
|
|
||||
|
void abort() |
||||
|
{ |
||||
|
Guard l(x_all); |
||||
|
if (m_aborted) |
||||
|
return; |
||||
|
// cdebug << "Attempting to abort";
|
||||
|
m_abort = true; |
||||
|
for (unsigned timeout = 0; timeout < 100 && !m_aborted; ++timeout) |
||||
|
std::this_thread::sleep_for(chrono::milliseconds(30)); |
||||
|
// if (!m_aborted)
|
||||
|
// cwarn << "Couldn't abort. Abandoning OpenCL process.";
|
||||
|
} |
||||
|
|
||||
|
void reset() |
||||
|
{ |
||||
|
m_aborted = m_abort = false; |
||||
|
} |
||||
|
|
||||
|
protected: |
||||
|
virtual bool found(uint64_t const* _nonces, uint32_t _count) override |
||||
|
{ |
||||
|
// dev::operator <<(std::cerr << "Found nonces: ", vector<uint64_t>(_nonces, _nonces + _count)) << std::endl;
|
||||
|
for (uint32_t i = 0; i < _count; ++i) |
||||
|
{ |
||||
|
if (m_owner->report(_nonces[i])) |
||||
|
{ |
||||
|
m_aborted = true; |
||||
|
return true; |
||||
|
} |
||||
|
} |
||||
|
return false; |
||||
|
} |
||||
|
|
||||
|
virtual bool searched(uint64_t _startNonce, uint32_t _count) override |
||||
|
{ |
||||
|
Guard l(x_all); |
||||
|
// std::cerr << "Searched " << _count << " from " << _startNonce << std::endl;
|
||||
|
m_owner->accumulateHashes(_count); |
||||
|
m_last = _startNonce + _count; |
||||
|
if (m_abort) |
||||
|
{ |
||||
|
m_aborted = true; |
||||
|
return true; |
||||
|
} |
||||
|
return false; |
||||
|
} |
||||
|
|
||||
|
private: |
||||
|
Mutex x_all; |
||||
|
uint64_t m_last; |
||||
|
bool m_abort = false; |
||||
|
bool m_aborted = true; |
||||
|
Ethash::GPUMiner* m_owner = nullptr; |
||||
|
}; |
||||
|
|
||||
|
unsigned Ethash::GPUMiner::s_deviceId = 0; |
||||
|
|
||||
|
Ethash::GPUMiner::GPUMiner(ConstructionInfo const& _ci): |
||||
|
Miner(_ci), |
||||
|
m_hook(new EthashCLHook(this)) |
||||
|
{ |
||||
|
} |
||||
|
|
||||
|
Ethash::GPUMiner::~GPUMiner() |
||||
|
{ |
||||
|
pause(); |
||||
|
delete m_miner; |
||||
|
delete m_hook; |
||||
|
} |
||||
|
|
||||
|
bool Ethash::GPUMiner::report(uint64_t _nonce) |
||||
|
{ |
||||
|
Nonce n = (Nonce)(u64)_nonce; |
||||
|
Result r = EthashAux::eval(work().seedHash, work().headerHash, n); |
||||
|
if (r.value < work().boundary) |
||||
|
return submitProof(Solution{n, r.mixHash}); |
||||
|
return false; |
||||
|
} |
||||
|
|
||||
|
void Ethash::GPUMiner::kickOff() |
||||
|
{ |
||||
|
m_hook->reset(); |
||||
|
startWorking(); |
||||
|
} |
||||
|
|
||||
|
void Ethash::GPUMiner::workLoop() |
||||
|
{ |
||||
|
// take local copy of work since it may end up being overwritten by kickOff/pause.
|
||||
|
WorkPackage w = work(); |
||||
|
if (!m_miner || m_minerSeed != w.seedHash) |
||||
|
{ |
||||
|
m_minerSeed = w.seedHash; |
||||
|
|
||||
|
delete m_miner; |
||||
|
m_miner = new ethash_cl_miner; |
||||
|
|
||||
|
auto p = EthashAux::params(m_minerSeed); |
||||
|
auto cb = [&](void* d) { EthashAux::full(m_minerSeed, bytesRef((byte*)d, p.full_size)); }; |
||||
|
m_miner->init(p, cb, 32, s_deviceId); |
||||
|
} |
||||
|
|
||||
|
uint64_t upper64OfBoundary = (uint64_t)(u64)((u256)w.boundary >> 192); |
||||
|
m_miner->search(w.headerHash.data(), upper64OfBoundary, *m_hook); |
||||
|
} |
||||
|
|
||||
|
void Ethash::GPUMiner::pause() |
||||
|
{ |
||||
|
m_hook->abort(); |
||||
|
stopWorking(); |
||||
|
} |
||||
|
|
||||
|
std::string Ethash::GPUMiner::platformInfo() |
||||
|
{ |
||||
|
return ethash_cl_miner::platform_info(); |
||||
|
} |
||||
|
|
||||
|
#endif |
||||
|
|
||||
|
} |
||||
|
} |
||||
@ -0,0 +1,140 @@ |
|||||
|
/*
|
||||
|
This file is part of cpp-ethereum. |
||||
|
|
||||
|
cpp-ethereum 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. |
||||
|
|
||||
|
cpp-ethereum 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.h
|
||||
|
* @author Gav Wood <i@gavwood.com> |
||||
|
* @date 2014 |
||||
|
* |
||||
|
* A proof of work algorithm. |
||||
|
*/ |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <chrono> |
||||
|
#include <thread> |
||||
|
#include <cstdint> |
||||
|
#include <libdevcore/CommonIO.h> |
||||
|
#include "Common.h" |
||||
|
#include "BlockInfo.h" |
||||
|
#include "Miner.h" |
||||
|
|
||||
|
class ethash_cl_miner; |
||||
|
|
||||
|
namespace dev |
||||
|
{ |
||||
|
namespace eth |
||||
|
{ |
||||
|
|
||||
|
class EthashCLHook; |
||||
|
|
||||
|
class Ethash |
||||
|
{ |
||||
|
public: |
||||
|
using Miner = GenericMiner<Ethash>; |
||||
|
|
||||
|
struct Solution |
||||
|
{ |
||||
|
Nonce nonce; |
||||
|
h256 mixHash; |
||||
|
}; |
||||
|
|
||||
|
struct Result |
||||
|
{ |
||||
|
h256 value; |
||||
|
h256 mixHash; |
||||
|
}; |
||||
|
|
||||
|
struct WorkPackage |
||||
|
{ |
||||
|
WorkPackage() = default; |
||||
|
|
||||
|
void reset() { headerHash = h256(); } |
||||
|
operator bool() const { return headerHash != h256(); } |
||||
|
|
||||
|
h256 boundary; |
||||
|
h256 headerHash; ///< When h256() means "pause until notified a new work package is available".
|
||||
|
h256 seedHash; |
||||
|
}; |
||||
|
|
||||
|
static const WorkPackage NullWorkPackage; |
||||
|
|
||||
|
static std::string name(); |
||||
|
static unsigned revision(); |
||||
|
static void prep(BlockInfo const& _header); |
||||
|
static bool verify(BlockInfo const& _header); |
||||
|
static bool preVerify(BlockInfo const& _header); |
||||
|
static WorkPackage package(BlockInfo const& _header); |
||||
|
static void assignResult(Solution const& _r, BlockInfo& _header) { _header.nonce = _r.nonce; _header.mixHash = _r.mixHash; } |
||||
|
|
||||
|
class CPUMiner: public Miner, Worker |
||||
|
{ |
||||
|
public: |
||||
|
CPUMiner(ConstructionInfo const& _ci): Miner(_ci), Worker("miner" + toString(index())) {} |
||||
|
|
||||
|
static unsigned instances() { return std::thread::hardware_concurrency(); } |
||||
|
static std::string platformInfo(); |
||||
|
static void setDefaultDevice(unsigned) {} |
||||
|
|
||||
|
protected: |
||||
|
void kickOff() override |
||||
|
{ |
||||
|
stopWorking(); |
||||
|
startWorking(); |
||||
|
} |
||||
|
|
||||
|
void pause() override { stopWorking(); } |
||||
|
|
||||
|
private: |
||||
|
void workLoop() override; |
||||
|
static unsigned s_deviceId; |
||||
|
}; |
||||
|
|
||||
|
#if ETH_ETHASHCL || !ETH_TRUE |
||||
|
class GPUMiner: public Miner, Worker |
||||
|
{ |
||||
|
friend class dev::eth::EthashCLHook; |
||||
|
|
||||
|
public: |
||||
|
GPUMiner(ConstructionInfo const& _ci); |
||||
|
~GPUMiner(); |
||||
|
|
||||
|
static unsigned instances() { return 1; } |
||||
|
static std::string platformInfo(); |
||||
|
static void setDefaultDevice(unsigned _id) { s_deviceId = _id; } |
||||
|
|
||||
|
protected: |
||||
|
void kickOff() override; |
||||
|
void pause() override; |
||||
|
|
||||
|
private: |
||||
|
void workLoop() override; |
||||
|
bool report(uint64_t _nonce); |
||||
|
|
||||
|
using Miner::accumulateHashes; |
||||
|
|
||||
|
EthashCLHook* m_hook = nullptr; |
||||
|
ethash_cl_miner* m_miner = nullptr; |
||||
|
|
||||
|
h256 m_minerSeed; ///< Last seed in m_miner
|
||||
|
static unsigned s_deviceId; |
||||
|
}; |
||||
|
#else |
||||
|
using GPUMiner = CPUMiner; |
||||
|
#endif |
||||
|
}; |
||||
|
|
||||
|
} |
||||
|
} |
||||
@ -0,0 +1,214 @@ |
|||||
|
/*
|
||||
|
This file is part of cpp-ethereum. |
||||
|
|
||||
|
cpp-ethereum 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. |
||||
|
|
||||
|
cpp-ethereum 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 EthashAux.cpp
|
||||
|
* @author Gav Wood <i@gavwood.com> |
||||
|
* @date 2014 |
||||
|
*/ |
||||
|
|
||||
|
#include "EthashAux.h" |
||||
|
|
||||
|
#include <boost/detail/endian.hpp> |
||||
|
#include <boost/filesystem.hpp> |
||||
|
#include <chrono> |
||||
|
#include <array> |
||||
|
#include <random> |
||||
|
#include <thread> |
||||
|
#include <libdevcore/Common.h> |
||||
|
#include <libdevcore/Guards.h> |
||||
|
#include <libdevcore/Log.h> |
||||
|
#include <libdevcrypto/CryptoPP.h> |
||||
|
#include <libdevcrypto/SHA3.h> |
||||
|
#include <libdevcrypto/FileSystem.h> |
||||
|
#include <libethcore/Params.h> |
||||
|
#include "BlockInfo.h" |
||||
|
using namespace std; |
||||
|
using namespace chrono; |
||||
|
using namespace dev; |
||||
|
using namespace eth; |
||||
|
|
||||
|
#define ETH_IGNORE_EXCEPTIONS(X) try { X; } catch (...) {} |
||||
|
|
||||
|
EthashAux* dev::eth::EthashAux::s_this = nullptr; |
||||
|
|
||||
|
EthashAux::~EthashAux() |
||||
|
{ |
||||
|
while (!m_lights.empty()) |
||||
|
killCache(m_lights.begin()->first); |
||||
|
} |
||||
|
|
||||
|
ethash_params EthashAux::params(BlockInfo const& _header) |
||||
|
{ |
||||
|
return params((unsigned)_header.number); |
||||
|
} |
||||
|
|
||||
|
ethash_params EthashAux::params(unsigned _n) |
||||
|
{ |
||||
|
ethash_params p; |
||||
|
p.cache_size = ethash_get_cachesize(_n); |
||||
|
p.full_size = ethash_get_datasize(_n); |
||||
|
return p; |
||||
|
} |
||||
|
|
||||
|
h256 EthashAux::seedHash(unsigned _number) |
||||
|
{ |
||||
|
unsigned epoch = _number / ETHASH_EPOCH_LENGTH; |
||||
|
RecursiveGuard l(get()->x_this); |
||||
|
if (epoch >= get()->m_seedHashes.size()) |
||||
|
{ |
||||
|
h256 ret; |
||||
|
unsigned n = 0; |
||||
|
if (!get()->m_seedHashes.empty()) |
||||
|
{ |
||||
|
ret = get()->m_seedHashes.back(); |
||||
|
n = get()->m_seedHashes.size() - 1; |
||||
|
} |
||||
|
get()->m_seedHashes.resize(epoch + 1); |
||||
|
cdebug << "Searching for seedHash of epoch " << epoch; |
||||
|
for (; n <= epoch; ++n, ret = sha3(ret)) |
||||
|
{ |
||||
|
get()->m_seedHashes[n] = ret; |
||||
|
cdebug << "Epoch" << n << "is" << ret.abridged(); |
||||
|
} |
||||
|
} |
||||
|
return get()->m_seedHashes[epoch]; |
||||
|
} |
||||
|
|
||||
|
ethash_params EthashAux::params(h256 const& _seedHash) |
||||
|
{ |
||||
|
RecursiveGuard l(get()->x_this); |
||||
|
unsigned epoch = 0; |
||||
|
try |
||||
|
{ |
||||
|
epoch = get()->m_epochs.at(_seedHash); |
||||
|
} |
||||
|
catch (...) |
||||
|
{ |
||||
|
cdebug << "Searching for seedHash " << _seedHash.abridged(); |
||||
|
for (h256 h; h != _seedHash && epoch < 2048; ++epoch, h = sha3(h), get()->m_epochs[h] = epoch) {} |
||||
|
if (epoch == 2048) |
||||
|
{ |
||||
|
std::ostringstream error; |
||||
|
error << "apparent block number for " << _seedHash.abridged() << " is too high; max is " << (ETHASH_EPOCH_LENGTH * 2048); |
||||
|
throw std::invalid_argument(error.str()); |
||||
|
} |
||||
|
} |
||||
|
return params(epoch * ETHASH_EPOCH_LENGTH); |
||||
|
} |
||||
|
|
||||
|
void EthashAux::killCache(h256 const& _s) |
||||
|
{ |
||||
|
RecursiveGuard l(x_this); |
||||
|
if (m_lights.count(_s)) |
||||
|
{ |
||||
|
ethash_delete_light(m_lights.at(_s)); |
||||
|
m_lights.erase(_s); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
void const* EthashAux::light(BlockInfo const& _header) |
||||
|
{ |
||||
|
return light(_header.seedHash()); |
||||
|
} |
||||
|
|
||||
|
void const* EthashAux::light(h256 const& _seedHash) |
||||
|
{ |
||||
|
RecursiveGuard l(get()->x_this); |
||||
|
if (!get()->m_lights.count(_seedHash)) |
||||
|
{ |
||||
|
ethash_params p = params(_seedHash); |
||||
|
get()->m_lights[_seedHash] = ethash_new_light(&p, _seedHash.data()); |
||||
|
} |
||||
|
return get()->m_lights[_seedHash]; |
||||
|
} |
||||
|
|
||||
|
bytesConstRef EthashAux::full(BlockInfo const& _header, bytesRef _dest) |
||||
|
{ |
||||
|
return full(_header.seedHash(), _dest); |
||||
|
} |
||||
|
|
||||
|
bytesConstRef EthashAux::full(h256 const& _seedHash, bytesRef _dest) |
||||
|
{ |
||||
|
RecursiveGuard l(get()->x_this); |
||||
|
if (get()->m_fulls.count(_seedHash) && _dest) |
||||
|
{ |
||||
|
assert(get()->m_fulls.size() <= _dest.size()); |
||||
|
get()->m_fulls.at(_seedHash).copyTo(_dest); |
||||
|
return _dest; |
||||
|
} |
||||
|
if (!get()->m_fulls.count(_seedHash)) |
||||
|
{ |
||||
|
// @memoryleak @bug place it on a pile for deletion - perhaps use shared_ptr.
|
||||
|
/* if (!m_fulls.empty())
|
||||
|
{ |
||||
|
delete [] m_fulls.begin()->second.data(); |
||||
|
m_fulls.erase(m_fulls.begin()); |
||||
|
}*/ |
||||
|
|
||||
|
try { |
||||
|
boost::filesystem::create_directories(getDataDir("ethash")); |
||||
|
} catch (...) {} |
||||
|
|
||||
|
auto info = rlpList(Ethash::revision(), _seedHash); |
||||
|
std::string oldMemoFile = getDataDir("ethash") + "/full"; |
||||
|
std::string memoFile = getDataDir("ethash") + "/full-R" + toString(ETHASH_REVISION) + "-" + toHex(_seedHash.ref().cropped(0, 8)); |
||||
|
if (boost::filesystem::exists(oldMemoFile) && contents(oldMemoFile + ".info") == info) |
||||
|
{ |
||||
|
// memofile valid - rename.
|
||||
|
boost::filesystem::rename(oldMemoFile, memoFile); |
||||
|
} |
||||
|
|
||||
|
ETH_IGNORE_EXCEPTIONS(boost::filesystem::remove(oldMemoFile)); |
||||
|
ETH_IGNORE_EXCEPTIONS(boost::filesystem::remove(oldMemoFile + ".info")); |
||||
|
|
||||
|
ethash_params p = params(_seedHash); |
||||
|
assert(!_dest || _dest.size() >= p.full_size); // must be big enough.
|
||||
|
|
||||
|
bytesRef r = contentsNew(memoFile, _dest); |
||||
|
if (!r) |
||||
|
{ |
||||
|
// file didn't exist.
|
||||
|
if (_dest) |
||||
|
// buffer was passed in - no insertion into cache nor need to allocate
|
||||
|
r = _dest; |
||||
|
else |
||||
|
r = bytesRef(new byte[p.full_size], p.full_size); |
||||
|
ethash_prep_full(r.data(), &p, light(_seedHash)); |
||||
|
writeFile(memoFile, r); |
||||
|
} |
||||
|
if (_dest) |
||||
|
return _dest; |
||||
|
get()->m_fulls[_seedHash] = r; |
||||
|
} |
||||
|
return get()->m_fulls[_seedHash]; |
||||
|
} |
||||
|
|
||||
|
Ethash::Result EthashAux::eval(BlockInfo const& _header, Nonce const& _nonce) |
||||
|
{ |
||||
|
return eval(_header.seedHash(), _header.headerHash(WithoutNonce), _nonce); |
||||
|
} |
||||
|
|
||||
|
Ethash::Result EthashAux::eval(h256 const& _seedHash, h256 const& _headerHash, Nonce const& _nonce) |
||||
|
{ |
||||
|
auto p = EthashAux::params(_seedHash); |
||||
|
ethash_return_value r; |
||||
|
if (EthashAux::get()->m_fulls.count(_seedHash)) |
||||
|
ethash_compute_full(&r, EthashAux::get()->full(_seedHash).data(), &p, _headerHash.data(), (uint64_t)(u64)_nonce); |
||||
|
else |
||||
|
ethash_compute_light(&r, EthashAux::get()->light(_seedHash), &p, _headerHash.data(), (uint64_t)(u64)_nonce); |
||||
|
// cdebug << "EthashAux::eval sha3(cache):" << sha3(EthashAux::get()->cache(_header)) << "hh:" << _header.headerHash(WithoutNonce) << "nonce:" << _nonce << " => " << h256(r.result, h256::ConstructFromPointer);
|
||||
|
return Ethash::Result{h256(r.result, h256::ConstructFromPointer), h256(r.mix_hash, h256::ConstructFromPointer)}; |
||||
|
} |
||||
@ -0,0 +1,67 @@ |
|||||
|
/*
|
||||
|
This file is part of cpp-ethereum. |
||||
|
|
||||
|
cpp-ethereum 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. |
||||
|
|
||||
|
cpp-ethereum 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 EthashAux.cpp
|
||||
|
* @author Gav Wood <i@gavwood.com> |
||||
|
* @date 2014 |
||||
|
*/ |
||||
|
|
||||
|
#include <libethash/ethash.h> |
||||
|
#include "Ethash.h" |
||||
|
|
||||
|
namespace dev |
||||
|
{ |
||||
|
namespace eth{ |
||||
|
|
||||
|
class EthashAux |
||||
|
{ |
||||
|
public: |
||||
|
~EthashAux(); |
||||
|
|
||||
|
static EthashAux* get() { if (!s_this) s_this = new EthashAux(); return s_this; } |
||||
|
|
||||
|
using LightType = void const*; |
||||
|
using FullType = void const*; |
||||
|
|
||||
|
static h256 seedHash(unsigned _number); |
||||
|
static ethash_params params(BlockInfo const& _header); |
||||
|
static ethash_params params(h256 const& _seedHash); |
||||
|
static ethash_params params(unsigned _n); |
||||
|
static LightType light(BlockInfo const& _header); |
||||
|
static LightType light(h256 const& _seedHash); |
||||
|
static bytesConstRef full(BlockInfo const& _header, bytesRef _dest = bytesRef()); |
||||
|
static bytesConstRef full(h256 const& _header, bytesRef _dest = bytesRef()); |
||||
|
|
||||
|
static Ethash::Result eval(BlockInfo const& _header) { return eval(_header, _header.nonce); } |
||||
|
static Ethash::Result eval(BlockInfo const& _header, Nonce const& _nonce); |
||||
|
static Ethash::Result eval(h256 const& _seedHash, h256 const& _headerHash, Nonce const& _nonce); |
||||
|
|
||||
|
private: |
||||
|
EthashAux() {} |
||||
|
|
||||
|
void killCache(h256 const& _s); |
||||
|
|
||||
|
static EthashAux* s_this; |
||||
|
RecursiveMutex x_this; |
||||
|
|
||||
|
std::map<h256, LightType> m_lights; |
||||
|
std::map<h256, bytesRef> m_fulls; |
||||
|
std::map<h256, unsigned> m_epochs; |
||||
|
h256s m_seedHashes; |
||||
|
}; |
||||
|
|
||||
|
} |
||||
|
} |
||||
@ -1,220 +0,0 @@ |
|||||
/*
|
|
||||
This file is part of cpp-ethereum. |
|
||||
|
|
||||
cpp-ethereum 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. |
|
||||
|
|
||||
cpp-ethereum 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 Ethasher.cpp
|
|
||||
* @author Gav Wood <i@gavwood.com> |
|
||||
* @date 2014 |
|
||||
*/ |
|
||||
|
|
||||
#include <boost/detail/endian.hpp> |
|
||||
#include <boost/filesystem.hpp> |
|
||||
#include <chrono> |
|
||||
#include <array> |
|
||||
#include <random> |
|
||||
#include <thread> |
|
||||
#include <libdevcore/Common.h> |
|
||||
#include <libdevcore/Guards.h> |
|
||||
#include <libdevcore/Log.h> |
|
||||
#include <libdevcrypto/CryptoPP.h> |
|
||||
#include <libdevcrypto/SHA3.h> |
|
||||
#include <libdevcrypto/FileSystem.h> |
|
||||
#include <libethcore/Params.h> |
|
||||
#include "BlockInfo.h" |
|
||||
#include "Ethasher.h" |
|
||||
using namespace std; |
|
||||
using namespace chrono; |
|
||||
using namespace dev; |
|
||||
using namespace eth; |
|
||||
|
|
||||
Ethasher* dev::eth::Ethasher::s_this = nullptr; |
|
||||
|
|
||||
Ethasher::~Ethasher() |
|
||||
{ |
|
||||
while (!m_lights.empty()) |
|
||||
killCache(m_lights.begin()->first); |
|
||||
} |
|
||||
|
|
||||
void Ethasher::killCache(h256 const& _s) |
|
||||
{ |
|
||||
RecursiveGuard l(x_this); |
|
||||
if (m_lights.count(_s)) |
|
||||
{ |
|
||||
ethash_delete_light(m_lights.at(_s)); |
|
||||
m_lights.erase(_s); |
|
||||
} |
|
||||
} |
|
||||
|
|
||||
void const* Ethasher::light(BlockInfo const& _header) |
|
||||
{ |
|
||||
RecursiveGuard l(x_this); |
|
||||
if (_header.number > c_ethashEpochLength * 2048) |
|
||||
{ |
|
||||
std::ostringstream error; |
|
||||
error << "block number is too high; max is " << c_ethashEpochLength * 2048 << "(was " << _header.number << ")"; |
|
||||
throw std::invalid_argument( error.str() ); |
|
||||
} |
|
||||
|
|
||||
if (!m_lights.count(_header.seedHash())) |
|
||||
{ |
|
||||
ethash_params p = params((unsigned)_header.number); |
|
||||
m_lights[_header.seedHash()] = ethash_new_light(&p, _header.seedHash().data()); |
|
||||
} |
|
||||
return m_lights[_header.seedHash()]; |
|
||||
} |
|
||||
|
|
||||
#define IGNORE_EXCEPTIONS(X) try { X; } catch (...) {} |
|
||||
|
|
||||
bytesConstRef Ethasher::full(BlockInfo const& _header) |
|
||||
{ |
|
||||
RecursiveGuard l(x_this); |
|
||||
if (!m_fulls.count(_header.seedHash())) |
|
||||
{ |
|
||||
// @memoryleak @bug place it on a pile for deletion - perhaps use shared_ptr.
|
|
||||
/* if (!m_fulls.empty())
|
|
||||
{ |
|
||||
delete [] m_fulls.begin()->second.data(); |
|
||||
m_fulls.erase(m_fulls.begin()); |
|
||||
}*/ |
|
||||
|
|
||||
try { |
|
||||
boost::filesystem::create_directories(getDataDir("ethash")); |
|
||||
} catch (...) {} |
|
||||
|
|
||||
auto info = rlpList(c_ethashRevision, _header.seedHash()); |
|
||||
std::string oldMemoFile = getDataDir("ethash") + "/full"; |
|
||||
std::string memoFile = getDataDir("ethash") + "/full-R" + toString(c_ethashRevision) + "-" + toHex(_header.seedHash().ref().cropped(0, 8)); |
|
||||
if (boost::filesystem::exists(oldMemoFile) && contents(oldMemoFile + ".info") == info) |
|
||||
{ |
|
||||
// memofile valid - rename.
|
|
||||
boost::filesystem::rename(oldMemoFile, memoFile); |
|
||||
} |
|
||||
|
|
||||
IGNORE_EXCEPTIONS(boost::filesystem::remove(oldMemoFile)); |
|
||||
IGNORE_EXCEPTIONS(boost::filesystem::remove(oldMemoFile + ".info")); |
|
||||
|
|
||||
m_fulls[_header.seedHash()] = contentsNew(memoFile); |
|
||||
if (!m_fulls[_header.seedHash()]) |
|
||||
{ |
|
||||
ethash_params p = params((unsigned)_header.number); |
|
||||
m_fulls[_header.seedHash()] = bytesRef(new byte[p.full_size], p.full_size); |
|
||||
auto c = light(_header); |
|
||||
ethash_prep_full(m_fulls[_header.seedHash()].data(), &p, c); |
|
||||
writeFile(memoFile, m_fulls[_header.seedHash()]); |
|
||||
} |
|
||||
} |
|
||||
return m_fulls[_header.seedHash()]; |
|
||||
} |
|
||||
|
|
||||
ethash_params Ethasher::params(BlockInfo const& _header) |
|
||||
{ |
|
||||
return params((unsigned)_header.number); |
|
||||
} |
|
||||
|
|
||||
void Ethasher::readFull(BlockInfo const& _header, void* _dest) |
|
||||
{ |
|
||||
if (!m_fulls.count(_header.seedHash())) |
|
||||
{ |
|
||||
// @memoryleak @bug place it on a pile for deletion - perhaps use shared_ptr.
|
|
||||
/* if (!m_fulls.empty())
|
|
||||
{ |
|
||||
delete [] m_fulls.begin()->second.data(); |
|
||||
m_fulls.erase(m_fulls.begin()); |
|
||||
}*/ |
|
||||
|
|
||||
try { |
|
||||
boost::filesystem::create_directories(getDataDir("ethash")); |
|
||||
} catch (...) {} |
|
||||
|
|
||||
auto info = rlpList(c_ethashRevision, _header.seedHash()); |
|
||||
std::string oldMemoFile = getDataDir("ethash") + "/full"; |
|
||||
std::string memoFile = getDataDir("ethash") + "/full-R" + toString(c_ethashRevision) + "-" + toHex(_header.seedHash().ref().cropped(0, 8)); |
|
||||
if (boost::filesystem::exists(oldMemoFile) && contents(oldMemoFile + ".info") == info) |
|
||||
{ |
|
||||
// memofile valid - rename.
|
|
||||
boost::filesystem::rename(oldMemoFile, memoFile); |
|
||||
} |
|
||||
|
|
||||
IGNORE_EXCEPTIONS(boost::filesystem::remove(oldMemoFile)); |
|
||||
IGNORE_EXCEPTIONS(boost::filesystem::remove(oldMemoFile + ".info")); |
|
||||
|
|
||||
ethash_params p = params((unsigned)_header.number); |
|
||||
bytesRef r = contentsNew(memoFile, bytesRef((byte*)_dest, p.full_size)); |
|
||||
if (!r) |
|
||||
{ |
|
||||
auto c = light(_header); |
|
||||
ethash_prep_full(_dest, &p, c); |
|
||||
writeFile(memoFile, bytesConstRef((byte*)_dest, p.full_size)); |
|
||||
} |
|
||||
} |
|
||||
} |
|
||||
|
|
||||
ethash_params Ethasher::params(unsigned _n) |
|
||||
{ |
|
||||
ethash_params p; |
|
||||
p.cache_size = ethash_get_cachesize(_n); |
|
||||
p.full_size = ethash_get_datasize(_n); |
|
||||
return p; |
|
||||
} |
|
||||
|
|
||||
bool Ethasher::verify(BlockInfo const& _header) |
|
||||
{ |
|
||||
if (_header.number >= c_ethashEpochLength * 2048) |
|
||||
return false; |
|
||||
|
|
||||
h256 boundary = u256((bigint(1) << 256) / _header.difficulty); |
|
||||
|
|
||||
bool quick = ethash_quick_check_difficulty( |
|
||||
_header.headerHash(WithoutNonce).data(), |
|
||||
(uint64_t)(u64)_header.nonce, |
|
||||
_header.mixHash.data(), |
|
||||
boundary.data()); |
|
||||
|
|
||||
#if !ETH_DEBUG |
|
||||
if (!quick) |
|
||||
return false; |
|
||||
#endif |
|
||||
|
|
||||
auto result = eval(_header); |
|
||||
bool slow = result.value <= boundary && result.mixHash == _header.mixHash; |
|
||||
|
|
||||
#if ETH_DEBUG |
|
||||
if (!quick && slow) |
|
||||
{ |
|
||||
cwarn << "WARNING: evaluated result gives true whereas ethash_quick_check_difficulty gives false."; |
|
||||
cwarn << "headerHash:" << _header.headerHash(WithoutNonce); |
|
||||
cwarn << "nonce:" << _header.nonce; |
|
||||
cwarn << "mixHash:" << _header.mixHash; |
|
||||
cwarn << "difficulty:" << _header.difficulty; |
|
||||
cwarn << "boundary:" << boundary; |
|
||||
cwarn << "result.value:" << result.value; |
|
||||
cwarn << "result.mixHash:" << result.mixHash; |
|
||||
} |
|
||||
#endif |
|
||||
|
|
||||
return slow; |
|
||||
} |
|
||||
|
|
||||
Ethasher::Result Ethasher::eval(BlockInfo const& _header, Nonce const& _nonce) |
|
||||
{ |
|
||||
auto p = Ethasher::params(_header); |
|
||||
ethash_return_value r; |
|
||||
if (Ethasher::get()->m_fulls.count(_header.seedHash())) |
|
||||
ethash_compute_full(&r, Ethasher::get()->full(_header).data(), &p, _header.headerHash(WithoutNonce).data(), (uint64_t)(u64)_nonce); |
|
||||
else |
|
||||
ethash_compute_light(&r, Ethasher::get()->light(_header), &p, _header.headerHash(WithoutNonce).data(), (uint64_t)(u64)_nonce); |
|
||||
// cdebug << "Ethasher::eval sha3(cache):" << sha3(Ethasher::get()->cache(_header)) << "hh:" << _header.headerHash(WithoutNonce) << "nonce:" << _nonce << " => " << h256(r.result, h256::ConstructFromPointer);
|
|
||||
return Result{h256(r.result, h256::ConstructFromPointer), h256(r.mix_hash, h256::ConstructFromPointer)}; |
|
||||
} |
|
||||
@ -1,109 +0,0 @@ |
|||||
/*
|
|
||||
This file is part of cpp-ethereum. |
|
||||
|
|
||||
cpp-ethereum 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. |
|
||||
|
|
||||
cpp-ethereum 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 Ethasher.h
|
|
||||
* @author Gav Wood <i@gavwood.com> |
|
||||
* @date 2014 |
|
||||
* |
|
||||
* ProofOfWork algorithm. |
|
||||
*/ |
|
||||
|
|
||||
#pragma once |
|
||||
|
|
||||
#include <chrono> |
|
||||
#include <thread> |
|
||||
#include <cstdint> |
|
||||
#include <libdevcore/Guards.h> |
|
||||
#include <libdevcore/Log.h> |
|
||||
#include <libdevcrypto/SHA3.h> |
|
||||
#include <libethash/ethash.h> // TODO: REMOVE once everything merged into this class and an opaque API can be provided. |
|
||||
static const unsigned c_ethashRevision = ETHASH_REVISION; |
|
||||
static const unsigned c_ethashEpochLength = ETHASH_EPOCH_LENGTH; |
|
||||
#include "Common.h" |
|
||||
#include "BlockInfo.h" |
|
||||
|
|
||||
namespace dev |
|
||||
{ |
|
||||
namespace eth |
|
||||
{ |
|
||||
|
|
||||
class Ethasher |
|
||||
{ |
|
||||
public: |
|
||||
Ethasher() {} |
|
||||
~Ethasher(); |
|
||||
|
|
||||
static Ethasher* get() { if (!s_this) s_this = new Ethasher(); return s_this; } |
|
||||
|
|
||||
using LightType = void const*; |
|
||||
using FullType = void const*; |
|
||||
|
|
||||
LightType light(BlockInfo const& _header); |
|
||||
bytesConstRef full(BlockInfo const& _header); |
|
||||
static ethash_params params(BlockInfo const& _header); |
|
||||
static ethash_params params(unsigned _n); |
|
||||
|
|
||||
void readFull(BlockInfo const& _header, void* _dest); |
|
||||
|
|
||||
struct Result |
|
||||
{ |
|
||||
h256 value; |
|
||||
h256 mixHash; |
|
||||
}; |
|
||||
|
|
||||
static Result eval(BlockInfo const& _header) { return eval(_header, _header.nonce); } |
|
||||
static Result eval(BlockInfo const& _header, Nonce const& _nonce); |
|
||||
static bool verify(BlockInfo const& _header); |
|
||||
|
|
||||
class Miner |
|
||||
{ |
|
||||
public: |
|
||||
Miner(BlockInfo const& _header): |
|
||||
m_headerHash(_header.headerHash(WithoutNonce)), |
|
||||
m_params(Ethasher::params(_header)), |
|
||||
m_datasetPointer(Ethasher::get()->full(_header).data()) |
|
||||
{} |
|
||||
|
|
||||
inline h256 mine(uint64_t _nonce) |
|
||||
{ |
|
||||
ethash_compute_full(&m_ethashReturn, m_datasetPointer, &m_params, m_headerHash.data(), _nonce); |
|
||||
// cdebug << "Ethasher::mine hh:" << m_headerHash << "nonce:" << (Nonce)(u64)_nonce << " => " << h256(m_ethashReturn.result, h256::ConstructFromPointer);
|
|
||||
return h256(m_ethashReturn.result, h256::ConstructFromPointer); |
|
||||
} |
|
||||
|
|
||||
inline h256 lastMixHash() const |
|
||||
{ |
|
||||
return h256(m_ethashReturn.mix_hash, h256::ConstructFromPointer); |
|
||||
} |
|
||||
|
|
||||
private: |
|
||||
ethash_return_value m_ethashReturn; |
|
||||
h256 m_headerHash; |
|
||||
ethash_params m_params; |
|
||||
void const* m_datasetPointer; |
|
||||
}; |
|
||||
|
|
||||
private: |
|
||||
void killCache(h256 const& _s); |
|
||||
|
|
||||
static Ethasher* s_this; |
|
||||
RecursiveMutex x_this; |
|
||||
std::map<h256, LightType> m_lights; |
|
||||
std::map<h256, bytesRef> m_fulls; |
|
||||
}; |
|
||||
|
|
||||
} |
|
||||
} |
|
||||
@ -0,0 +1,172 @@ |
|||||
|
/*
|
||||
|
This file is part of cpp-ethereum. |
||||
|
|
||||
|
cpp-ethereum 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. |
||||
|
|
||||
|
cpp-ethereum 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 Miner.h
|
||||
|
* @author Gav Wood <i@gavwood.com> |
||||
|
* @date 2015 |
||||
|
*/ |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <thread> |
||||
|
#include <list> |
||||
|
#include <atomic> |
||||
|
#include <libdevcore/Common.h> |
||||
|
#include <libdevcore/Worker.h> |
||||
|
#include <libethcore/Common.h> |
||||
|
|
||||
|
namespace dev |
||||
|
{ |
||||
|
|
||||
|
namespace eth |
||||
|
{ |
||||
|
|
||||
|
/**
|
||||
|
* @brief Describes the progress of a mining operation. |
||||
|
*/ |
||||
|
struct MiningProgress |
||||
|
{ |
||||
|
// MiningProgress& operator+=(MiningProgress const& _mp) { hashes += _mp.hashes; ms = std::max(ms, _mp.ms); return *this; }
|
||||
|
uint64_t hashes = 0; ///< Total number of hashes computed.
|
||||
|
uint64_t ms = 0; ///< Total number of milliseconds of mining thus far.
|
||||
|
uint64_t rate() const { return hashes * 1000 / ms; } |
||||
|
}; |
||||
|
|
||||
|
struct MineInfo: public MiningProgress {}; |
||||
|
|
||||
|
inline std::ostream& operator<<(std::ostream& _out, MiningProgress _p) |
||||
|
{ |
||||
|
_out << _p.rate() << " H/s = " << _p.hashes << " hashes / " << (double(_p.ms) / 1000) << " s"; |
||||
|
return _out; |
||||
|
} |
||||
|
|
||||
|
template <class PoW> class GenericMiner; |
||||
|
|
||||
|
/**
|
||||
|
* @brief Class for hosting one or more Miners. |
||||
|
* @warning Must be implemented in a threadsafe manner since it will be called from multiple |
||||
|
* miner threads. |
||||
|
*/ |
||||
|
template <class PoW> class GenericFarmFace |
||||
|
{ |
||||
|
public: |
||||
|
using WorkPackage = typename PoW::WorkPackage; |
||||
|
using Solution = typename PoW::Solution; |
||||
|
using Miner = GenericMiner<PoW>; |
||||
|
|
||||
|
/**
|
||||
|
* @brief Called from a Miner to note a WorkPackage has a solution. |
||||
|
* @param _p The solution. |
||||
|
* @param _wp The WorkPackage that the Solution is for; this will be reset if the work is accepted. |
||||
|
* @param _finder The miner that found it. |
||||
|
* @return true iff the solution was good (implying that mining should be . |
||||
|
*/ |
||||
|
virtual bool submitProof(Solution const& _p, Miner* _finder) = 0; |
||||
|
}; |
||||
|
|
||||
|
/**
|
||||
|
* @brief A miner - a member and adoptee of the Farm. |
||||
|
* @warning Not threadsafe. It is assumed Farm will synchronise calls to/from this class. |
||||
|
*/ |
||||
|
template <class PoW> class GenericMiner |
||||
|
{ |
||||
|
public: |
||||
|
using WorkPackage = typename PoW::WorkPackage; |
||||
|
using Solution = typename PoW::Solution; |
||||
|
using FarmFace = GenericFarmFace<PoW>; |
||||
|
using ConstructionInfo = std::pair<FarmFace*, unsigned>; |
||||
|
|
||||
|
GenericMiner(ConstructionInfo const& _ci): |
||||
|
m_farm(_ci.first), |
||||
|
m_index(_ci.second) |
||||
|
{} |
||||
|
|
||||
|
// API FOR THE FARM TO CALL IN WITH
|
||||
|
|
||||
|
void setWork(WorkPackage const& _work = WorkPackage()) |
||||
|
{ |
||||
|
auto old = m_work; |
||||
|
{ |
||||
|
Guard l(x_work); |
||||
|
m_work = _work; |
||||
|
} |
||||
|
if (!!_work) |
||||
|
{ |
||||
|
pause(); |
||||
|
kickOff(); |
||||
|
} |
||||
|
else if (!_work && !!old) |
||||
|
pause(); |
||||
|
m_hashCount = 0; |
||||
|
} |
||||
|
|
||||
|
uint64_t hashCount() const { return m_hashCount; } |
||||
|
|
||||
|
void resetHashCount() { m_hashCount = 0; } |
||||
|
|
||||
|
unsigned index() const { return m_index; } |
||||
|
|
||||
|
protected: |
||||
|
|
||||
|
// REQUIRED TO BE REIMPLEMENTED BY A SUBCLASS:
|
||||
|
|
||||
|
/**
|
||||
|
* @brief Begin working on a given work package, discarding any previous work. |
||||
|
* @param _work The package for which to find a solution. |
||||
|
*/ |
||||
|
virtual void kickOff() = 0; |
||||
|
|
||||
|
/**
|
||||
|
* @brief No work left to be done. Pause until told to kickOff(). |
||||
|
*/ |
||||
|
virtual void pause() = 0; |
||||
|
|
||||
|
// AVAILABLE FOR A SUBCLASS TO CALL:
|
||||
|
|
||||
|
/**
|
||||
|
* @brief Notes that the Miner found a solution. |
||||
|
* @param _s The solution. |
||||
|
* @return true if the solution was correct and that the miner should pause. |
||||
|
*/ |
||||
|
bool submitProof(Solution const& _s) |
||||
|
{ |
||||
|
if (!m_farm) |
||||
|
return true; |
||||
|
if (m_farm->submitProof(_s, this)) |
||||
|
{ |
||||
|
Guard l(x_work); |
||||
|
m_work.reset(); |
||||
|
return true; |
||||
|
} |
||||
|
return false; |
||||
|
} |
||||
|
|
||||
|
WorkPackage const& work() const { Guard l(x_work); return m_work; } |
||||
|
|
||||
|
void accumulateHashes(unsigned _n) { m_hashCount += _n; } |
||||
|
|
||||
|
private: |
||||
|
FarmFace* m_farm = nullptr; |
||||
|
unsigned m_index; |
||||
|
|
||||
|
uint64_t m_hashCount = 0; |
||||
|
|
||||
|
WorkPackage m_work; |
||||
|
mutable Mutex x_work; |
||||
|
}; |
||||
|
|
||||
|
} |
||||
|
} |
||||
@ -0,0 +1,205 @@ |
|||||
|
/*
|
||||
|
This file is part of cpp-ethereum. |
||||
|
|
||||
|
cpp-ethereum 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. |
||||
|
|
||||
|
cpp-ethereum 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 Farm.h
|
||||
|
* @author Gav Wood <i@gavwood.com> |
||||
|
* @date 2015 |
||||
|
*/ |
||||
|
|
||||
|
#pragma once |
||||
|
|
||||
|
#include <thread> |
||||
|
#include <list> |
||||
|
#include <atomic> |
||||
|
#include <libdevcore/Common.h> |
||||
|
#include <libdevcore/Worker.h> |
||||
|
#include <libethcore/Common.h> |
||||
|
#include <libethcore/Miner.h> |
||||
|
#include <libethcore/BlockInfo.h> |
||||
|
#include <libethcore/ProofOfWork.h> |
||||
|
|
||||
|
namespace dev |
||||
|
{ |
||||
|
|
||||
|
namespace eth |
||||
|
{ |
||||
|
|
||||
|
/**
|
||||
|
* @brief A collective of Miners. |
||||
|
* Miners ask for work, then submit proofs |
||||
|
* @threadsafe |
||||
|
*/ |
||||
|
template <class PoW> |
||||
|
class GenericFarm: public GenericFarmFace<PoW> |
||||
|
{ |
||||
|
public: |
||||
|
using WorkPackage = typename PoW::WorkPackage; |
||||
|
using Solution = typename PoW::Solution; |
||||
|
using Miner = GenericMiner<PoW>; |
||||
|
|
||||
|
~GenericFarm() |
||||
|
{ |
||||
|
stop(); |
||||
|
} |
||||
|
|
||||
|
/**
|
||||
|
* @brief Sets the current mining mission. |
||||
|
* @param _bi The block (header) we wish to be mining. |
||||
|
*/ |
||||
|
void setWork(BlockInfo const& _bi) |
||||
|
{ |
||||
|
WriteGuard l(x_minerWork); |
||||
|
m_header = _bi; |
||||
|
auto p = PoW::package(m_header); |
||||
|
if (p.headerHash == m_work.headerHash) |
||||
|
return; |
||||
|
m_work = p; |
||||
|
for (auto const& m: m_miners) |
||||
|
m->setWork(m_work); |
||||
|
resetTimer(); |
||||
|
} |
||||
|
|
||||
|
/**
|
||||
|
* @brief (Re)start miners for CPU only. |
||||
|
* @returns true if started properly. |
||||
|
*/ |
||||
|
bool startCPU() { return start<typename PoW::CPUMiner>(); } |
||||
|
|
||||
|
/**
|
||||
|
* @brief (Re)start miners for GPU only. |
||||
|
* @returns true if started properly. |
||||
|
*/ |
||||
|
bool startGPU() { return start<typename PoW::GPUMiner>(); } |
||||
|
|
||||
|
/**
|
||||
|
* @brief Stop all mining activities. |
||||
|
*/ |
||||
|
void stop() |
||||
|
{ |
||||
|
WriteGuard l(x_minerWork); |
||||
|
m_miners.clear(); |
||||
|
m_work.reset(); |
||||
|
m_isMining = false; |
||||
|
} |
||||
|
|
||||
|
bool isMining() const |
||||
|
{ |
||||
|
return m_isMining; |
||||
|
} |
||||
|
|
||||
|
/**
|
||||
|
* @brief Get information on the progress of mining this work package. |
||||
|
* @return The progress with mining so far. |
||||
|
*/ |
||||
|
MiningProgress const& miningProgress() const |
||||
|
{ |
||||
|
MiningProgress p; |
||||
|
p.ms = std::chrono::duration_cast<std::chrono::milliseconds>(std::chrono::steady_clock::now() - m_lastStart).count(); |
||||
|
{ |
||||
|
ReadGuard l2(x_minerWork); |
||||
|
for (auto const& i: m_miners) |
||||
|
p.hashes += i->hashCount(); |
||||
|
} |
||||
|
ReadGuard l(x_progress); |
||||
|
m_progress = p; |
||||
|
return m_progress; |
||||
|
} |
||||
|
|
||||
|
/**
|
||||
|
* @brief Reset the mining progess counter. |
||||
|
*/ |
||||
|
void resetMiningProgress() |
||||
|
{ |
||||
|
ETH_READ_GUARDED(x_minerWork) |
||||
|
for (auto const& i: m_miners) |
||||
|
i->resetHashCount(); |
||||
|
resetTimer(); |
||||
|
} |
||||
|
|
||||
|
using SolutionFound = std::function<bool(Solution const&)>; |
||||
|
|
||||
|
/**
|
||||
|
* @brief Provides a valid header based upon that received previously with setWork(). |
||||
|
* @param _bi The now-valid header. |
||||
|
* @return true if the header was good and that the Farm should pause until more work is submitted. |
||||
|
*/ |
||||
|
void onSolutionFound(SolutionFound const& _handler) { m_onSolutionFound = _handler; } |
||||
|
|
||||
|
WorkPackage work() const { ReadGuard l(x_minerWork); return m_work; } |
||||
|
|
||||
|
private: |
||||
|
/**
|
||||
|
* @brief Called from a Miner to note a WorkPackage has a solution. |
||||
|
* @param _p The solution. |
||||
|
* @param _wp The WorkPackage that the Solution is for. |
||||
|
* @return true iff the solution was good (implying that mining should be . |
||||
|
*/ |
||||
|
bool submitProof(Solution const& _s, Miner* _m) override |
||||
|
{ |
||||
|
if (m_onSolutionFound && m_onSolutionFound(_s)) |
||||
|
{ |
||||
|
WriteGuard ul(x_minerWork); |
||||
|
for (std::shared_ptr<Miner> const& m: m_miners) |
||||
|
if (m.get() != _m) |
||||
|
m->setWork(); |
||||
|
m_work.reset(); |
||||
|
return true; |
||||
|
} |
||||
|
return false; |
||||
|
} |
||||
|
|
||||
|
/**
|
||||
|
* @brief Start a number of miners. |
||||
|
*/ |
||||
|
template <class MinerType> |
||||
|
bool start() |
||||
|
{ |
||||
|
WriteGuard l(x_minerWork); |
||||
|
if (!m_miners.empty() && !!std::dynamic_pointer_cast<MinerType>(m_miners[0])) |
||||
|
return true; |
||||
|
m_miners.clear(); |
||||
|
m_miners.reserve(MinerType::instances()); |
||||
|
for (unsigned i = 0; i < MinerType::instances(); ++i) |
||||
|
{ |
||||
|
m_miners.push_back(std::shared_ptr<Miner>(new MinerType(std::make_pair(this, i)))); |
||||
|
m_miners.back()->setWork(m_work); |
||||
|
} |
||||
|
m_isMining = true; |
||||
|
resetTimer(); |
||||
|
return true; |
||||
|
} |
||||
|
|
||||
|
void resetTimer() |
||||
|
{ |
||||
|
m_lastStart = std::chrono::steady_clock::now(); |
||||
|
} |
||||
|
|
||||
|
mutable SharedMutex x_minerWork; |
||||
|
std::vector<std::shared_ptr<Miner>> m_miners; |
||||
|
WorkPackage m_work; |
||||
|
BlockInfo m_header; |
||||
|
|
||||
|
std::atomic<bool> m_isMining = {false}; |
||||
|
|
||||
|
mutable SharedMutex x_progress; |
||||
|
mutable MiningProgress m_progress; |
||||
|
std::chrono::steady_clock::time_point m_lastStart; |
||||
|
|
||||
|
SolutionFound m_onSolutionFound; |
||||
|
}; |
||||
|
|
||||
|
} |
||||
|
} |
||||
@ -1,96 +0,0 @@ |
|||||
/*
|
|
||||
This file is part of cpp-ethereum. |
|
||||
|
|
||||
cpp-ethereum 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. |
|
||||
|
|
||||
cpp-ethereum 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 Miner.cpp
|
|
||||
* @author Gav Wood <i@gavwood.com> |
|
||||
* @author Giacomo Tazzari |
|
||||
* @date 2014 |
|
||||
*/ |
|
||||
|
|
||||
#include "Miner.h" |
|
||||
|
|
||||
#include <libdevcore/CommonIO.h> |
|
||||
#include "State.h" |
|
||||
using namespace std; |
|
||||
using namespace dev; |
|
||||
using namespace dev::eth; |
|
||||
|
|
||||
Miner::~Miner() {} |
|
||||
|
|
||||
LocalMiner::LocalMiner(MinerHost* _host, unsigned _id): |
|
||||
AsyncMiner(_host, _id), |
|
||||
Worker("miner-" + toString(_id)) |
|
||||
{ |
|
||||
m_pow.reset(_host->turbo() ? new Ethash : (Ethash*)new EthashCPU); |
|
||||
} |
|
||||
|
|
||||
void LocalMiner::setup(MinerHost* _host, unsigned _id) |
|
||||
{ |
|
||||
AsyncMiner::setup(_host, _id); |
|
||||
setName("miner-" + toString(m_id)); |
|
||||
m_pow.reset(_host->turbo() ? new Ethash : (Ethash*)new EthashCPU); |
|
||||
} |
|
||||
|
|
||||
void LocalMiner::doWork() |
|
||||
{ |
|
||||
// Do some mining.
|
|
||||
if (m_miningStatus != Waiting && m_miningStatus != Mined) |
|
||||
{ |
|
||||
if (m_miningStatus == Preparing) |
|
||||
{ |
|
||||
m_host->setupState(m_mineState); |
|
||||
if (m_host->force() || m_mineState.pending().size()) |
|
||||
m_miningStatus = Mining; |
|
||||
else |
|
||||
m_miningStatus = Waiting; |
|
||||
|
|
||||
{ |
|
||||
Guard l(x_mineInfo); |
|
||||
m_mineProgress.best = (double)-1; |
|
||||
m_mineProgress.hashes = 0; |
|
||||
m_mineProgress.ms = 0; |
|
||||
} |
|
||||
} |
|
||||
|
|
||||
if (m_miningStatus == Mining) |
|
||||
{ |
|
||||
// Mine for a while.
|
|
||||
MineInfo mineInfo = m_mineState.mine(m_pow.get()); |
|
||||
|
|
||||
{ |
|
||||
Guard l(x_mineInfo); |
|
||||
m_mineProgress.best = min(m_mineProgress.best, mineInfo.best); |
|
||||
m_mineProgress.current = mineInfo.best; |
|
||||
m_mineProgress.requirement = mineInfo.requirement; |
|
||||
m_mineProgress.ms += 100; |
|
||||
m_mineProgress.hashes += mineInfo.hashes; |
|
||||
m_mineHistory.push_back(mineInfo); |
|
||||
} |
|
||||
if (mineInfo.completed) |
|
||||
{ |
|
||||
m_mineState.completeMine(); |
|
||||
m_host->onComplete(); |
|
||||
m_miningStatus = Mined; |
|
||||
} |
|
||||
else |
|
||||
m_host->onProgressed(); |
|
||||
} |
|
||||
} |
|
||||
else |
|
||||
{ |
|
||||
this_thread::sleep_for(chrono::milliseconds(100)); |
|
||||
} |
|
||||
} |
|
||||
@ -1,178 +0,0 @@ |
|||||
/*
|
|
||||
This file is part of cpp-ethereum. |
|
||||
|
|
||||
cpp-ethereum 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. |
|
||||
|
|
||||
cpp-ethereum 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 Miner.h
|
|
||||
* @author Alex Leverington <nessence@gmail.com> |
|
||||
* @author Gav Wood <i@gavwood.com> |
|
||||
* @date 2014 |
|
||||
*/ |
|
||||
|
|
||||
#pragma once |
|
||||
|
|
||||
#include <thread> |
|
||||
#include <list> |
|
||||
#include <atomic> |
|
||||
#include <libdevcore/Common.h> |
|
||||
#include <libdevcore/Worker.h> |
|
||||
#include <libethcore/Common.h> |
|
||||
#include "State.h" |
|
||||
|
|
||||
namespace dev |
|
||||
{ |
|
||||
|
|
||||
namespace eth |
|
||||
{ |
|
||||
|
|
||||
/**
|
|
||||
* @brief Describes the progress of a mining operation. |
|
||||
*/ |
|
||||
struct MineProgress |
|
||||
{ |
|
||||
void combine(MineProgress const& _m) { requirement = std::max(requirement, _m.requirement); best = std::min(best, _m.best); current = std::max(current, _m.current); hashes += _m.hashes; ms = std::max(ms, _m.ms); } |
|
||||
double requirement = 0; ///< The PoW requirement - as the second logarithm of the minimum acceptable hash.
|
|
||||
double best = 1e99; ///< The PoW achievement - as the second logarithm of the minimum found hash.
|
|
||||
double current = 0; ///< The most recent PoW achievement - as the second logarithm of the presently found hash.
|
|
||||
unsigned hashes = 0; ///< Total number of hashes computed.
|
|
||||
unsigned ms = 0; ///< Total number of milliseconds of mining thus far.
|
|
||||
}; |
|
||||
|
|
||||
/**
|
|
||||
* @brief Class for hosting one or more Miners. |
|
||||
* @warning Must be implemented in a threadsafe manner since it will be called from multiple |
|
||||
* miner threads. |
|
||||
*/ |
|
||||
class MinerHost |
|
||||
{ |
|
||||
public: |
|
||||
virtual void setupState(State& _s) = 0; ///< Reset the given State object to the one that should be being mined.
|
|
||||
virtual void onProgressed() {} ///< Called once some progress has been made.
|
|
||||
virtual void onComplete() {} ///< Called once a block is found.
|
|
||||
virtual bool force() const = 0; ///< @returns true iff the Miner should mine regardless of the number of transactions.
|
|
||||
virtual bool turbo() const = 0; ///< @returns true iff the Miner should use GPU if possible.
|
|
||||
}; |
|
||||
|
|
||||
class Miner |
|
||||
{ |
|
||||
public: |
|
||||
virtual ~Miner(); |
|
||||
|
|
||||
virtual void noteStateChange() = 0; |
|
||||
virtual bool isComplete() const = 0; |
|
||||
virtual bytes const& blockData() const = 0; |
|
||||
}; |
|
||||
|
|
||||
class AsyncMiner: public Miner |
|
||||
{ |
|
||||
public: |
|
||||
/// Null constructor.
|
|
||||
AsyncMiner(): m_host(nullptr) {} |
|
||||
|
|
||||
/// Constructor.
|
|
||||
AsyncMiner(MinerHost* _host, unsigned _id = 0): m_host(_host), m_id(_id) {} |
|
||||
|
|
||||
/// Setup its basics.
|
|
||||
void setup(MinerHost* _host, unsigned _id = 0) { m_host = _host; m_id = _id; } |
|
||||
|
|
||||
/// Start mining.
|
|
||||
virtual void start() {} |
|
||||
|
|
||||
/// Stop mining.
|
|
||||
virtual void stop() {} |
|
||||
|
|
||||
/// @returns true iff the mining has been start()ed. It may still not be actually mining, depending on the host's turbo() & force().
|
|
||||
virtual bool isRunning() const { return false; } |
|
||||
|
|
||||
protected: |
|
||||
MinerHost* m_host = nullptr; ///< Our host.
|
|
||||
unsigned m_id = 0; ///< Our unique id.
|
|
||||
}; |
|
||||
|
|
||||
/**
|
|
||||
* @brief Implements Miner. |
|
||||
* To begin mining, use start() & stop(). noteStateChange() can be used to reset the mining and set up the |
|
||||
* State object according to the host. Use isRunning() to determine if the miner has been start()ed. |
|
||||
* Use isComplete() to determine if the miner has finished mining. |
|
||||
* |
|
||||
* blockData() can be used to retrieve the complete block, ready for insertion into the BlockChain. |
|
||||
* |
|
||||
* Information on the mining can be queried through miningProgress() and miningHistory(). |
|
||||
* @threadsafe |
|
||||
* @todo Signal Miner to restart once with condition variables. |
|
||||
*/ |
|
||||
class LocalMiner: public AsyncMiner, Worker |
|
||||
{ |
|
||||
public: |
|
||||
/// Null constructor.
|
|
||||
LocalMiner() {} |
|
||||
|
|
||||
/// Constructor.
|
|
||||
LocalMiner(MinerHost* _host, unsigned _id = 0); |
|
||||
|
|
||||
/// Move-constructor.
|
|
||||
LocalMiner(LocalMiner&& _m): Worker((Worker&&)_m) { std::swap(m_host, _m.m_host); std::swap(m_pow, _m.m_pow); } |
|
||||
|
|
||||
/// Move-assignment.
|
|
||||
LocalMiner& operator=(LocalMiner&& _m) { Worker::operator=((Worker&&)_m); std::swap(m_host, _m.m_host); std::swap(m_pow, _m.m_pow); return *this; } |
|
||||
|
|
||||
/// Destructor. Stops miner.
|
|
||||
~LocalMiner() { stop(); } |
|
||||
|
|
||||
/// Setup its basics.
|
|
||||
void setup(MinerHost* _host, unsigned _id = 0); |
|
||||
|
|
||||
/// Start mining.
|
|
||||
void start() { startWorking(); } |
|
||||
|
|
||||
/// Stop mining.
|
|
||||
void stop() { stopWorking(); } |
|
||||
|
|
||||
/// Call to notify Miner of a state change.
|
|
||||
virtual void noteStateChange() override { m_miningStatus = Preparing; } |
|
||||
|
|
||||
/// @returns true iff the mining has been start()ed. It may still not be actually mining, depending on the host's turbo() & force().
|
|
||||
bool isRunning() const override { return isWorking(); } |
|
||||
|
|
||||
/// @returns true if mining is complete.
|
|
||||
virtual bool isComplete() const override { return m_miningStatus == Mined; } |
|
||||
|
|
||||
/// @returns the internal State object.
|
|
||||
virtual bytes const& blockData() const override { return m_mineState.blockData(); } |
|
||||
|
|
||||
/// Check the progress of the mining.
|
|
||||
MineProgress miningProgress() const { Guard l(x_mineInfo); return m_mineProgress; } |
|
||||
|
|
||||
/// Get and clear the mining history.
|
|
||||
std::list<MineInfo> miningHistory() { Guard l(x_mineInfo); auto ret = m_mineHistory; m_mineHistory.clear(); return ret; } |
|
||||
|
|
||||
/// @returns the state on which we mined.
|
|
||||
State const& state() const { return m_mineState; } |
|
||||
|
|
||||
private: |
|
||||
/// Do some work on the mining.
|
|
||||
virtual void doWork(); |
|
||||
|
|
||||
enum MiningStatus { Waiting, Preparing, Mining, Mined, Stopping, Stopped }; |
|
||||
MiningStatus m_miningStatus = Waiting; ///< TODO: consider mutex/atomic variable.
|
|
||||
State m_mineState; ///< The state on which we are mining, generally equivalent to m_postMine.
|
|
||||
std::unique_ptr<EthashPoW> m_pow; ///< Our miner.
|
|
||||
|
|
||||
mutable Mutex x_mineInfo; ///< Lock for the mining progress & history.
|
|
||||
MineProgress m_mineProgress; ///< What's our progress?
|
|
||||
std::list<MineInfo> m_mineHistory; ///< What the history of our mining?
|
|
||||
}; |
|
||||
|
|
||||
} |
|
||||
} |
|
||||
@ -0,0 +1,351 @@ |
|||||
|
/* |
||||
|
This file is part of cpp-ethereum. |
||||
|
|
||||
|
cpp-ethereum 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. |
||||
|
|
||||
|
cpp-ethereum 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 NetworkDeployment.js |
||||
|
* @author Arkadiy Paronyan arkadiy@ethdev.com |
||||
|
* @author Yann yann@ethdev.com |
||||
|
* @date 2015 |
||||
|
* Ethereum IDE client. |
||||
|
*/ |
||||
|
Qt.include("TransactionHelper.js") |
||||
|
|
||||
|
|
||||
|
var jsonRpcRequestId = 1; |
||||
|
function deployProject(force) { |
||||
|
saveAll(); //TODO: ask user
|
||||
|
deploymentDialog.open(); |
||||
|
} |
||||
|
|
||||
|
function startDeployProject(erasePrevious) |
||||
|
{ |
||||
|
var date = new Date(); |
||||
|
var deploymentId = date.toLocaleString(Qt.locale(), "ddMMyyHHmmsszzz"); |
||||
|
if (!erasePrevious) |
||||
|
{ |
||||
|
finalizeDeployment(deploymentId, projectModel.deploymentAddresses); |
||||
|
return; |
||||
|
} |
||||
|
|
||||
|
var jsonRpcUrl = "http://127.0.0.1:8080"; |
||||
|
console.log("Deploying " + deploymentId + " to " + jsonRpcUrl); |
||||
|
deploymentStarted(); |
||||
|
|
||||
|
var ctrNames = Object.keys(codeModel.contracts); |
||||
|
var ctrAddresses = {}; |
||||
|
deployContracts(0, ctrAddresses, ctrNames, function (){ |
||||
|
finalizeDeployment(deploymentId, ctrAddresses); |
||||
|
}); |
||||
|
} |
||||
|
|
||||
|
function deployContracts(ctrIndex, ctrAddresses, ctrNames, callBack) |
||||
|
{ |
||||
|
var code = codeModel.contracts[ctrNames[ctrIndex]].codeHex; |
||||
|
var requests = [{ |
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_sendTransaction", |
||||
|
params: [ { "from": deploymentDialog.currentAccount, "gas": deploymentDialog.gasToUse, "code": code } ], |
||||
|
id: 0 |
||||
|
}]; |
||||
|
rpcCall(requests, function (httpCall, response){ |
||||
|
var txt = qsTr("Please wait while " + ctrNames[ctrIndex] + " is published ...") |
||||
|
deploymentStepChanged(txt); |
||||
|
console.log(txt); |
||||
|
ctrAddresses[ctrNames[ctrIndex]] = JSON.parse(response)[0].result |
||||
|
deploymentDialog.waitForTrCountToIncrement(function(status) { |
||||
|
if (status === -1) |
||||
|
{ |
||||
|
trCountIncrementTimeOut(); |
||||
|
return; |
||||
|
} |
||||
|
ctrIndex++; |
||||
|
if (ctrIndex < ctrNames.length) |
||||
|
deployContracts(ctrIndex, ctrAddresses, ctrNames, callBack); |
||||
|
else |
||||
|
callBack(); |
||||
|
}); |
||||
|
}); |
||||
|
} |
||||
|
|
||||
|
function finalizeDeployment(deploymentId, addresses) { |
||||
|
deploymentStepChanged(qsTr("Packaging application ...")); |
||||
|
var deploymentDir = projectPath + deploymentId + "/"; |
||||
|
projectModel.deploymentDir = deploymentDir; |
||||
|
fileIo.makeDir(deploymentDir); |
||||
|
for (var i = 0; i < projectListModel.count; i++) { |
||||
|
var doc = projectListModel.get(i); |
||||
|
if (doc.isContract) |
||||
|
continue; |
||||
|
if (doc.isHtml) { |
||||
|
//inject the script to access contract API
|
||||
|
//TODO: use a template
|
||||
|
var html = fileIo.readFile(doc.path); |
||||
|
var insertAt = html.indexOf("<head>") |
||||
|
if (insertAt < 0) |
||||
|
insertAt = 0; |
||||
|
else |
||||
|
insertAt += 6; |
||||
|
html = html.substr(0, insertAt) + |
||||
|
"<script src=\"deployment.js\"></script>" + |
||||
|
html.substr(insertAt); |
||||
|
fileIo.writeFile(deploymentDir + doc.fileName, html); |
||||
|
} |
||||
|
else |
||||
|
fileIo.copyFile(doc.path, deploymentDir + doc.fileName); |
||||
|
} |
||||
|
//write deployment js
|
||||
|
var deploymentJs = |
||||
|
"// Autogenerated by Mix\n" + |
||||
|
"contracts = {};\n"; |
||||
|
for (var c in codeModel.contracts) { |
||||
|
var contractAccessor = "contracts[\"" + codeModel.contracts[c].contract.name + "\"]"; |
||||
|
deploymentJs += contractAccessor + " = {\n" + |
||||
|
"\tinterface: " + codeModel.contracts[c].contractInterface + ",\n" + |
||||
|
"\taddress: \"" + addresses[c] + "\"\n" + |
||||
|
"};\n" + |
||||
|
contractAccessor + ".contractClass = web3.eth.contract(" + contractAccessor + ".interface);\n" + |
||||
|
contractAccessor + ".contract = new " + contractAccessor + ".contractClass(" + contractAccessor + ".address);\n"; |
||||
|
} |
||||
|
fileIo.writeFile(deploymentDir + "deployment.js", deploymentJs); |
||||
|
deploymentAddresses = addresses; |
||||
|
saveProject(); |
||||
|
|
||||
|
var packageRet = fileIo.makePackage(deploymentDir); |
||||
|
deploymentDialog.packageHash = packageRet[0]; |
||||
|
deploymentDialog.packageBase64 = packageRet[1]; |
||||
|
deploymentDialog.localPackageUrl = packageRet[2] + "?hash=" + packageRet[0]; |
||||
|
|
||||
|
var applicationUrlEth = deploymentDialog.applicationUrlEth; |
||||
|
|
||||
|
applicationUrlEth = formatAppUrl(applicationUrlEth); |
||||
|
deploymentStepChanged(qsTr("Registering application on the Ethereum network ...")); |
||||
|
checkEthPath(applicationUrlEth, function () { |
||||
|
deploymentComplete(); |
||||
|
deployResourcesDialog.text = qsTr("Register Web Application to finalize deployment."); |
||||
|
deployResourcesDialog.open(); |
||||
|
}); |
||||
|
} |
||||
|
|
||||
|
function checkEthPath(dappUrl, callBack) |
||||
|
{ |
||||
|
if (dappUrl.length === 1) |
||||
|
registerContentHash(deploymentDialog.eth, callBack); // we directly create a dapp under the root registrar.
|
||||
|
else |
||||
|
{ |
||||
|
// the first owned reigstrar must have been created to follow the path.
|
||||
|
var str = createString(dappUrl[0]); |
||||
|
var requests = []; |
||||
|
requests.push({ |
||||
|
//register()
|
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_call", |
||||
|
params: [ { "gas": 150000, "from": deploymentDialog.currentAccount, "to": '0x' + deploymentDialog.eth, "data": "0x6be16bed" + str.encodeValueAsString() } ], |
||||
|
id: jsonRpcRequestId++ |
||||
|
}); |
||||
|
rpcCall(requests, function (httpRequest, response) { |
||||
|
var res = JSON.parse(response); |
||||
|
var addr = normalizeAddress(res[0].result); |
||||
|
if (addr.replace(/0+/g, "") === "") |
||||
|
{ |
||||
|
var errorTxt = qsTr("Path does not exists " + JSON.stringify(dappUrl) + ". Please register using Registration Dapp. Aborting."); |
||||
|
deploymentError(errorTxt); |
||||
|
console.log(errorTxt); |
||||
|
} |
||||
|
else |
||||
|
{ |
||||
|
dappUrl.splice(0, 1); |
||||
|
checkRegistration(dappUrl, addr, callBack); |
||||
|
} |
||||
|
}); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
function checkRegistration(dappUrl, addr, callBack) |
||||
|
{ |
||||
|
if (dappUrl.length === 1) |
||||
|
registerContentHash(addr, callBack); // We do not create the register for the last part, just registering the content hash.
|
||||
|
else |
||||
|
{ |
||||
|
var txt = qsTr("Checking " + JSON.stringify(dappUrl) + " ... in registrar " + addr); |
||||
|
deploymentStepChanged(txt); |
||||
|
console.log(txt); |
||||
|
var requests = []; |
||||
|
var registrar = {} |
||||
|
var str = createString(dappUrl[0]); |
||||
|
requests.push({ |
||||
|
//getOwner()
|
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_call", |
||||
|
params: [ { "gas" : 2000, "from": deploymentDialog.currentAccount, "to": '0x' + addr, "data": "0x893d20e8" } ], |
||||
|
id: jsonRpcRequestId++ |
||||
|
}); |
||||
|
|
||||
|
requests.push({ |
||||
|
//register()
|
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_call", |
||||
|
params: [ { "from": deploymentDialog.currentAccount, "to": '0x' + addr, "data": "0x6be16bed" + str.encodeValueAsString() } ], |
||||
|
id: jsonRpcRequestId++ |
||||
|
}); |
||||
|
|
||||
|
rpcCall(requests, function (httpRequest, response) { |
||||
|
var res = JSON.parse(response); |
||||
|
var nextAddr = normalizeAddress(res[1].result); |
||||
|
var errorTxt; |
||||
|
if (res[1].result === "0x") |
||||
|
{ |
||||
|
errorTxt = qsTr("Error when creating new owned regsitrar. Please use the regsitration Dapp. Aborting"); |
||||
|
deploymentError(errorTxt); |
||||
|
console.log(errorTxt); |
||||
|
} |
||||
|
else if (normalizeAddress(deploymentDialog.currentAccount) !== normalizeAddress(res[0].result)) |
||||
|
{ |
||||
|
errorTxt = qsTr("You are not the owner of " + dappUrl[0] + ". Aborting"); |
||||
|
deploymentError(errorTxt); |
||||
|
console.log(errorTxt); |
||||
|
} |
||||
|
else if (nextAddr.replace(/0+/g, "") !== "") |
||||
|
{ |
||||
|
dappUrl.splice(0, 1); |
||||
|
checkRegistration(dappUrl, nextAddr, callBack); |
||||
|
} |
||||
|
else |
||||
|
{ |
||||
|
var txt = qsTr("Registering sub domain " + dappUrl[0] + " ..."); |
||||
|
console.log(txt); |
||||
|
deploymentStepChanged(txt); |
||||
|
//current registrar is owned => ownedregistrar creation and continue.
|
||||
|
requests = []; |
||||
|
|
||||
|
requests.push({ |
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_sendTransaction", |
||||
|
params: [ { "from": deploymentDialog.currentAccount, "gas": 20000, "code": "0x60056013565b61059e8061001d6000396000f35b33600081905550560060003560e060020a90048063019848921461009a578063449c2090146100af5780635d574e32146100cd5780635fd4b08a146100e1578063618242da146100f65780636be16bed1461010b5780636c4489b414610129578063893d20e8146101585780639607730714610173578063c284bc2a14610187578063e50f599a14610198578063e5811b35146101af578063ec7b9200146101cd57005b6100a560043561031b565b8060005260206000f35b6100ba6004356103a0565b80600160a060020a031660005260206000f35b6100db600435602435610537565b60006000f35b6100ec600435610529565b8060005260206000f35b6101016004356103dd565b8060005260206000f35b6101166004356103bd565b80600160a060020a031660005260206000f35b61013460043561034b565b82600160a060020a031660005281600160a060020a03166020528060405260606000f35b610160610341565b80600160a060020a031660005260206000f35b6101816004356024356102b4565b60006000f35b6101926004356103fd565b60006000f35b6101a96004356024356044356101f2565b60006000f35b6101ba6004356101eb565b80600160a060020a031660005260206000f35b6101d8600435610530565b80600160a060020a031660005260206000f35b6000919050565b600054600160a060020a031633600160a060020a031614610212576102af565b8160026000858152602001908152602001600020819055508061023457610287565b81600160a060020a0316837f680ad70765443c2967675ab0fb91a46350c01c6df59bf9a41ff8a8dd097464ec60006000a3826001600084600160a060020a03168152602001908152602001600020819055505b827f18d67da0cd86808336a3aa8912f6ea70c5250f1a98b586d1017ef56fe199d4fc60006000a25b505050565b600054600160a060020a031633600160a060020a0316146102d457610317565b806002600084815260200190815260200160002060010181905550817f18d67da0cd86808336a3aa8912f6ea70c5250f1a98b586d1017ef56fe199d4fc60006000a25b5050565b60006001600083600160a060020a03168152602001908152602001600020549050919050565b6000600054905090565b6000600060006002600085815260200190815260200160002054925060026000858152602001908152602001600020600101549150600260008581526020019081526020016000206002015490509193909250565b600060026000838152602001908152602001600020549050919050565b600060026000838152602001908152602001600020600101549050919050565b600060026000838152602001908152602001600020600201549050919050565b600054600160a060020a031633600160a060020a03161461041d57610526565b80600160006002600085815260200190815260200160002054600160a060020a031681526020019081526020016000205414610458576104d2565b6002600082815260200190815260200160002054600160a060020a0316817f680ad70765443c2967675ab0fb91a46350c01c6df59bf9a41ff8a8dd097464ec60006000a36000600160006002600085815260200190815260200160002054600160a060020a03168152602001908152602001600020819055505b6002600082815260200190815260200160002060008101600090556001810160009055600281016000905550807f18d67da0cd86808336a3aa8912f6ea70c5250f1a98b586d1017ef56fe199d4fc60006000a25b50565b6000919050565b6000919050565b600054600160a060020a031633600160a060020a0316146105575761059a565b806002600084815260200190815260200160002060020181905550817f18d67da0cd86808336a3aa8912f6ea70c5250f1a98b586d1017ef56fe199d4fc60006000a25b505056" } ], |
||||
|
id: jsonRpcRequestId++ |
||||
|
}); |
||||
|
|
||||
|
rpcCall(requests, function(httpRequest, response) { |
||||
|
var newCtrAddress = normalizeAddress(JSON.parse(response)[0].result); |
||||
|
requests = []; |
||||
|
var txt = qsTr("Please wait " + dappUrl[0] + " is registering ..."); |
||||
|
deploymentStepChanged(txt); |
||||
|
console.log(txt); |
||||
|
deploymentDialog.waitForTrCountToIncrement(function(status) { |
||||
|
if (status === -1) |
||||
|
{ |
||||
|
trCountIncrementTimeOut(); |
||||
|
return; |
||||
|
} |
||||
|
var crLevel = createString(dappUrl[0]).encodeValueAsString(); |
||||
|
requests.push({ |
||||
|
//setRegister()
|
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_sendTransaction", |
||||
|
params: [ { "from": deploymentDialog.currentAccount, "gas": 30000, "to": '0x' + addr, "data": "0x96077307" + crLevel + deploymentDialog.pad(newCtrAddress) } ], |
||||
|
id: jsonRpcRequestId++ |
||||
|
}); |
||||
|
|
||||
|
rpcCall(requests, function(request, response){ |
||||
|
dappUrl.splice(0, 1); |
||||
|
checkRegistration(dappUrl, newCtrAddress, callBack); |
||||
|
}); |
||||
|
}); |
||||
|
}); |
||||
|
} |
||||
|
}); |
||||
|
} |
||||
|
} |
||||
|
|
||||
|
function trCountIncrementTimeOut() |
||||
|
{ |
||||
|
var error = qsTr("Something went wrong during the deployment. Please verify the amount of gas for this transaction and check your balance.") |
||||
|
console.log(error); |
||||
|
deploymentError(error); |
||||
|
} |
||||
|
|
||||
|
function registerContentHash(registrar, callBack) |
||||
|
{ |
||||
|
var txt = qsTr("Finalizing Dapp registration ..."); |
||||
|
deploymentStepChanged(txt); |
||||
|
console.log(txt); |
||||
|
var requests = []; |
||||
|
var paramTitle = clientModel.encodeAbiString(projectModel.projectTitle); |
||||
|
requests.push({ |
||||
|
//setContent()
|
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_sendTransaction", |
||||
|
params: [ { "from": deploymentDialog.currentAccount, "gas": 30000, "gasPrice": "10", "to": '0x' + registrar, "data": "0x5d574e32" + paramTitle + deploymentDialog.packageHash } ], |
||||
|
id: jsonRpcRequestId++ |
||||
|
}); |
||||
|
rpcCall(requests, function (httpRequest, response) { |
||||
|
callBack(); |
||||
|
}); |
||||
|
} |
||||
|
|
||||
|
function registerToUrlHint() |
||||
|
{ |
||||
|
deploymentStepChanged(qsTr("Registering application Resources (" + deploymentDialog.applicationUrlHttp) + ") ..."); |
||||
|
var requests = []; |
||||
|
var paramUrlHttp = createString(deploymentDialog.applicationUrlHttp); |
||||
|
requests.push({ |
||||
|
//urlHint => suggestUrl
|
||||
|
jsonrpc: "2.0", |
||||
|
method: "eth_sendTransaction", |
||||
|
params: [ { "to": '0x' + deploymentDialog.urlHintContract, "gas": 30000, "data": "0x4983e19c" + deploymentDialog.packageHash + paramUrlHttp.encodeValueAsString() } ], |
||||
|
id: jsonRpcRequestId++ |
||||
|
}); |
||||
|
|
||||
|
rpcCall(requests, function (httpRequest, response) { |
||||
|
deploymentComplete(); |
||||
|
}); |
||||
|
} |
||||
|
|
||||
|
function normalizeAddress(addr) |
||||
|
{ |
||||
|
addr = addr.replace('0x', ''); |
||||
|
if (addr.length <= 40) |
||||
|
return addr; |
||||
|
var left = addr.length - 40; |
||||
|
return addr.substring(left); |
||||
|
} |
||||
|
|
||||
|
function formatAppUrl(url) |
||||
|
{ |
||||
|
if (url.toLowerCase().indexOf("eth://") === 0) |
||||
|
url = url.substring(6); |
||||
|
if (url.toLowerCase().indexOf(projectModel.projectTitle + ".") === 0) |
||||
|
url = url.substring(projectModel.projectTitle.length + 1); |
||||
|
if (url === "") |
||||
|
return [projectModel.projectTitle]; |
||||
|
|
||||
|
var ret; |
||||
|
if (url.indexOf("/") === -1) |
||||
|
ret = url.split('.').reverse(); |
||||
|
else |
||||
|
{ |
||||
|
var slash = url.indexOf("/"); |
||||
|
var left = url.substring(0, slash); |
||||
|
var leftA = left.split("."); |
||||
|
leftA.reverse(); |
||||
|
|
||||
|
var right = url.substring(slash + 1); |
||||
|
var rightA = right.split('/'); |
||||
|
ret = leftA.concat(rightA); |
||||
|
} |
||||
|
if (ret[0].toLowerCase() === "eth") |
||||
|
ret.splice(0, 1); |
||||
|
ret.push(projectModel.projectTitle); |
||||
|
return ret; |
||||
|
} |
||||
@ -0,0 +1,20 @@ |
|||||
|
function test_getDefaultMiner() |
||||
|
{ |
||||
|
newProject(); |
||||
|
var state = mainApplication.projectModel.stateListModel.get(0); |
||||
|
compare(state.miner.secret, "cb73d9408c4720e230387d956eb0f829d8a4dd2c1055f96257167e14e7169074"); |
||||
|
} |
||||
|
|
||||
|
function test_selectMiner() |
||||
|
{ |
||||
|
newProject(); |
||||
|
mainApplication.projectModel.stateListModel.editState(0); |
||||
|
var account = mainApplication.projectModel.stateDialog.newAccAction.add(); |
||||
|
account = mainApplication.projectModel.stateDialog.newAccAction.add(); |
||||
|
mainApplication.projectModel.stateDialog.minerComboBox.currentIndex = 2; |
||||
|
ts.waitForRendering(mainApplication.projectModel.stateDialog.minerComboBox, 3000); |
||||
|
mainApplication.projectModel.stateDialog.acceptAndClose(); |
||||
|
var state = mainApplication.projectModel.stateListModel.get(0); |
||||
|
compare(state.miner.secret, account.secret); |
||||
|
} |
||||
|
|
||||
@ -0,0 +1,18 @@ |
|||||
|
function test_contractRename() |
||||
|
{ |
||||
|
newProject(); |
||||
|
tryCompare(mainApplication.mainContent.projectNavigator.sections.itemAt(0).model.get(0), "name", "Contract"); |
||||
|
editContract("contract Renamed {}"); |
||||
|
if (!ts.waitForSignal(mainApplication.clientModel, "runComplete()", 5000)) |
||||
|
fail("Error running transaction"); |
||||
|
wait(1000); |
||||
|
tryCompare(mainApplication.mainContent.projectNavigator.sections.itemAt(0).model.get(0), "name", "Renamed"); |
||||
|
mainApplication.projectModel.stateListModel.editState(0); |
||||
|
mainApplication.projectModel.stateDialog.model.editTransaction(2); |
||||
|
var transactionDialog = mainApplication.projectModel.stateDialog.transactionDialog; |
||||
|
tryCompare(transactionDialog, "contractId", "Renamed"); |
||||
|
tryCompare(transactionDialog, "functionId", "Renamed"); |
||||
|
transactionDialog.close(); |
||||
|
mainApplication.projectModel.stateDialog.close(); |
||||
|
tryCompare(mainApplication.mainContent.rightPane.transactionLog.transactionModel.get(2), "contract", "Renamed"); |
||||
|
} |
||||
Loading…
Reference in new issue