Browse Source

change defaults to handle cuda 9+, disable heavy and SM 2.x

Heavy is the only algo using thrust which is generally broken on new cuda releases

mjollnir dropped too... never seen this coin anyway...
pull/2/head
Tanguy Pruvot 7 years ago
parent
commit
11a512f2a7
  1. 7
      Makefile.am
  2. 4
      README.txt
  3. 3
      bench.cpp
  4. 19
      ccminer.cpp
  5. 14
      ccminer.vcxproj
  6. 89
      heavy/heavy.cu

7
Makefile.am

@ -116,7 +116,6 @@ nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\" @@ -116,7 +116,6 @@ nvcc_ARCH = -gencode=arch=compute_50,code=\"sm_50,compute_50\"
nvcc_ARCH += -gencode=arch=compute_52,code=\"sm_52,compute_52\"
#nvcc_ARCH += -gencode=arch=compute_35,code=\"sm_35,compute_35\"
#nvcc_ARCH += -gencode=arch=compute_30,code=\"sm_30,compute_30\"
#nvcc_ARCH += -gencode=arch=compute_20,code=\"sm_21,compute_20\"
nvcc_FLAGS = $(nvcc_ARCH) @CUDA_INCLUDES@ -I. @CUDA_CFLAGS@
nvcc_FLAGS += $(JANSSON_INCLUDES) --ptxas-options="-v"
@ -176,15 +175,15 @@ JHA/cuda_jha_compactionTest.o: JHA/cuda_jha_compactionTest.cu @@ -176,15 +175,15 @@ JHA/cuda_jha_compactionTest.o: JHA/cuda_jha_compactionTest.cu
# This object does not use cuda device code but call the different kernels (autotune)
scrypt/salsa_kernel.o: scrypt/salsa_kernel.cu
$(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_20,code=\"sm_21,compute_20\" -o $@ -c $<
$(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $<
# These kernels are for older devices (SM)
scrypt/test_kernel.o: scrypt/test_kernel.cu
$(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_20,code=\"sm_20,compute_20\" -o $@ -c $<
$(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $<
scrypt/fermi_kernel.o: scrypt/fermi_kernel.cu
$(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_20,code=\"sm_21,compute_20\" -o $@ -c $<
$(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $<
scrypt/kepler_kernel.o: scrypt/kepler_kernel.cu
$(NVCC) $(JANSSON_INCLUDES) -I. @CUDA_INCLUDES@ @CUDA_CFLAGS@ -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o $@ -c $<

4
README.txt

@ -89,7 +89,6 @@ its command line interface and options. @@ -89,7 +89,6 @@ its command line interface and options.
fresh use to mine Freshcoin
fugue256 use to mine Fuguecoin
groestl use to mine Groestlcoin
heavy use to mine Heavycoin
hsr use to mine Hshare
jackpot use to mine Sweepcoin
keccak use to mine Maxcoin
@ -99,7 +98,6 @@ its command line interface and options. @@ -99,7 +98,6 @@ its command line interface and options.
lyra2 use to mine CryptoCoin
lyra2v2 use to mine Vertcoin
lyra2z use to mine Zerocoin (XZC)
mjollnir use to mine Mjollnircoin
myr-gr use to mine Myriad-Groest
neoscrypt use to mine FeatherCoin
nist5 use to mine TalkCoin
@ -142,8 +140,6 @@ its command line interface and options. @@ -142,8 +140,6 @@ its command line interface and options.
--cuda-schedule Set device threads scheduling mode (default: auto)
-f, --diff-factor Divide difficulty by this factor (default 1.0)
-m, --diff-multiplier Multiply difficulty by this value (default 1.0)
--vote=VOTE block reward vote (for HeavyCoin)
--trust-pool trust the max block reward vote (maxvote) sent by the pool
-o, --url=URL URL of mining server
-O, --userpass=U:P username:password pair for mining server
-u, --user=USERNAME username for mining server

3
bench.cpp

@ -64,7 +64,9 @@ void algo_free_all(int thr_id) @@ -64,7 +64,9 @@ void algo_free_all(int thr_id)
free_fresh(thr_id);
free_fugue256(thr_id);
free_groestlcoin(thr_id);
#ifdef WITH_HEAVY_ALGO
free_heavy(thr_id);
#endif
free_hmq17(thr_id);
free_hsr(thr_id);
free_jackpot(thr_id);
@ -125,6 +127,7 @@ bool bench_algo_switch_next(int thr_id) @@ -125,6 +127,7 @@ bool bench_algo_switch_next(int thr_id)
// skip some duplicated algos
if (algo == ALGO_C11) algo++; // same as x11
if (algo == ALGO_DMD_GR) algo++; // same as groestl
if (algo == ALGO_HEAVY) algo++; // dead
if (algo == ALGO_MJOLLNIR) algo++; // same as heavy
if (algo == ALGO_KECCAKC) algo++; // same as keccak
if (algo == ALGO_WHIRLCOIN) algo++; // same as whirlpool

19
ccminer.cpp

@ -251,9 +251,11 @@ Options:\n\ @@ -251,9 +251,11 @@ Options:\n\
dmd-gr Diamond-Groestl\n\
fresh Freshcoin (shavite 80)\n\
fugue256 Fuguecoin\n\
groestl Groestlcoin\n\
heavy Heavycoin\n\
hmq1725 Doubloons / Espers\n\
groestl Groestlcoin\n"
#ifdef WITH_HEAVY_ALGO
" heavy Heavycoin\n"
#endif
" hmq1725 Doubloons / Espers\n\
jackpot JHA v8\n\
keccak Deprecated Keccak-256\n\
keccakc Keccak-256 (CreativeCoin)\n\
@ -262,7 +264,6 @@ Options:\n\ @@ -262,7 +264,6 @@ Options:\n\
lyra2 CryptoCoin\n\
lyra2v2 VertCoin\n\
lyra2z ZeroCoin (3rd impl)\n\
mjollnir Mjollnircoin\n\
myr-gr Myriad-Groestl\n\
neoscrypt FeatherCoin, Phoenix, UFO...\n\
nist5 NIST5 (TalkCoin)\n\
@ -304,8 +305,6 @@ Options:\n\ @@ -304,8 +305,6 @@ Options:\n\
--cuda-schedule Set device threads scheduling mode (default: auto)\n\
-f, --diff-factor Divide difficulty by this factor (default 1.0) \n\
-m, --diff-multiplier Multiply difficulty by this value (default 1.0) \n\
--vote=VOTE vote (for HeavyCoin)\n\
--trust-pool trust the max block reward vote (maxvote) sent by the pool\n\
-o, --url=URL URL of mining server\n\
-O, --userpass=U:P username:password pair for mining server\n\
-u, --user=USERNAME username for mining server\n\
@ -1555,10 +1554,12 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) @@ -1555,10 +1554,12 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
case ALGO_SIA:
// getwork over stratum, no merkle to generate
break;
#ifdef WITH_HEAVY_ALGO
case ALGO_HEAVY:
case ALGO_MJOLLNIR:
heavycoin_hash(merkle_root, sctx->job.coinbase, (int)sctx->job.coinbase_size);
break;
#endif
case ALGO_FUGUE256:
case ALGO_GROESTL:
case ALGO_KECCAK:
@ -1573,9 +1574,11 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) @@ -1573,9 +1574,11 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work)
for (i = 0; i < sctx->job.merkle_count; i++) {
memcpy(merkle_root + 32, sctx->job.merkle[i], 32);
#ifdef WITH_HEAVY_ALGO
if (opt_algo == ALGO_HEAVY || opt_algo == ALGO_MJOLLNIR)
heavycoin_hash(merkle_root, merkle_root, 64);
else
#endif
sha256d(merkle_root, merkle_root, 64);
}
@ -2368,14 +2371,14 @@ static void *miner_thread(void *userdata) @@ -2368,14 +2371,14 @@ static void *miner_thread(void *userdata)
case ALGO_HSR:
rc = scanhash_hsr(thr_id, &work, max_nonce, &hashes_done);
break;
#ifdef WITH_HEAVY_ALGO
case ALGO_HEAVY:
rc = scanhash_heavy(thr_id, &work, max_nonce, &hashes_done, work.maxvote, HEAVYCOIN_BLKHDR_SZ);
break;
case ALGO_MJOLLNIR:
rc = scanhash_heavy(thr_id, &work, max_nonce, &hashes_done, 0, MNR_BLKHDR_SZ);
break;
#endif
case ALGO_KECCAK:
case ALGO_KECCAKC:
rc = scanhash_keccak256(thr_id, &work, max_nonce, &hashes_done);

14
ccminer.vcxproj

@ -39,7 +39,7 @@ @@ -39,7 +39,7 @@
</PropertyGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
<ImportGroup Label="ExtensionSettings">
<Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 8.0.props" />
<Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 9.0.props" />
</ImportGroup>
<ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
<Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
@ -155,7 +155,7 @@ @@ -155,7 +155,7 @@
<MaxRegCount>80</MaxRegCount>
<PtxAsOptionV>true</PtxAsOptionV>
<Keep>true</Keep>
<CodeGeneration>compute_50,sm_50;compute_52,sm_52;compute_30,sm_30;compute_20,sm_21</CodeGeneration>
<CodeGeneration>compute_50,sm_50;compute_52,sm_52;compute_30,sm_30</CodeGeneration>
<AdditionalOptions>--ptxas-options="-O2" --Wno-deprecated-gpu-targets %(AdditionalOptions)</AdditionalOptions>
<Optimization>O2</Optimization>
</CudaCompile>
@ -198,7 +198,7 @@ @@ -198,7 +198,7 @@
<MaxRegCount>80</MaxRegCount>
<PtxAsOptionV>true</PtxAsOptionV>
<Keep>true</Keep>
<CodeGeneration>compute_52,sm_52;compute_50,sm_50;compute_35,sm_35;compute_30,sm_30;compute_20,sm_21</CodeGeneration>
<CodeGeneration>compute_52,sm_52;compute_50,sm_50;compute_35,sm_35;compute_30,sm_30</CodeGeneration>
<Include>$(NVTOOLSEXT_PATH)\include</Include>
<Optimization>O3</Optimization>
<TargetMachinePlatform>64</TargetMachinePlatform>
@ -410,10 +410,10 @@ @@ -410,10 +410,10 @@
<CudaCompile Include="scrypt\keccak.cu" />
<CudaCompile Include="scrypt\sha256.cu" />
<CudaCompile Include="scrypt\salsa_kernel.cu">
<CodeGeneration>compute_20,sm_21</CodeGeneration>
<CodeGeneration>compute_30,sm_30</CodeGeneration>
</CudaCompile>
<CudaCompile Include="scrypt\fermi_kernel.cu">
<CodeGeneration>compute_20,sm_21</CodeGeneration>
<CodeGeneration>compute_30,sm_30</CodeGeneration>
</CudaCompile>
<CudaCompile Include="scrypt\kepler_kernel.cu">
<CodeGeneration>compute_30,sm_30</CodeGeneration>
@ -425,7 +425,7 @@ @@ -425,7 +425,7 @@
<CodeGeneration>compute_35,sm_35;compute_50,sm_50;compute_52,sm_52</CodeGeneration>
</CudaCompile>
<CudaCompile Include="scrypt\test_kernel.cu">
<CodeGeneration>compute_20,sm_21</CodeGeneration>
<CodeGeneration>compute_30,sm_30</CodeGeneration>
</CudaCompile>
<CudaCompile Include="scrypt\titan_kernel.cu">
<CodeGeneration>compute_35,sm_35;compute_50,sm_50</CodeGeneration>
@ -606,7 +606,7 @@ @@ -606,7 +606,7 @@
</ItemGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
<ImportGroup Label="ExtensionTargets">
<Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 8.0.targets" />
<Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 9.0.targets" />
</ImportGroup>
<!-- Copy the required dlls -->
<Target Name="AfterBuild">

89
heavy/heavy.cu

@ -3,6 +3,21 @@ @@ -3,6 +3,21 @@
#include <cuda.h>
#include <map>
#ifndef WITH_HEAVY_ALGO
#include <unistd.h>
#include "miner.h"
// nonce array also used in other algos
uint32_t *heavy_nonceVector[MAX_GPUS];
int scanhash_heavy(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done, uint32_t maxvote, int blocklen)
{
applog(LOG_ERR, "heavy algo not included in this build!");
sleep(3);
return -1;
}
void free_heavy(int thr_id) {}
#else
// include thrust if possible
#if defined(__GNUC__) && __GNUC__ == 5 && __GNUC_MINOR__ >= 2 && CUDA_VERSION < 7000
#warning "Heavy: incompatible GCC version!"
@ -17,16 +32,11 @@ @@ -17,16 +32,11 @@
#endif
#include "miner.h"
extern "C" {
#include "sph/sph_keccak.h"
#include "sph/sph_blake.h"
#include "sph/sph_groestl.h"
}
#include "hefty1.h"
#include "heavy/heavy.h"
#include "cuda_helper.h"
// nonce array also used in other algos
uint32_t *heavy_nonceVector[MAX_GPUS];
extern uint32_t *d_hash2output[MAX_GPUS];
extern uint32_t *d_hash3output[MAX_GPUS];
extern uint32_t *d_hash4output[MAX_GPUS];
@ -35,35 +45,8 @@ extern uint32_t *d_hash5output[MAX_GPUS]; @@ -35,35 +45,8 @@ extern uint32_t *d_hash5output[MAX_GPUS];
#define HEAVYCOIN_BLKHDR_SZ 84
#define MNR_BLKHDR_SZ 80
// nonce-array für die threads
uint32_t *heavy_nonceVector[MAX_GPUS];
extern uint32_t *heavy_heftyHashes[MAX_GPUS];
/* Combines top 64-bits from each hash into a single hash */
static void combine_hashes(uint32_t *out, const uint32_t *hash1, const uint32_t *hash2, const uint32_t *hash3, const uint32_t *hash4)
{
const uint32_t *hash[4] = { hash1, hash2, hash3, hash4 };
int bits;
unsigned int i;
uint32_t mask;
unsigned int k;
/* Transpose first 64 bits of each hash into out */
memset(out, 0, 32);
bits = 0;
for (i = 7; i >= 6; i--) {
for (mask = 0x80000000; mask; mask >>= 1) {
for (k = 0; k < 4; k++) {
out[(255 - bits)/32] <<= 1;
if ((hash[k][i] & mask) != 0)
out[(255 - bits)/32] |= 1;
bits++;
}
}
}
}
#ifdef _MSC_VER
#include <intrin.h>
static uint32_t __inline bitsset( uint32_t x )
@ -349,6 +332,42 @@ extern "C" void free_heavy(int thr_id) @@ -349,6 +332,42 @@ extern "C" void free_heavy(int thr_id)
cudaDeviceSynchronize();
}
#endif
extern "C" {
#include "sph/sph_keccak.h"
#include "sph/sph_blake.h"
#include "sph/sph_groestl.h"
}
#include "hefty1.h"
#include "heavy/heavy.h"
/* Combines top 64-bits from each hash into a single hash */
__host__
static void combine_hashes(uint32_t *out, const uint32_t *hash1, const uint32_t *hash2, const uint32_t *hash3, const uint32_t *hash4)
{
const uint32_t *hash[4] = { hash1, hash2, hash3, hash4 };
int bits;
unsigned int i;
uint32_t mask;
unsigned int k;
/* Transpose first 64 bits of each hash into out */
memset(out, 0, 32);
bits = 0;
for (i = 7; i >= 6; i--) {
for (mask = 0x80000000; mask; mask >>= 1) {
for (k = 0; k < 4; k++) {
out[(255 - bits) / 32] <<= 1;
if ((hash[k][i] & mask) != 0)
out[(255 - bits) / 32] |= 1;
bits++;
}
}
}
}
// CPU hash function
__host__
void heavycoin_hash(uchar* output, const uchar* input, int len)
{

Loading…
Cancel
Save