Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
whqwjb
go-ethereum
提交
a9d6846f
G
go-ethereum
项目概览
whqwjb
/
go-ethereum
与 Fork 源项目一致
从无法访问的项目Fork
通知
1
Star
0
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
G
go-ethereum
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
提交
Issue看板
体验新版 GitCode,发现更多精彩内容 >>
提交
a9d6846f
编写于
6月 16, 2015
作者:
G
Gustav Simonsson
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
Update ethash Godep
上级
8f372c86
变更
16
展开全部
隐藏空白更改
内联
并排
Showing
16 changed file
with
76 addition
and
13958 deletion
+76
-13958
Godeps/Godeps.json
Godeps/Godeps.json
+2
-2
Godeps/_workspace/src/github.com/ethereum/ethash/CMakeLists.txt
.../_workspace/src/github.com/ethereum/ethash/CMakeLists.txt
+0
-7
Godeps/_workspace/src/github.com/ethereum/ethash/ethash.go
Godeps/_workspace/src/github.com/ethereum/ethash/ethash.go
+7
-0
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/CMakeLists.txt
...ithub.com/ethereum/ethash/src/libethash-cl/CMakeLists.txt
+0
-47
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/bin2h.cmake
...c/github.com/ethereum/ethash/src/libethash-cl/bin2h.cmake
+0
-86
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/cl.hpp
...ce/src/github.com/ethereum/ethash/src/libethash-cl/cl.hpp
+0
-12906
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner.cpp
....com/ethereum/ethash/src/libethash-cl/ethash_cl_miner.cpp
+0
-384
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner.h
...ub.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner.h
+0
-57
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner_kernel.cl
...thereum/ethash/src/libethash-cl/ethash_cl_miner_kernel.cl
+0
-460
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/endian.h
...ace/src/github.com/ethereum/ethash/src/libethash/endian.h
+3
-0
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/internal.c
...e/src/github.com/ethereum/ethash/src/libethash/internal.c
+12
-1
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.c
...rkspace/src/github.com/ethereum/ethash/src/libethash/io.c
+19
-2
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.h
...rkspace/src/github.com/ethereum/ethash/src/libethash/io.h
+17
-0
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_posix.c
...e/src/github.com/ethereum/ethash/src/libethash/io_posix.c
+9
-0
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_win32.c
...e/src/github.com/ethereum/ethash/src/libethash/io_win32.c
+2
-2
Godeps/_workspace/src/github.com/ethereum/ethash/test/c/test.cpp
..._workspace/src/github.com/ethereum/ethash/test/c/test.cpp
+5
-4
未找到文件。
Godeps/Godeps.json
浏览文件 @
a9d6846f
...
...
@@ -21,8 +21,8 @@
},
{
"ImportPath"
:
"github.com/ethereum/ethash"
,
"Comment"
:
"v23.1-2
06-gf0e6321
"
,
"Rev"
:
"
f0e63218b721dc2f696920a92d5de1f6364e9bf
7"
"Comment"
:
"v23.1-2
22-g5cfdcba
"
,
"Rev"
:
"
5cfdcba92e634db228d1ddb140e3b7a3c4b3817
7"
},
{
"ImportPath"
:
"github.com/howeyc/fsnotify"
,
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/CMakeLists.txt
浏览文件 @
a9d6846f
...
...
@@ -9,13 +9,6 @@ if (WIN32 AND WANT_CRYPTOPP)
endif
()
add_subdirectory
(
src/libethash
)
# bin2h.cmake doesn't work
if
(
NOT OpenCL_FOUND
)
find_package
(
OpenCL
)
endif
()
if
(
OpenCL_FOUND
)
add_subdirectory
(
src/libethash-cl
)
endif
()
add_subdirectory
(
src/benchmark EXCLUDE_FROM_ALL
)
add_subdirectory
(
test/c
)
Godeps/_workspace/src/github.com/ethereum/ethash/ethash.go
浏览文件 @
a9d6846f
...
...
@@ -8,6 +8,7 @@ int ethashGoCallback_cgo(unsigned);
import
"C"
import
(
"bytes"
"errors"
"fmt"
"io/ioutil"
...
...
@@ -119,6 +120,12 @@ func (l *Light) Verify(block pow.Block) bool {
if
!
ret
.
success
{
return
false
}
// avoid mixdigest malleability as it's not included in a block's "hashNononce"
if
!
bytes
.
Equal
(
block
.
MixDigest
()
.
Bytes
(),
C
.
GoBytes
(
unsafe
.
Pointer
(
&
ret
.
mix_hash
),
C
.
int
(
32
)))
{
return
false
}
// Make sure cache is live until after the C call.
// This is important because a GC might happen and execute
// the finalizer before the call completes.
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/CMakeLists.txt
已删除
100644 → 0
浏览文件 @
8f372c86
cmake_minimum_required
(
VERSION 2.8
)
set
(
LIBRARY ethash-cl
)
set
(
CMAKE_BUILD_TYPE Release
)
include
(
bin2h.cmake
)
bin2h
(
SOURCE_FILE ethash_cl_miner_kernel.cl
VARIABLE_NAME ethash_cl_miner_kernel
HEADER_FILE
${
CMAKE_CURRENT_BINARY_DIR
}
/ethash_cl_miner_kernel.h
)
if
(
NOT MSVC
)
# Initialize CXXFLAGS for c++11
set
(
CMAKE_CXX_FLAGS
"-Wall -std=c++11"
)
set
(
CMAKE_CXX_FLAGS_DEBUG
"-O0 -g"
)
set
(
CMAKE_CXX_FLAGS_MINSIZEREL
"-Os -DNDEBUG"
)
set
(
CMAKE_CXX_FLAGS_RELEASE
"-O3 -DNDEBUG"
)
set
(
CMAKE_CXX_FLAGS_RELWITHDEBINFO
"-O2 -g"
)
# Compiler-specific C++11 activation.
if
(
"
${
CMAKE_CXX_COMPILER_ID
}
"
MATCHES
"GNU"
)
execute_process
(
COMMAND
${
CMAKE_CXX_COMPILER
}
-dumpversion OUTPUT_VARIABLE GCC_VERSION
)
if
(
NOT
(
GCC_VERSION VERSION_GREATER 4.7 OR GCC_VERSION VERSION_EQUAL 4.7
))
message
(
FATAL_ERROR
"
${
PROJECT_NAME
}
requires g++ 4.7 or greater."
)
endif
()
elseif
(
"
${
CMAKE_CXX_COMPILER_ID
}
"
MATCHES
"Clang"
)
set
(
CMAKE_CXX_FLAGS
"
${
CMAKE_CXX_FLAGS
}
-stdlib=libc++"
)
else
()
message
(
FATAL_ERROR
"Your C++ compiler does not support C++11."
)
endif
()
endif
()
set
(
OpenCL_FOUND TRUE
)
set
(
OpenCL_INCLUDE_DIRS /usr/include/CL
)
set
(
OpenCL_LIBRARIES -lOpenCL
)
if
(
NOT OpenCL_FOUND
)
find_package
(
OpenCL
)
endif
()
if
(
OpenCL_FOUND
)
set
(
CMAKE_CXX_FLAGS
"-std=c++11 -Wall -Wno-unknown-pragmas -Wextra -Werror -pedantic -fPIC
${
CMAKE_CXX_FLAGS
}
"
)
include_directories
(
${
OpenCL_INCLUDE_DIRS
}
${
CMAKE_CURRENT_BINARY_DIR
}
)
include_directories
(
..
)
add_library
(
${
LIBRARY
}
ethash_cl_miner.cpp ethash_cl_miner.h cl.hpp
)
TARGET_LINK_LIBRARIES
(
${
LIBRARY
}
${
OpenCL_LIBRARIES
}
ethash
)
endif
()
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/bin2h.cmake
已删除
100644 → 0
浏览文件 @
8f372c86
# 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
()
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/cl.hpp
已删除
100644 → 0
浏览文件 @
8f372c86
此差异已折叠。
点击以展开。
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner.cpp
已删除
100644 → 0
浏览文件 @
8f372c86
/*
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 <iostream>
#include <assert.h>
#include <queue>
#include <vector>
#include <libethash/util.h>
#include <libethash/ethash.h>
#include <libethash/internal.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
using
namespace
std
;
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
(
unsigned
_platformId
,
unsigned
_deviceId
)
{
std
::
vector
<
cl
::
Platform
>
platforms
;
cl
::
Platform
::
get
(
&
platforms
);
if
(
platforms
.
empty
())
{
cout
<<
"No OpenCL platforms found."
<<
endl
;
return
std
::
string
();
}
// get GPU device of the selected platform
std
::
vector
<
cl
::
Device
>
devices
;
unsigned
platform_num
=
std
::
min
<
unsigned
>
(
_platformId
,
platforms
.
size
()
-
1
);
platforms
[
platform_num
].
getDevices
(
CL_DEVICE_TYPE_ALL
,
&
devices
);
if
(
devices
.
empty
())
{
cout
<<
"No OpenCL devices found."
<<
endl
;
return
std
::
string
();
}
// use selected default device
unsigned
device_num
=
std
::
min
<
unsigned
>
(
_deviceId
,
devices
.
size
()
-
1
);
cl
::
Device
&
device
=
devices
[
device_num
];
std
::
string
device_version
=
device
.
getInfo
<
CL_DEVICE_VERSION
>
();
return
"{
\"
platform
\"
:
\"
"
+
platforms
[
platform_num
].
getInfo
<
CL_PLATFORM_NAME
>
()
+
"
\"
,
\"
device
\"
:
\"
"
+
device
.
getInfo
<
CL_DEVICE_NAME
>
()
+
"
\"
,
\"
version
\"
:
\"
"
+
device_version
+
"
\"
}"
;
}
unsigned
ethash_cl_miner
::
get_num_devices
(
unsigned
_platformId
)
{
std
::
vector
<
cl
::
Platform
>
platforms
;
cl
::
Platform
::
get
(
&
platforms
);
if
(
platforms
.
empty
())
{
cout
<<
"No OpenCL platforms found."
<<
endl
;
return
0
;
}
std
::
vector
<
cl
::
Device
>
devices
;
unsigned
platform_num
=
std
::
min
<
unsigned
>
(
_platformId
,
platforms
.
size
()
-
1
);
platforms
[
platform_num
].
getDevices
(
CL_DEVICE_TYPE_ALL
,
&
devices
);
if
(
devices
.
empty
())
{
cout
<<
"No OpenCL devices found."
<<
endl
;
return
0
;
}
return
devices
.
size
();
}
void
ethash_cl_miner
::
finish
()
{
if
(
m_queue
())
m_queue
.
finish
();
}
bool
ethash_cl_miner
::
init
(
uint64_t
block_number
,
std
::
function
<
void
(
void
*
)
>
_fillDAG
,
unsigned
workgroup_size
,
unsigned
_platformId
,
unsigned
_deviceId
)
{
// store params
m_fullSize
=
ethash_get_datasize
(
block_number
);
// get all platforms
std
::
vector
<
cl
::
Platform
>
platforms
;
cl
::
Platform
::
get
(
&
platforms
);
if
(
platforms
.
empty
())
{
cout
<<
"No OpenCL platforms found."
<<
endl
;
return
false
;
}
// use selected platform
_platformId
=
std
::
min
<
unsigned
>
(
_platformId
,
platforms
.
size
()
-
1
);
cout
<<
"Using platform: "
<<
platforms
[
_platformId
].
getInfo
<
CL_PLATFORM_NAME
>
().
c_str
()
<<
endl
;
// get GPU device of the default platform
std
::
vector
<
cl
::
Device
>
devices
;
platforms
[
_platformId
].
getDevices
(
CL_DEVICE_TYPE_ALL
,
&
devices
);
if
(
devices
.
empty
())
{
cout
<<
"No OpenCL devices found."
<<
endl
;
return
false
;
}
// use selected device
cl
::
Device
&
device
=
devices
[
std
::
min
<
unsigned
>
(
_deviceId
,
devices
.
size
()
-
1
)];
std
::
string
device_version
=
device
.
getInfo
<
CL_DEVICE_VERSION
>
();
cout
<<
"Using device: "
<<
device
.
getInfo
<
CL_DEVICE_NAME
>
().
c_str
()
<<
"("
<<
device_version
.
c_str
()
<<
")"
<<
endl
;
if
(
strncmp
(
"OpenCL 1.0"
,
device_version
.
c_str
(),
10
)
==
0
)
{
cout
<<
"OpenCL 1.0 is not supported."
<<
endl
;
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
)(
m_fullSize
/
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
)
{
cout
<<
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
,
m_fullSize
);
// 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
,
m_fullSize
);
// 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
}
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner.h
已删除
100644 → 0
浏览文件 @
8f372c86
#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
(
uint64_t
block_number
,
std
::
function
<
void
(
void
*
)
>
_fillDAG
,
unsigned
workgroup_size
=
64
,
unsigned
_platformId
=
0
,
unsigned
_deviceId
=
0
);
static
std
::
string
platform_info
(
unsigned
_platformId
=
0
,
unsigned
_deviceId
=
0
);
static
unsigned
get_num_devices
(
unsigned
_platformId
=
0
);
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
};
uint64_t
m_fullSize
;
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
;
};
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash-cl/ethash_cl_miner_kernel.cl
已删除
100644 → 0
浏览文件 @
8f372c86
//
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
;
}
}
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/endian.h
浏览文件 @
a9d6846f
...
...
@@ -32,6 +32,9 @@
#include <libkern/OSByteOrder.h>
#define ethash_swap_u32(input_) OSSwapInt32(input_)
#define ethash_swap_u64(input_) OSSwapInt64(input_)
#elif defined(__FreeBSD__) || defined(__DragonFly__) || defined(__NetBSD__)
#define ethash_swap_u32(input_) bswap32(input_)
#define ethash_swap_u64(input_) bswap64(input_)
#else // posix
#include <byteswap.h>
#define ethash_swap_u32(input_) __bswap_32(input_)
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/internal.c
浏览文件 @
a9d6846f
...
...
@@ -364,6 +364,7 @@ static bool ethash_mmap(struct ethash_full* ret, FILE* f)
{
int
fd
;
char
*
mmapped_data
;
errno
=
0
;
ret
->
file
=
f
;
if
((
fd
=
ethash_fileno
(
ret
->
file
))
==
-
1
)
{
return
false
;
...
...
@@ -400,38 +401,48 @@ ethash_full_t ethash_full_new_internal(
ret
->
file_size
=
(
size_t
)
full_size
;
switch
(
ethash_io_prepare
(
dirname
,
seed_hash
,
&
f
,
(
size_t
)
full_size
,
false
))
{
case
ETHASH_IO_FAIL
:
// ethash_io_prepare will do all ETHASH_CRITICAL() logging in fail case
goto
fail_free_full
;
case
ETHASH_IO_MEMO_MATCH
:
if
(
!
ethash_mmap
(
ret
,
f
))
{
ETHASH_CRITICAL
(
"mmap failure()"
);
goto
fail_close_file
;
}
return
ret
;
case
ETHASH_IO_MEMO_SIZE_MISMATCH
:
// if a DAG of same filename but unexpected size is found, silently force new file creation
if
(
ethash_io_prepare
(
dirname
,
seed_hash
,
&
f
,
(
size_t
)
full_size
,
true
)
!=
ETHASH_IO_MEMO_MISMATCH
)
{
ETHASH_CRITICAL
(
"Could not recreate DAG file after finding existing DAG with unexpected size."
);
goto
fail_free_full
;
}
// fallthrough to the mismatch case here, DO NOT go through match
case
ETHASH_IO_MEMO_MISMATCH
:
if
(
!
ethash_mmap
(
ret
,
f
))
{
ETHASH_CRITICAL
(
"mmap failure()"
);
goto
fail_close_file
;
}
break
;
}
if
(
!
ethash_compute_full_data
(
ret
->
data
,
full_size
,
light
,
callback
))
{
ETHASH_CRITICAL
(
"Failure at computing DAG data."
);
goto
fail_free_full_data
;
}
// after the DAG has been filled then we finalize it by writting the magic number at the beginning
if
(
fseek
(
f
,
0
,
SEEK_SET
)
!=
0
)
{
ETHASH_CRITICAL
(
"Could not seek to DAG file start to write magic number."
);
goto
fail_free_full_data
;
}
uint64_t
const
magic_num
=
ETHASH_DAG_MAGIC_NUM
;
if
(
fwrite
(
&
magic_num
,
ETHASH_DAG_MAGIC_NUM_SIZE
,
1
,
f
)
!=
1
)
{
ETHASH_CRITICAL
(
"Could not write magic number to DAG's beginning."
);
goto
fail_free_full_data
;
}
if
(
fflush
(
f
)
!=
0
)
{
// make sure the magic number IS there
ETHASH_CRITICAL
(
"Could not flush memory mapped data to DAG file. Insufficient space?"
);
goto
fail_free_full_data
;
}
fflush
(
f
);
// make sure the magic number IS there
return
ret
;
fail_free_full_data:
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.c
浏览文件 @
a9d6846f
...
...
@@ -21,6 +21,7 @@
#include "io.h"
#include <string.h>
#include <stdio.h>
#include <errno.h>
enum
ethash_io_rc
ethash_io_prepare
(
char
const
*
dirname
,
...
...
@@ -32,15 +33,19 @@ enum ethash_io_rc ethash_io_prepare(
{
char
mutable_name
[
DAG_MUTABLE_NAME_MAX_SIZE
];
enum
ethash_io_rc
ret
=
ETHASH_IO_FAIL
;
// reset errno before io calls
errno
=
0
;
// assert directory exists
if
(
!
ethash_mkdir
(
dirname
))
{
ETHASH_CRITICAL
(
"Could not create the ethash directory"
);
goto
end
;
}
ethash_io_mutable_name
(
ETHASH_REVISION
,
&
seedhash
,
mutable_name
);
char
*
tmpfile
=
ethash_io_create_filename
(
dirname
,
mutable_name
,
strlen
(
mutable_name
));
if
(
!
tmpfile
)
{
ETHASH_CRITICAL
(
"Could not create the full DAG pathname"
);
goto
end
;
}
...
...
@@ -52,6 +57,7 @@ enum ethash_io_rc ethash_io_prepare(
size_t
found_size
;
if
(
!
ethash_file_size
(
f
,
&
found_size
))
{
fclose
(
f
);
ETHASH_CRITICAL
(
"Could not query size of DAG file:
\"
%s
\"
"
,
tmpfile
);
goto
free_memo
;
}
if
(
file_size
!=
found_size
-
ETHASH_DAG_MAGIC_NUM_SIZE
)
{
...
...
@@ -64,6 +70,7 @@ enum ethash_io_rc ethash_io_prepare(
if
(
fread
(
&
magic_num
,
ETHASH_DAG_MAGIC_NUM_SIZE
,
1
,
f
)
!=
1
)
{
// I/O error
fclose
(
f
);
ETHASH_CRITICAL
(
"Could not read from DAG file:
\"
%s
\"
"
,
tmpfile
);
ret
=
ETHASH_IO_MEMO_SIZE_MISMATCH
;
goto
free_memo
;
}
...
...
@@ -80,15 +87,25 @@ enum ethash_io_rc ethash_io_prepare(
// file does not exist, will need to be created
f
=
ethash_fopen
(
tmpfile
,
"wb+"
);
if
(
!
f
)
{
ETHASH_CRITICAL
(
"Could not create DAG file:
\"
%s
\"
"
,
tmpfile
);
goto
free_memo
;
}
// make sure it's of the proper size
if
(
fseek
(
f
,
(
long
int
)(
file_size
+
ETHASH_DAG_MAGIC_NUM_SIZE
-
1
),
SEEK_SET
)
!=
0
)
{
fclose
(
f
);
ETHASH_CRITICAL
(
"Could not seek to the end of DAG file:
\"
%s
\"
. Insufficient space?"
,
tmpfile
);
goto
free_memo
;
}
if
(
fputc
(
'\n'
,
f
)
==
EOF
)
{
fclose
(
f
);
ETHASH_CRITICAL
(
"Could not write in the end of DAG file:
\"
%s
\"
. Insufficient space?"
,
tmpfile
);
goto
free_memo
;
}
if
(
fflush
(
f
)
!=
0
)
{
fclose
(
f
);
ETHASH_CRITICAL
(
"Could not flush at end of DAG file:
\"
%s
\"
. Insufficient space?"
,
tmpfile
);
goto
free_memo
;
}
fputc
(
'\n'
,
f
);
fflush
(
f
);
ret
=
ETHASH_IO_MEMO_MISMATCH
;
goto
set_file
;
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io.h
浏览文件 @
a9d6846f
...
...
@@ -54,6 +54,23 @@ enum ethash_io_rc {
#define snprintf(...) sprintf_s(__VA_ARGS__)
#endif
/**
* Logs a critical error in important parts of ethash. Should mostly help
* figure out what kind of problem (I/O, memory e.t.c.) causes a NULL
* ethash_full_t
*/
#ifdef ETHASH_PRINT_CRITICAL_OUTPUT
#define ETHASH_CRITICAL(...) \
do \
{ \
printf("ETHASH CRITICAL ERROR: "__VA_ARGS__); \
printf("\n"); \
fflush(stdout); \
} while (0)
#else
#define ETHASH_CRITICAL(...)
#endif
/**
* Prepares io for ethash
*
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_posix.c
浏览文件 @
a9d6846f
...
...
@@ -26,6 +26,8 @@
#include <libgen.h>
#include <stdio.h>
#include <unistd.h>
#include <stdlib.h>
#include <pwd.h>
FILE
*
ethash_fopen
(
char
const
*
file_name
,
char
const
*
mode
)
{
...
...
@@ -89,6 +91,13 @@ bool ethash_get_default_dirname(char* strbuf, size_t buffsize)
static
const
char
dir_suffix
[]
=
".ethash/"
;
strbuf
[
0
]
=
'\0'
;
char
*
home_dir
=
getenv
(
"HOME"
);
if
(
!
home_dir
||
strlen
(
home_dir
)
==
0
)
{
struct
passwd
*
pwd
=
getpwuid
(
getuid
());
if
(
pwd
)
home_dir
=
pwd
->
pw_dir
;
}
size_t
len
=
strlen
(
home_dir
);
if
(
!
ethash_strncat
(
strbuf
,
buffsize
,
home_dir
,
len
))
{
return
false
;
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/src/libethash/io_win32.c
浏览文件 @
a9d6846f
...
...
@@ -87,9 +87,9 @@ bool ethash_file_size(FILE* f, size_t* ret_size)
bool
ethash_get_default_dirname
(
char
*
strbuf
,
size_t
buffsize
)
{
static
const
char
dir_suffix
[]
=
"
Appdata
\\
Ethash
\\
"
;
static
const
char
dir_suffix
[]
=
"Ethash
\\
"
;
strbuf
[
0
]
=
'\0'
;
if
(
!
SUCCEEDED
(
SHGetFolderPath
W
(
NULL
,
CSIDL_PROFILE
,
NULL
,
0
,
(
W
CHAR
*
)
strbuf
)))
{
if
(
!
SUCCEEDED
(
SHGetFolderPath
A
(
NULL
,
CSIDL_LOCAL_APPDATA
,
NULL
,
0
,
(
CHAR
*
)
strbuf
)))
{
return
false
;
}
if
(
!
ethash_strncat
(
strbuf
,
buffsize
,
"
\\
"
,
1
))
{
...
...
Godeps/_workspace/src/github.com/ethereum/ethash/test/c/test.cpp
浏览文件 @
a9d6846f
...
...
@@ -292,12 +292,13 @@ BOOST_AUTO_TEST_CASE(test_ethash_io_memo_file_size_mismatch) {
BOOST_AUTO_TEST_CASE
(
test_ethash_get_default_dirname
)
{
char
result
[
256
];
// this is really not an easy thing to test for in a unit test, so yeah it does look ugly
// this is really not an easy thing to test for in a unit test
// TODO: Improve this test ...
#ifdef _WIN32
char
homedir
[
256
];
BOOST_REQUIRE
(
SUCCEEDED
(
SHGetFolderPath
W
(
NULL
,
CSIDL_PROFILE
,
NULL
,
0
,
(
W
CHAR
*
)
homedir
)));
BOOST_REQUIRE
(
SUCCEEDED
(
SHGetFolderPath
A
(
NULL
,
CSIDL_PROFILE
,
NULL
,
0
,
(
CHAR
*
)
homedir
)));
BOOST_REQUIRE
(
ethash_get_default_dirname
(
result
,
256
));
std
::
string
res
=
std
::
string
(
homedir
)
+
std
::
string
(
"
\\
App
data
\\
Ethash
\\
"
);
std
::
string
res
=
std
::
string
(
homedir
)
+
std
::
string
(
"
\\
App
Data
\\
Local
\\
Ethash
\\
"
);
#else
char
*
homedir
=
getenv
(
"HOME"
);
BOOST_REQUIRE
(
ethash_get_default_dirname
(
result
,
256
));
...
...
@@ -305,7 +306,7 @@ BOOST_AUTO_TEST_CASE(test_ethash_get_default_dirname) {
#endif
BOOST_CHECK_MESSAGE
(
strcmp
(
res
.
c_str
(),
result
)
==
0
,
"Expected
\"
"
+
res
+
"
\"
but got
\"
"
+
std
::
string
(
result
)
+
"
\"
"
);
);
}
BOOST_AUTO_TEST_CASE
(
light_and_full_client_checks
)
{
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录