From ecd7559e9bb687558314eb67a36b97305119de19 Mon Sep 17 00:00:00 2001 From: hackyminer Date: Sat, 5 Sep 2020 14:57:52 +0900 Subject: [PATCH 01/16] OpenCL: reduce GenerateDAG() time using shared __local memory --- libethash-cl/kernels/cl/ethash.cl | 28 +++++++++++++++++++++------- 1 file changed, 21 insertions(+), 7 deletions(-) diff --git a/libethash-cl/kernels/cl/ethash.cl b/libethash-cl/kernels/cl/ethash.cl index b793dcc90..6bf715155 100644 --- a/libethash-cl/kernels/cl/ethash.cl +++ b/libethash-cl/kernels/cl/ethash.cl @@ -436,23 +436,37 @@ static void SHA3_512(uint2 *s) __kernel void GenerateDAG(uint start, __global const uint16 *_Cache, __global uint16 *_DAG0, __global uint16 *_DAG1, uint light_size) { __global const Node *Cache = (__global const Node *) _Cache; - uint NodeIdx = start + get_global_id(0); + const uint gid = get_global_id(0); + uint NodeIdx = start + gid; + const uint thread_id = gid & 3; + + __local Node sharebuf[WORKSIZE]; + __local uint indexbuf[WORKSIZE]; + __local Node *dagNode = sharebuf + (get_local_id(0) / 4) * 4; + __local uint *indexes = indexbuf + (get_local_id(0) / 4) * 4; + __global const Node *parentNode; Node DAGNode = Cache[NodeIdx % light_size]; DAGNode.dwords[0] ^= NodeIdx; SHA3_512(DAGNode.qwords); + dagNode[thread_id] = DAGNode; + barrier(CLK_LOCAL_MEM_FENCE); for (uint i = 0; i < 256; ++i) { - uint ParentIdx = fnv(NodeIdx ^ i, DAGNode.dwords[i & 15]) % light_size; - __global const Node *ParentNode = Cache + ParentIdx; + uint ParentIdx = fnv(NodeIdx ^ i, dagNode[thread_id].dwords[i & 15]) % light_size; + indexes[thread_id] = ParentIdx; + barrier(CLK_LOCAL_MEM_FENCE); -#pragma unroll - for (uint x = 0; x < 4; ++x) { - DAGNode.dqwords[x] *= (uint4)(FNV_PRIME); - DAGNode.dqwords[x] ^= ParentNode->dqwords[x]; + for (uint t = 0; t < 4; ++t) { + uint parentIndex = indexes[t]; + parentNode = Cache + parentIndex; + + dagNode[t].dqwords[thread_id] = fnv(dagNode[t].dqwords[thread_id], parentNode->dqwords[thread_id]); + barrier(CLK_LOCAL_MEM_FENCE); } } + DAGNode = dagNode[thread_id]; SHA3_512(DAGNode.qwords); From 1536abba52dbc7f0812dec37a448568cd53563d9 Mon Sep 17 00:00:00 2001 From: ianzur Date: Thu, 10 Sep 2020 16:34:11 -0500 Subject: [PATCH 02/16] config option to use system OpenCL (for ROCm users) --- CMakeLists.txt | 10 +++++++++- docs/BUILD.md | 3 +++ libethash-cl/CMakeLists.txt | 4 ++-- 3 files changed, 14 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index cdb5e6082..6a2a07a7f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -31,6 +31,7 @@ option(ETHDBUS "Build with D-Bus support" OFF) option(APICORE "Build with API Server support" ON) option(BINKERN "Install AMD binary kernels" ON) option(DEVBUILD "Log developer metrics" OFF) +option(USE_SYS_OPENCL "Build with system OpenCL" OFF) # propagates CMake configuration options to the compiler function(configureProject) @@ -55,6 +56,9 @@ function(configureProject) if (DEVBUILD) add_definitions(-DDEV_BUILD) endif() + if (USE_SYS_OPENCL) + add_definitions(-DUSE_SYS_OPENCL) + endif() endfunction() hunter_add_package(Boost COMPONENTS system filesystem thread) @@ -68,6 +72,10 @@ find_package(ethash CONFIG REQUIRED) configureProject() +if(APPLE) + set(USE_SYS_OPENCL ON) +endif() + message("----------------------------------------------------------------------------") message("-- CMake ${CMAKE_VERSION}") message("-- Build ${CMAKE_BUILD_TYPE} / ${CMAKE_SYSTEM_NAME}") @@ -79,6 +87,7 @@ message("-- ETHDBUS Build D-Bus components ${ETHD message("-- APICORE Build API Server components ${APICORE}") message("-- BINKERN Install AMD binary kernels ${BINKERN}") message("-- DEVBUILD Build with dev logging ${DEVBUILD}") +message("-- USE_SYS_OPENCL Build with system OpenCL ${USE_SYS_OPENCL}") message("----------------------------------------------------------------------------") message("") @@ -87,7 +96,6 @@ if(UNIX AND NOT APPLE) set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -static-libstdc++") endif() - cable_add_buildinfo_library(PROJECT_NAME ${PROJECT_NAME}) add_subdirectory(libdevcore) diff --git a/docs/BUILD.md b/docs/BUILD.md index 44649083e..32f9ceff7 100644 --- a/docs/BUILD.md +++ b/docs/BUILD.md @@ -43,6 +43,8 @@ you have to install the OpenGL libraries. E.g. on Ubuntu run: sudo apt-get install mesa-common-dev ``` +If you want to use locally installed [ROCm-OpenCL](https://rocmdocs.amd.com/en/latest/) package, use build flag `-DUSE_SYS_OPENCL=ON` with cmake config. + ### macOS 1. GCC version >= TBF @@ -143,6 +145,7 @@ cmake .. -DETHASHCUDA=ON -DETHASHCL=OFF * `-DAPICORE=ON` - enable API Server, `ON` by default. * `-DBINKERN=ON` - install AMD binary kernels, `ON` by default. * `-DETHDBUS=ON` - enable D-Bus support, `OFF` by default. +* `-DUSE_SYS_OPENCL=ON` - Use system OpenCL, `OFF` by default, unless on macOS. Specify to use local **ROCm-OpenCL** package. ## Disable Hunter diff --git a/libethash-cl/CMakeLists.txt b/libethash-cl/CMakeLists.txt index 7aabf0445..f190b97ba 100644 --- a/libethash-cl/CMakeLists.txt +++ b/libethash-cl/CMakeLists.txt @@ -20,8 +20,8 @@ set(SOURCES ${CMAKE_CURRENT_BINARY_DIR}/ethash.h ) -if(APPLE) - # On macOS use system OpenCL library. +if(USE_SYS_OPENCL) + # On macOS or using ROCm-OpenCL, use system OpenCL library. find_package(OpenCL REQUIRED) else() hunter_add_package(OpenCL) From 515df7aec9316066397443a2b07ce698e5fafc78 Mon Sep 17 00:00:00 2001 From: hackyminer Date: Fri, 11 Sep 2020 21:29:39 +0900 Subject: [PATCH 03/16] fix epoch zero bug --- libpoolprotocols/PoolManager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libpoolprotocols/PoolManager.cpp b/libpoolprotocols/PoolManager.cpp index 12a5584f6..d9192ad6f 100644 --- a/libpoolprotocols/PoolManager.cpp +++ b/libpoolprotocols/PoolManager.cpp @@ -182,7 +182,7 @@ void PoolManager::setClientHandlers() // If epoch is valued in workpackage take it if (wp.epoch == -1) { - if (m_currentWp.block > 0) + if (m_currentWp.block >= 0) m_currentWp.epoch = m_currentWp.block / 30000; else m_currentWp.epoch = ethash::find_epoch_number( From 861087f7f4ec75405438f35923427c54c56e5ab3 Mon Sep 17 00:00:00 2001 From: hackyminer Date: Sat, 12 Sep 2020 14:12:14 +0900 Subject: [PATCH 04/16] MESA define cl_amd_media_ops but no amd_bitalign() defined. * define/restore OPENCL_PLATFORM_* * define amd_bitalign() for MESA/Clover See also https://github.com/openwall/john/issues/3454#issuecomment-436899959 --- libethash-cl/kernels/cl/ethash.cl | 22 +++++++++++++++++++++- 1 file changed, 21 insertions(+), 1 deletion(-) diff --git a/libethash-cl/kernels/cl/ethash.cl b/libethash-cl/kernels/cl/ethash.cl index ce4586b67..2cbd6a32e 100644 --- a/libethash-cl/kernels/cl/ethash.cl +++ b/libethash-cl/kernels/cl/ethash.cl @@ -15,7 +15,11 @@ // You should have received a copy of the GNU General Public License // along with Gateless Gate Sharp. If not, see . - +#define OPENCL_PLATFORM_UNKNOWN 0 +#define OPENCL_PLATFORM_AMD 1 +#define OPENCL_PLATFORM_CLOVER 2 +#define OPENCL_PLATFORM_NVIDIA 3 +#define OPENCL_PLATFORM_INTEL 4 #if (defined(__Tahiti__) || defined(__Pitcairn__) || defined(__Capeverde__) || defined(__Oland__) || defined(__Hainan__)) #define LEGACY @@ -26,6 +30,22 @@ #endif #if defined(cl_amd_media_ops) +#if PLATFORM == OPENCL_PLATFORM_CLOVER +/* + * MESA define cl_amd_media_ops but no amd_bitalign() defined. + * https://github.com/openwall/john/issues/3454#issuecomment-436899959 + */ +uint2 amd_bitalign(uint2 src0, uint2 src1, uint2 src2) +{ + uint2 dst; + __asm("v_alignbit_b32 %0, %2, %3, %4\n" + "v_alignbit_b32 %1, %5, %6, %7" + : "=v" (dst.x), "=v" (dst.y) + : "v" (src0.x), "v" (src1.x), "v" (src2.x), + "v" (src0.y), "v" (src1.y), "v" (src2.y)); + return dst; +} +#endif #pragma OPENCL EXTENSION cl_amd_media_ops : enable #elif defined(cl_nv_pragma_unroll) uint amd_bitalign(uint src0, uint src1, uint src2) From 6a8ed69feeec2de587687eccc4e2e188cbfabd4f Mon Sep 17 00:00:00 2001 From: hackyminer Date: Thu, 17 Sep 2020 15:57:49 +0900 Subject: [PATCH 05/16] set PLATFORM as intended --- libethash-cl/CLMiner.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libethash-cl/CLMiner.cpp b/libethash-cl/CLMiner.cpp index 38a2b5ce4..62d1dc231 100644 --- a/libethash-cl/CLMiner.cpp +++ b/libethash-cl/CLMiner.cpp @@ -768,7 +768,7 @@ bool CLMiner::initEpoch_internal() addDefinition(code, "WORKSIZE", m_settings.localWorkSize); addDefinition(code, "ACCESSES", 64); addDefinition(code, "MAX_OUTPUTS", c_maxSearchResults); - addDefinition(code, "PLATFORM", m_deviceDescriptor.clPlatformId); + addDefinition(code, "PLATFORM", static_cast(m_deviceDescriptor.clPlatformType)); addDefinition(code, "COMPUTE", computeCapability); if (m_deviceDescriptor.clPlatformType == ClPlatformTypeEnum::Clover) From e7369aafc18c271130b716961c0610ce790e0309 Mon Sep 17 00:00:00 2001 From: Joe Clapis Date: Sun, 20 Sep 2020 13:22:04 -0400 Subject: [PATCH 06/16] Added rejected and failed share counts to the API's HTTP site --- libapicore/ApiServer.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/libapicore/ApiServer.cpp b/libapicore/ApiServer.cpp index 2e6934b78..a79b253ee 100644 --- a/libapicore/ApiServer.cpp +++ b/libapicore/ApiServer.cpp @@ -1169,7 +1169,11 @@ std::string ApiConnection::getHttpMinerStatDetail() _ret << "" << dev::getFormattedHashes(hashrate) << ""; - _ret << "" << device["mining"]["shares"][0].asString() << ""; + + string solString = "A" + device["mining"]["shares"][0].asString() + + ":R" + device["mining"]["shares"][1].asString() + + ":F" + device["mining"]["shares"][2].asString(); + _ret << "" << solString << ""; _ret << "" << device["hardware"]["sensors"][0].asString() << ""; _ret << "" << device["hardware"]["sensors"][1].asString() << ""; _ret << "" << device["hardware"]["sensors"][2].asString() << ""; From 34f891b3503973d2c7412dc0adf9f778d06736af Mon Sep 17 00:00:00 2001 From: Joe Clapis Date: Sun, 20 Sep 2020 13:51:14 -0400 Subject: [PATCH 07/16] Cleaned up the reported power numbers in the API's HTTP site --- libapicore/ApiServer.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/libapicore/ApiServer.cpp b/libapicore/ApiServer.cpp index a79b253ee..d23df8dfa 100644 --- a/libapicore/ApiServer.cpp +++ b/libapicore/ApiServer.cpp @@ -1176,7 +1176,10 @@ std::string ApiConnection::getHttpMinerStatDetail() _ret << "" << solString << ""; _ret << "" << device["hardware"]["sensors"][0].asString() << ""; _ret << "" << device["hardware"]["sensors"][1].asString() << ""; - _ret << "" << device["hardware"]["sensors"][2].asString() << ""; + + stringstream powerStream; // Round the power to 2 decimal places to remove floating point garbage + powerStream << fixed << setprecision(2) << device["hardware"]["sensors"][2].asString(); + _ret << "" << powerStream.str() << ""; _ret << ""; // Close row } From 29fcb8bb9b38038c3751aae38be821d9a64e267f Mon Sep 17 00:00:00 2001 From: Joe Clapis Date: Sun, 20 Sep 2020 14:01:50 -0400 Subject: [PATCH 08/16] Cleaned the reported power in the CLI display when HWMON is enabled --- libethcore/Miner.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libethcore/Miner.h b/libethcore/Miner.h index 082f203b3..2f3b98f21 100644 --- a/libethcore/Miner.h +++ b/libethcore/Miner.h @@ -147,7 +147,7 @@ struct HwSensorsType { string _ret = to_string(tempC) + "C " + to_string(fanP) + "%"; if (powerW) - _ret.append(boost::str(boost::format("%f") % powerW)); + _ret.append(" " + boost::str(boost::format("%0.2f") % powerW) + "W"); return _ret; }; }; From 433ec1b10f5f6a9e32c43a8a8480dd47e2edb840 Mon Sep 17 00:00:00 2001 From: Joe Clapis Date: Sun, 20 Sep 2020 14:07:11 -0400 Subject: [PATCH 09/16] Fixed API HTTP power report so it properly truncates to 2 decimal places --- libapicore/ApiServer.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libapicore/ApiServer.cpp b/libapicore/ApiServer.cpp index d23df8dfa..473e8a079 100644 --- a/libapicore/ApiServer.cpp +++ b/libapicore/ApiServer.cpp @@ -1178,7 +1178,7 @@ std::string ApiConnection::getHttpMinerStatDetail() _ret << "" << device["hardware"]["sensors"][1].asString() << ""; stringstream powerStream; // Round the power to 2 decimal places to remove floating point garbage - powerStream << fixed << setprecision(2) << device["hardware"]["sensors"][2].asString(); + powerStream << fixed << setprecision(2) << device["hardware"]["sensors"][2].asDouble(); _ret << "" << powerStream.str() << ""; _ret << ""; // Close row From c23dae17c65faccfeef72c9c8fadeba91c6222f2 Mon Sep 17 00:00:00 2001 From: Benjamin Kurczyk Date: Tue, 29 Sep 2020 08:53:23 +0200 Subject: [PATCH 10/16] Update POOL_EXAMPLES_ETH.md Small typo changes --- docs/POOL_EXAMPLES_ETH.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/POOL_EXAMPLES_ETH.md b/docs/POOL_EXAMPLES_ETH.md index 7c58615eb..d2a08a55a 100644 --- a/docs/POOL_EXAMPLES_ETH.md +++ b/docs/POOL_EXAMPLES_ETH.md @@ -67,7 +67,7 @@ The above samples produce the very same result. -P stratum://account%%2e1234.Worker:password@eu1.ethermine.org:4444 ``` -## Secure socket comunications for stratum only +## Secure socket communications for stratum only Ethminer supports secure socket communications (where pool implements and offers it) to avoid the risk of a [man-in-the-middle attack](https://en.wikipedia.org/wiki/Man-in-the-middle_attack) To enable it simply replace tcp with either: @@ -104,7 +104,7 @@ Here you can find a collection of samples to connect to most commonly used ethas * Stratum connection is **always to be preferred** over **getwork** when pool offers it due to its better network latency. * If possible the samples use a protocol which supports reporting of hashrate (`--report-hashrate`) if pool supports this. -**Check for updates in the pool connection settings visiting the pools homepage.** +**Check for updates in the pool connection settings visiting the pool's homepage.** ## Variables From d6d4197381815281c1205bb5955464a545120879 Mon Sep 17 00:00:00 2001 From: Menshykov Date: Fri, 9 Oct 2020 13:57:09 +0300 Subject: [PATCH 11/16] Update POOL_EXAMPLES_ETH.md --- docs/POOL_EXAMPLES_ETH.md | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/docs/POOL_EXAMPLES_ETH.md b/docs/POOL_EXAMPLES_ETH.md index d2a08a55a..ed56446e4 100644 --- a/docs/POOL_EXAMPLES_ETH.md +++ b/docs/POOL_EXAMPLES_ETH.md @@ -252,6 +252,11 @@ HINTS: ### nanopool.org +Notice ⚠ + +* Use "%40" for the @-sign in your email address +* Use "\%2e" for the .-sign on Linux in ETH_WALLET.WORKERNAME + With email: ``` @@ -272,9 +277,6 @@ Without email: -P stratum1+tcp://ETH_WALLET.WORKERNAME@eth-us-west1.nanopool.org:9999 ``` -HINTS: - -* Use "%40" for the @-sign in your email address ### nicehash.com From 752100a2ef19095387a24230e257d45f0dbc1c53 Mon Sep 17 00:00:00 2001 From: Marius van der Wijden Date: Tue, 13 Oct 2020 17:47:44 +0200 Subject: [PATCH 12/16] Readme.md: removed my ETH address Haven't done any work for ethminer in a long time so it is not fair that my address is still on the board. Thank you to everyone for the great times! --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 7ae2fabb2..5e239a13c 100644 --- a/README.md +++ b/README.md @@ -92,7 +92,7 @@ Ordered alphabetically. [Contributors statistics since 2015-08-20]. | EoD | [@EoD](https://github.com/EoD) | | | Genoil | [@Genoil](https://github.com/Genoil) | | | goobur | [@goobur](https://github.com/goobur) | | -| Marius van der Wijden | [@MariusVanDerWijden](https://github.com/MariusVanDerWijden) | ETH: 0x57d22b967c9dc64e5577f37edf1514c2d8985099 | +| Marius van der Wijden | [@MariusVanDerWijden](https://github.com/MariusVanDerWijden) | | | Paweł Bylica | [@chfast](https://github.com/chfast) | | | Philipp Andreas | [@smurfy](https://github.com/smurfy) | | | Stefan Oberhumer | [@StefanOberhumer](https://github.com/StefanOberhumer) | | From 8f48e43f9232655340badc626863e88cbc960b73 Mon Sep 17 00:00:00 2001 From: Alois Klink Date: Mon, 23 Nov 2020 17:29:23 +0000 Subject: [PATCH 13/16] Support Ampere graphic cards using Cuda 11.0/11.1 Ampere graphic cards require compute capability 8.0 and 8.6. Without this, trying to run this on my RTX 3090 fails. See https://arnon.dk/matching-sm-architectures-arch-and-gencode-for-various-nvidia-cards/ --- libethash-cuda/CMakeLists.txt | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/libethash-cuda/CMakeLists.txt b/libethash-cuda/CMakeLists.txt index 72eb897e6..2c27ad2ed 100644 --- a/libethash-cuda/CMakeLists.txt +++ b/libethash-cuda/CMakeLists.txt @@ -28,6 +28,14 @@ else() if(NOT CUDA_VERSION VERSION_LESS 10.0) list(APPEND CUDA_NVCC_FLAGS "-gencode arch=compute_75,code=sm_75") endif() + if(NOT CUDA_VERSION VERSION_LESS 11.0) + # NVIDIA A100 and NVIDIA DGX-A100 + list(APPEND CUDA_NVCC_FLAGS "-gencode arch=compute_80,code=sm_80") + endif() + if(NOT CUDA_VERSION VERSION_LESS 11.1) + # Tesla GA10x cards, RTX Ampere – RTX 3080/3090, RTX A6000, RTX A40 + list(APPEND CUDA_NVCC_FLAGS "-gencode arch=compute_86,code=sm_86") + endif() endif() file(GLOB sources "*.cpp" "*.cu") From c41fcefd4e55ba384356514a819114c661f38295 Mon Sep 17 00:00:00 2001 From: cryptomine Date: Fri, 8 Jan 2021 07:32:44 +0000 Subject: [PATCH 14/16] Update README.md Updated DAG size and GPU minimum memory (was 2GB, now it's 4GB as of epoch 382 on 2020-12-15) --- README.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index 5e239a13c..bf50c25cb 100644 --- a/README.md +++ b/README.md @@ -131,9 +131,9 @@ Because of the GDDR5X memory, which can't be fully utilized for ETH mining (yet) Only GCN 1.0 GPUs (78x0, 79x0, 270, 280), but in a different way. You'll see that on each new epoch (30K blocks), the hashrate will go down a little bit. -### Can I still mine ETH with my 2GB GPU? +### Can I still mine ETH with my 4GB GPU? -Not really, your VRAM must be above the DAG size (Currently about 2.15 GB.) to get best performance. Without it severe hash loss will occur. +Not really, your VRAM must be above the DAG size (Currently about 4.023 GB.) to get best performance. Without it severe hash loss will occur. ### What are the optimal launch parameters? From 608f4de6c4e995635cdfde80aed7435f3b1d7677 Mon Sep 17 00:00:00 2001 From: Jean Cyr Date: Sun, 10 Jan 2021 22:22:41 -0500 Subject: [PATCH 15/16] Fix fastexit not working on some AMD GPUs --- libethash-cl/kernels/cl/ethash.cl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libethash-cl/kernels/cl/ethash.cl b/libethash-cl/kernels/cl/ethash.cl index 9f0d6038d..d9614e9cb 100644 --- a/libethash-cl/kernels/cl/ethash.cl +++ b/libethash-cl/kernels/cl/ethash.cl @@ -266,7 +266,7 @@ struct SearchResults { __attribute__((reqd_work_group_size(WORKSIZE, 1, 1))) __kernel void search( - __global struct SearchResults* restrict g_output, + __global volatile struct SearchResults* restrict g_output, __constant uint2 const* g_header, __global ulong8 const* _g_dag0, __global ulong8 const* _g_dag1, From 9cb460b06f9da928636faa5f53dc78820a428f9b Mon Sep 17 00:00:00 2001 From: root Date: Mon, 8 Feb 2021 10:17:07 +0000 Subject: [PATCH 16/16] Avoiding buffer overflow in SHA3 --- libethash-cuda/ethash_cuda_miner_kernel.cu | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/libethash-cuda/ethash_cuda_miner_kernel.cu b/libethash-cuda/ethash_cuda_miner_kernel.cu index 1561d7939..4da35d7a8 100644 --- a/libethash-cuda/ethash_cuda_miner_kernel.cu +++ b/libethash-cuda/ethash_cuda_miner_kernel.cu @@ -52,11 +52,13 @@ __global__ void ethash_calculate_dag_item(uint32_t start) uint32_t const node_index = start + blockIdx.x * blockDim.x + threadIdx.x; if (((node_index >> 1) & (~1)) >= d_dag_size) return; - - hash128_t dag_node; + union { + hash128_t dag_node; + uint2 dag_node_mem[25]; + }; copy(dag_node.uint4s, d_light[node_index % d_light_size].uint4s, 4); dag_node.words[0] ^= node_index; - SHA3_512(dag_node.uint2s); + SHA3_512(dag_node_mem); const int thread_id = threadIdx.x & 3; @@ -78,7 +80,7 @@ __global__ void ethash_calculate_dag_item(uint32_t start) } } } - SHA3_512(dag_node.uint2s); + SHA3_512(dag_node_mem); hash64_t* dag_nodes = (hash64_t*)d_dag; copy(dag_nodes[node_index].uint4s, dag_node.uint4s, 4); }