Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Don't redefine and misuse ETHASH_MIX_BYTES #35

Merged
merged 1 commit into from
Jun 10, 2019
Merged

Don't redefine and misuse ETHASH_MIX_BYTES #35

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion libethash-cl/CLMiner.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
4 changes: 2 additions & 2 deletions libethash-cuda/CUDAMiner.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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));
Expand Down
2 changes: 1 addition & 1 deletion libethash/ethash.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
4 changes: 2 additions & 2 deletions libprogpow/ProgPow.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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";
Expand All @@ -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";
Expand Down