From c7055275d0a21a4bde84d53848b74ee49d372eb0 Mon Sep 17 00:00:00 2001 From: Solar Designer Date: Thu, 4 Apr 2019 23:31:56 +0200 Subject: [PATCH] Don't redefine and misuse ETHASH_MIX_BYTES --- libethash-cl/CLMiner.cpp | 2 +- libethash-cuda/CUDAMiner.cpp | 4 ++-- libethash/ethash.h | 2 +- libprogpow/ProgPow.cpp | 4 ++-- 4 files changed, 6 insertions(+), 6 deletions(-) diff --git a/libethash-cl/CLMiner.cpp b/libethash-cl/CLMiner.cpp index 2a591f845..05691e749 100644 --- a/libethash-cl/CLMiner.cpp +++ b/libethash-cl/CLMiner.cpp @@ -597,7 +597,7 @@ bool CLMiner::init(int epoch) m_globalWorkSize = ((m_globalWorkSize / m_workgroupSize) + 1) * m_workgroupSize; uint64_t dagBytes = ethash_get_datasize(light->light->block_number); - uint32_t dagElms = (unsigned)(dagBytes / ETHASH_MIX_BYTES); + uint32_t dagElms = (unsigned)(dagBytes / (PROGPOW_LANES * PROGPOW_DAG_LOADS * 4)); uint32_t lightWords = (unsigned)(light->data().size() / sizeof(node)); // patch source code diff --git a/libethash-cuda/CUDAMiner.cpp b/libethash-cuda/CUDAMiner.cpp index 91c1655ce..59e837bfb 100644 --- a/libethash-cuda/CUDAMiner.cpp +++ b/libethash-cuda/CUDAMiner.cpp @@ -124,7 +124,7 @@ void CUDAMiner::workLoop() if (old_period_seed != period_seed) { uint64_t dagBytes = ethash_get_datasize(w.height); - uint32_t dagElms = (unsigned)(dagBytes / ETHASH_MIX_BYTES); + uint32_t dagElms = (unsigned)(dagBytes / (PROGPOW_LANES * PROGPOW_DAG_LOADS * 4)); compileKernel(w.height, dagElms); } old_period_seed = period_seed; @@ -356,7 +356,7 @@ bool CUDAMiner::cuda_init( m_streams = new cudaStream_t[s_numStreams]; uint64_t dagBytes = ethash_get_datasize(_light->block_number); - uint32_t dagElms = (unsigned)(dagBytes / ETHASH_MIX_BYTES); + uint32_t dagElms = (unsigned)(dagBytes / (PROGPOW_LANES * PROGPOW_DAG_LOADS * 4)); uint32_t lightWords = (unsigned)(_lightBytes / sizeof(node)); CUDA_SAFE_CALL(cudaSetDevice(m_device_num)); diff --git a/libethash/ethash.h b/libethash/ethash.h index b56674885..657fc6386 100644 --- a/libethash/ethash.h +++ b/libethash/ethash.h @@ -32,7 +32,7 @@ #define ETHASH_CACHE_BYTES_INIT 1073741824U // 2**24 #define ETHASH_CACHE_BYTES_GROWTH 131072U // 2**17 #define ETHASH_EPOCH_LENGTH 30000U -#define ETHASH_MIX_BYTES 256 +#define ETHASH_MIX_BYTES 128 #define ETHASH_HASH_BYTES 64 #define ETHASH_DATASET_PARENTS 256 #define ETHASH_CACHE_ROUNDS 3 diff --git a/libprogpow/ProgPow.cpp b/libprogpow/ProgPow.cpp index d1ec3ec24..1907c5a30 100644 --- a/libprogpow/ProgPow.cpp +++ b/libprogpow/ProgPow.cpp @@ -86,7 +86,7 @@ std::string ProgPow::getKern(uint64_t block_number, kernel_t kern) if (kern == KERNEL_CUDA) { - ret << "typedef struct __align__(16) {uint32_t s[PROGPOW_DAG_LOADS];} dag_t;\n"; + ret << "typedef struct __align__(PROGPOW_DAG_LOADS * 4) {uint32_t s[PROGPOW_DAG_LOADS];} dag_t;\n"; ret << "\n"; ret << "// Inner loop for prog_seed " << prog_seed << "\n"; ret << "__device__ __forceinline__ void progPowLoop(const uint32_t loop,\n"; @@ -97,7 +97,7 @@ std::string ProgPow::getKern(uint64_t block_number, kernel_t kern) } else { - ret << "typedef struct __attribute__ ((aligned (16))) {uint32_t s[PROGPOW_DAG_LOADS];} dag_t;\n"; + ret << "typedef struct __attribute__ ((aligned (PROGPOW_DAG_LOADS * 4))) {uint32_t s[PROGPOW_DAG_LOADS];} dag_t;\n"; ret << "\n"; ret << "// Inner loop for prog_seed " << prog_seed << "\n"; ret << "void progPowLoop(const uint32_t loop,\n";