fzorb/REDACTED-rig
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Disable OPENCL double threading as default option as it seems gives lower hashrate.

Browse source
This commit is contained in:
Haifa Bogdan Adnan 2019-09-08 01:11:17 +03:00
parent e6fa35db42
commit f787b9f2cc
2 changed files with 4 additions and 6 deletions
Download patch file Download diff file Expand all files Collapse all files

2
CMakeLists.txt
View file

@ -8,7 +8,7 @@ option(WITH_TLS "Enable OpenSSL support" ON)
option(WITH_EMBEDDED_CONFIG "Enable internal embedded JSON config" OFF) option(WITH_EMBEDDED_CONFIG "Enable internal embedded JSON config" OFF)
option(WITH_CUDA "Enable CUDA support" ON) option(WITH_CUDA "Enable CUDA support" ON)
option(WITH_OPENCL "Enable OpenCL support" ON) option(WITH_OPENCL "Enable OpenCL support" ON)
option(WITH_OPENCL_DOUBLE_THREADS "Enable dual threads for OpenCL jobs" ON) option(WITH_OPENCL_DOUBLE_THREADS "Enable dual threads for OpenCL jobs" OFF)
option(WITH_CUDA_DOUBLE_THREADS "Enable dual threads for CUDA jobs" ON) option(WITH_CUDA_DOUBLE_THREADS "Enable dual threads for CUDA jobs" ON)
include (CheckIncludeFile) include (CheckIncludeFile)

8
src/crypto/argon2_hasher/hash/gpu/opencl/OpenCLKernel.cpp
View file

@ -862,7 +862,7 @@ __kernel void fill_blocks(__global ulong *chunk_0,
cur_idx = seg_idxs[0]; cur_idx = seg_idxs[0];
} }
ulong4 nextref = vload4(wave_id, memory + ref_idx * 2 * BLOCK_SIZE_ULONG); ref = vload4(wave_id, memory + ref_idx * 2 * BLOCK_SIZE_ULONG);
for (int i=0;idx < seg_length;i++, idx++) { for (int i=0;idx < seg_length;i++, idx++) {
next_block = memory + (cur_idx & 0x7FFFFFFF) * 2 * BLOCK_SIZE_ULONG; next_block = memory + (cur_idx & 0x7FFFFFFF) * 2 * BLOCK_SIZE_ULONG;
@ -870,7 +870,7 @@ __kernel void fill_blocks(__global ulong *chunk_0,
if(with_xor == 1) if(with_xor == 1)
next = vload4(wave_id, next_block); next = vload4(wave_id, next_block);
ref = nextref; tmp ^= ref;
if (idx < seg_length - 1) { if (idx < seg_length - 1) {
ref_idx = seg_refs[i + 1]; ref_idx = seg_refs[i + 1];
@ -882,11 +882,9 @@ __kernel void fill_blocks(__global ulong *chunk_0,
else else
cur_idx++; cur_idx++;
nextref = vload4(wave_id, memory + ref_idx * 2 * BLOCK_SIZE_ULONG); ref = vload4(wave_id, memory + ref_idx * 2 * BLOCK_SIZE_ULONG);
} }
tmp ^= ref;
vstore4(tmp, id, state); vstore4(tmp, id, state);
G1(state); G1(state);