diff --git a/JHA/jackpotcoin.cu b/JHA/jackpotcoin.cu index f13a7e8bc6..78e5bd02e7 100644 --- a/JHA/jackpotcoin.cu +++ b/JHA/jackpotcoin.cu @@ -30,8 +30,8 @@ extern void jackpot_compactTest_cpu_hash_64(int thr_id, uint32_t threads, uint32 extern uint32_t cuda_check_hash_branch(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_inputHash, int order); -// Original jackpothash Funktion aus einem miner Quelltext -extern "C" unsigned int jackpothash(void *state, const void *input) +// CPU HASH JHA v8 +extern "C" void jackpothash(void *state, const void *input) { uint32_t hash[16]; unsigned int rnd; @@ -71,8 +71,6 @@ extern "C" unsigned int jackpothash(void *state, const void *input) } } memcpy(state, hash, 32); - - return rnd; } static bool init[MAX_GPUS] = { 0 }; diff --git a/Makefile.am b/Makefile.am index 08971d0667..307d86e42e 100644 --- a/Makefile.am +++ b/Makefile.am @@ -45,7 +45,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \ crypto/cryptolight.cu crypto/cryptolight-core.cu crypto/cryptolight-cpu.cpp \ crypto/cryptonight.cu crypto/cryptonight-core.cu crypto/cryptonight-extra.cu \ crypto/cryptonight-cpu.cpp crypto/oaes_lib.cpp crypto/aesb.cpp crypto/cpu/c_keccak.c \ - JHA/jha.cu JHA/cuda_jha_keccak512.cu \ + JHA/jha.cu JHA/jackpotcoin.cu JHA/cuda_jha_keccak512.cu \ JHA/cuda_jha_compactionTest.cu cuda_checkhash.cu \ quark/cuda_jh512.cu quark/cuda_quark_blake512.cu quark/cuda_quark_groestl512.cu quark/cuda_skein512.cu \ quark/cuda_bmw512.cu quark/cuda_quark_keccak512.cu \ diff --git a/algos.h b/algos.h index df18131f53..e4dc4bb05a 100644 --- a/algos.h +++ b/algos.h @@ -22,6 +22,7 @@ enum sha_algos { ALGO_HEAVY, /* Heavycoin hash */ ALGO_HMQ1725, ALGO_KECCAK, + ALGO_JACKPOT, ALGO_JHA, ALGO_LBRY, ALGO_LUFFA, @@ -83,6 +84,7 @@ static const char *algo_names[] = { "heavy", "hmq1725", "keccak", + "jackpot", "jha", "lbry", "luffa", @@ -151,8 +153,8 @@ static inline int algo_to_int(char* arg) i = ALGO_LUFFA; else if (!strcasecmp("hmq17", arg)) i = ALGO_HMQ1725; - else if (!strcasecmp("jackpot", arg)) - i = ALGO_JHA; + //else if (!strcasecmp("jackpot", arg)) + // i = ALGO_JHA; else if (!strcasecmp("lyra2re", arg)) i = ALGO_LYRA2; else if (!strcasecmp("lyra2rev2", arg)) diff --git a/bench.cpp b/bench.cpp index 3e6822456f..b288089c3c 100644 --- a/bench.cpp +++ b/bench.cpp @@ -65,7 +65,7 @@ void algo_free_all(int thr_id) free_groestlcoin(thr_id); free_heavy(thr_id); free_hmq17(thr_id); - //free_jackpot(thr_id); + free_jackpot(thr_id); free_jha(thr_id); free_lbry(thr_id); free_luffa(thr_id); diff --git a/ccminer.cpp b/ccminer.cpp index 142fbfe7d9..690f38a0cc 100644 --- a/ccminer.cpp +++ b/ccminer.cpp @@ -240,7 +240,7 @@ Options:\n\ groestl Groestlcoin\n\ heavy Heavycoin\n\ hmq1725 Doubloons / Espers\n\ - jackpot Jackpot\n\ + jha JHA v8 (JackpotCoin)\n\ keccak Keccak-256 (Maxcoin)\n\ lbry LBRY Credits (Sha/Ripemd)\n\ luffa Joincoin\n\ @@ -1612,6 +1612,7 @@ static bool stratum_gen_work(struct stratum_ctx *sctx, struct work *work) switch (opt_algo) { case ALGO_HMQ1725: + case ALGO_JACKPOT: case ALGO_JHA: case ALGO_NEOSCRYPT: case ALGO_SCRYPT: @@ -2128,6 +2129,7 @@ static void *miner_thread(void *userdata) case ALGO_C11: case ALGO_DEEP: case ALGO_HEAVY: + case ALGO_JACKPOT: case ALGO_JHA: case ALGO_LYRA2v2: case ALGO_S3: @@ -2269,9 +2271,14 @@ static void *miner_thread(void *userdata) case ALGO_KECCAK: rc = scanhash_keccak256(thr_id, &work, max_nonce, &hashes_done); break; + + case ALGO_JACKPOT: + rc = scanhash_jackpot(thr_id, &work, max_nonce, &hashes_done); + break; case ALGO_JHA: rc = scanhash_jha(thr_id, &work, max_nonce, &hashes_done); break; + case ALGO_LBRY: rc = scanhash_lbry(thr_id, &work, max_nonce, &hashes_done); break; @@ -2426,7 +2433,7 @@ static void *miner_thread(void *userdata) /* hashrate factors for some algos */ double rate_factor = 1.0; switch (opt_algo) { - //case ALGO_JACKPOT: + case ALGO_JACKPOT: case ALGO_QUARK: // to stay comparable to other ccminer forks or pools rate_factor = 0.5; diff --git a/ccminer.vcxproj b/ccminer.vcxproj index c18d560fbb..5aad65e959 100644 --- a/ccminer.vcxproj +++ b/ccminer.vcxproj @@ -436,6 +436,7 @@ + -Xptxas "-abi=yes" %(AdditionalOptions) diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters index 4b3d828771..f37da2345f 100644 --- a/ccminer.vcxproj.filters +++ b/ccminer.vcxproj.filters @@ -577,6 +577,9 @@ Source Files\CUDA\JHA + + Source Files\CUDA\JHA + Source Files\CUDA diff --git a/miner.h b/miner.h index 9e30dac091..f720bc75f0 100644 --- a/miner.h +++ b/miner.h @@ -871,7 +871,7 @@ void fugue256_hash(unsigned char* output, const unsigned char* input, int len); void heavycoin_hash(unsigned char* output, const unsigned char* input, int len); void hmq17hash(void *output, const void *input); void keccak256_hash(void *state, const void *input); -unsigned int jackpothash(void *state, const void *input); +void jackpothash(void *state, const void *input); void groestlhash(void *state, const void *input); void jha_hash(void *output, const void *input); void lbry_hash(void *output, const void *input); diff --git a/quark/cuda_quark.h b/quark/cuda_quark.h index 18f37d5f95..fbb0c1da43 100644 --- a/quark/cuda_quark.h +++ b/quark/cuda_quark.h @@ -17,7 +17,7 @@ extern void quark_doublegroestl512_cpu_hash_64(int thr_id, uint32_t threads, uin extern void quark_groestl512_cpu_free(int thr_id); extern void quark_skein512_cpu_init(int thr_id, uint32_t threads); -extern void quark_skein512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); +extern void quark_skein512_cpu_hash_64(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); extern void quark_keccak512_cpu_init(int thr_id, uint32_t threads); extern void quark_keccak512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order); diff --git a/quark/cuda_skein512.cu b/quark/cuda_skein512.cu index bdb5d13e0e..7c4d99ab3a 100644 --- a/quark/cuda_skein512.cu +++ b/quark/cuda_skein512.cu @@ -463,7 +463,7 @@ __launch_bounds__(TPB52, 3) #else __launch_bounds__(TPB50, 5) #endif -void quark_skein512_gpu_hash_64(const uint32_t threads,uint64_t* __restrict__ g_hash, const uint32_t *const __restrict__ g_nonceVector) +void quark_skein512_gpu_hash_64(const uint32_t threads, const uint32_t startNonce, uint64_t* __restrict__ g_hash, const uint32_t *const __restrict__ g_nonceVector) { const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x); @@ -472,7 +472,7 @@ void quark_skein512_gpu_hash_64(const uint32_t threads,uint64_t* __restrict__ g_ // Skein uint2 p[8], h[9]; - const uint32_t hashPosition = (g_nonceVector == NULL) ? thread : g_nonceVector[thread]; + const uint32_t hashPosition = (g_nonceVector == NULL) ? thread : g_nonceVector[thread] - startNonce; uint64_t *Hash = &g_hash[hashPosition<<3]; @@ -756,7 +756,7 @@ void quark_skein512_gpu_hash_64(const uint32_t threads,uint64_t* __restrict__ g_ __host__ //void quark_skein512_cpu_hash_64(int thr_id,uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_hash) -void quark_skein512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) +void quark_skein512_cpu_hash_64(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_nonceVector, uint32_t *d_hash, int order) { uint32_t tpb = TPB52; int dev_id = device_map[thr_id]; @@ -764,7 +764,7 @@ void quark_skein512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNoun if (device_sm[dev_id] <= 500) tpb = TPB50; const dim3 grid((threads + tpb-1)/tpb); const dim3 block(tpb); - quark_skein512_gpu_hash_64 << > >(threads, (uint64_t*)d_hash, d_nonceVector); + quark_skein512_gpu_hash_64 <<>>(threads, startNonce, (uint64_t*)d_hash, d_nonceVector); }