v0.9.2

Whirlpool improvements for x16 (higher impact on x16s)
DumaxFr · Jun 24, 2018 · 52795ba · 52795ba
1 parent 010e019
commit 52795ba
Show file tree

Hide file tree

Showing 12 changed files with 484 additions and 57 deletions.
diff --git a/Algo512/cuda_b_whirlpool.cu b/Algo512/cuda_b_whirlpool.cu
diff --git a/Algo512/cuda_b_whirlpool.h b/Algo512/cuda_b_whirlpool.h
@@ -0,0 +1,17 @@
+#ifndef CUDA_BASE_WHIRLPOOL
+#define CUDA_BASE_WHIRLPOOL
+
+#include <stdint.h>
+
+void cuda_base_whirlpool_cpu_init();
+
+//void cuda_base_whirlpool_setBlock_80(void * pdata);
+//void cuda_base_whirlpool_cpu_hash_80(uint32_t threads, uint32_t startNounce, uint32_t * d_hash);
+
+void cuda_base_whirlpool_cpu_hash_64(const uint32_t threads, uint32_t * d_hash);
+void cuda_base_whirlpool_cpu_hash_64f(const uint32_t threads, const uint32_t * d_hash, const uint32_t startNonce, uint32_t * d_resNonce, const uint64_t target);
+
+
+#endif // !CUDA_BASE_SKEIN512
+
+
diff --git a/Makefile.am b/Makefile.am
@@ -58,6 +58,7 @@ ccminer_SOURCES	= elist.h miner.h compat.h \
 			  Algo512/cuda_b_shabal512.cu Algo512/cuda_b_shabal512.h\
 			  Algo512/cuda_b_shavite512.cu Algo512/cuda_b_shavite512.h\
 			  Algo512/cuda_b_skein512.cu Algo512/cuda_b_skein512.h\
+			  Algo512/cuda_b_whirlpool.cu Algo512/cuda_b_whirlpool.h\
 			  crypto/xmr-rpc.cpp crypto/wildkeccak-cpu.cpp crypto/wildkeccak.cu \
 			  crypto/cryptolight.cu crypto/cryptolight-core.cu crypto/cryptolight-cpu.cpp \
 			  crypto/cryptonight.cu crypto/cryptonight-core.cu crypto/cryptonight-extra.cu \

diff --git a/README.txt b/README.txt
@@ -1,5 +1,5 @@
 
-ccminer dumax-0.9.1 (June 19th 2018)                "Seed pack"
+ccminer dumax-0.9.2 (June 24th 2018)              "Reden'ption"
 ---------------------------------------------------------------
 
 ***************************************************************
@@ -9,6 +9,7 @@ If you find this tool useful and like to support its continuous
 DumaxFr@github:
   LUX  : LWcYLSx37F37gHjbmvhwQPQ8PhyDYibmpr
   RVN  : RQXpsvSaVrGYo4tyGityWDNBQMFcnqANyj
+  REDN : RFVNEtRzAfeCDphuh4f3gPhN7hdX71BtEH
   BTC  : 1AtQXFbnzYTsjzy2bzSH6nPGxqZ32NG42T
 
 tpruvot@github:
@@ -284,6 +285,9 @@ features.
 
 >>> RELEASE HISTORY <<<
 
+  June 24th 2018  dumax-0.9.2
+                  Improved x16r/x16s (Whirlpool-64f)
+
   June 19th 2018  dumax-0.9.1
                   Improved x16r/x16s (Keccak512-64)
 				  Added gpu temp colors (from BCT tpruvot's ccminer)

diff --git a/ccminer.vcxproj b/ccminer.vcxproj
@@ -344,6 +344,7 @@
     <ClInclude Include="Algo512\cuda_b_shabal512.h" />
     <ClInclude Include="Algo512\cuda_b_shavite512.h" />
     <ClInclude Include="Algo512\cuda_b_skein512.h" />
+    <ClInclude Include="Algo512\cuda_b_whirlpool.h" />
     <ClInclude Include="cuda_uint2_utils.h" />
     <ClInclude Include="compat\pthreads\pthread.h" />
     <ClCompile Include="compat\winansi.c" />
@@ -415,6 +416,7 @@
     <CudaCompile Include="Algo512\cuda_b_shabal512.cu" />
     <CudaCompile Include="Algo512\cuda_b_skein512.cu" />
     <CudaCompile Include="Algo512\cuda_b_jh512.cu" />
+    <CudaCompile Include="Algo512\cuda_b_whirlpool.cu" />
     <CudaCompile Include="crypto\cryptolight.cu" />
     <CudaCompile Include="crypto\cryptolight-core.cu">
       <MaxRegCount>64</MaxRegCount>

diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters
@@ -621,6 +621,7 @@
     <ClInclude Include="Algo512\cuda_b_jh512.h" />
     <ClInclude Include="phi\cuda_phi2_skein512.h" />
     <ClInclude Include="phi\cuda_phi2.h" />
+    <ClInclude Include="Algo512\cuda_b_whirlpool.h" />
   </ItemGroup>
   <ItemGroup>
     <CudaCompile Include="cuda.cpp">
@@ -1013,6 +1014,7 @@
     <CudaCompile Include="phi\cuda_phi2_lyra2.cu" />
     <CudaCompile Include="phi\cuda_phi2_br_streebog_echo512.cu" />
     <CudaCompile Include="phi\phi2.cu" />
+    <CudaCompile Include="Algo512\cuda_b_whirlpool.cu" />
   </ItemGroup>
   <ItemGroup>
     <Image Include="res\ccminer.ico">

diff --git a/compat/ccminer-config.h b/compat/ccminer-config.h
@@ -164,7 +164,7 @@
 #define PACKAGE_URL "https://github.com/DumaxFr/ccminer"
 
 /* Define to the version of this package. */
-#define PACKAGE_VERSION "dumax-0.9.1"
+#define PACKAGE_VERSION "dumax-0.9.2"
 
 /* If using the C implementation of alloca, define if you know the
    direction of stack growth for your system; otherwise it will be

diff --git a/configure.ac b/configure.ac
@@ -1,4 +1,4 @@
-AC_INIT([ccminer], [dumax-0.9.1], [], [ccminer], [https://github.com/DumaxFr/ccminer])
+AC_INIT([ccminer], [dumax-0.9.2], [], [ccminer], [https://github.com/DumaxFr/ccminer])
 
 AC_PREREQ([2.59c])
 AC_CANONICAL_SYSTEM

diff --git a/res/ccminer.aps b/res/ccminer.aps
diff --git a/res/ccminer.rc b/res/ccminer.rc
@@ -61,13 +61,13 @@ IDI_ICON1               ICON                    "ccminer.ico"
 //
 
 VS_VERSION_INFO VERSIONINFO
- FILEVERSION 0,9,1,0
- PRODUCTVERSION 0,9,1,0
+ FILEVERSION 0,9,2,0
+ PRODUCTVERSION 0,9,2,0
  FILEFLAGSMASK 0x3fL
 #ifdef _DEBUG
- FILEFLAGS 0x21L
+ FILEFLAGS 0x23L
 #else
- FILEFLAGS 0x20L
+ FILEFLAGS 0x22L
 #endif
  FILEOS 0x40004L
  FILETYPE 0x1L
@@ -77,10 +77,10 @@ BEGIN
     BEGIN
         BLOCK "040904e4"
         BEGIN
-            VALUE "FileVersion", "0.9.1.0"
+            VALUE "FileVersion", "0.9.2.0"
             VALUE "LegalCopyright", "Copyright (C) 2018"
-            VALUE "ProductName", "ccminer dumax"
-            VALUE "ProductVersion", "0.9.1.0"
+            VALUE "ProductName", "ccminer/dumax"
+            VALUE "ProductVersion", "0.9.2.0"
         END
     END
     BLOCK "VarFileInfo"

diff --git a/x16x/cuda_x16x.h b/x16x/cuda_x16x.h
@@ -1,7 +1,5 @@
 #include "x11/cuda_x11.h"
 
-//#include "aes/cuda_aes.cuh"
-
 #include "Algo512/cuda_b_blake512.h"
 #include "Algo512/cuda_b_bmw512.h"
 #include "Algo512/cuda_b_cubehash512.h"
@@ -15,26 +13,10 @@
 #include "Algo512/cuda_b_shabal512.h"
 #include "Algo512/cuda_b_shavite512.h"
 #include "Algo512/cuda_b_skein512.h"
+#include "Algo512/cuda_b_whirlpool.h"
 
-//extern void x13_hamsi512_cpu_init(int thr_id, uint32_t threads);
-//extern void x13_hamsi512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
-
-//extern void x13_fugue512_cpu_init(int thr_id, uint32_t threads);
-//extern void x13_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
-//extern void x13_fugue512_cpu_free(int thr_id);
-
-//extern void x14_shabal512_cpu_init(int thr_id, uint32_t threads);
-//extern void x14_shabal512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
-
-extern void x15_whirlpool_cpu_init(int thr_id, uint32_t threads, int flag);
-extern void x15_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
-extern void x15_whirlpool_cpu_free(int thr_id);
-
-//extern void x17_sha512_cpu_init(int thr_id, uint32_t threads);
-//extern void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_hash);
-
-extern void x17_haval256_cpu_init(int thr_id, uint32_t threads);
-extern void x17_haval256_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_hash, const int outlen);
+//extern void x17_haval256_cpu_init(int thr_id, uint32_t threads);
+//extern void x17_haval256_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_hash, const int outlen);
 
 // ---- 80 bytes kernels
 
@@ -44,13 +26,6 @@ void quark_bmw512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNonce,
 void groestl512_setBlock_80(int thr_id, uint32_t *endiandata);
 void groestl512_cuda_hash_80(const int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
 
-//void skein512_cpu_setBlock_80(void *pdata);
-//void skein512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_hash, int swap);
-
-//void qubit_luffa512_cpu_init(int thr_id, uint32_t threads);
-//void qubit_luffa512_cpu_setBlock_80(void *pdata);
-//void qubit_luffa512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_hash, int order);
-
 void jh512_setBlock_80(int thr_id, uint32_t *endiandata);
 void jh512_cuda_hash_80(const int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
 
@@ -60,23 +35,9 @@ void keccak512_cuda_hash_80(const int thr_id, const uint32_t threads, const uint
 void cubehash512_setBlock_80(int thr_id, uint32_t* endiandata);
 void cubehash512_cuda_hash_80(const int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
 
-//void x16_shabal512_setBlock_80(void *pdata);
-//void x16_shabal512_cuda_hash_80(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
-
 void x16_simd512_setBlock_80(void *pdata);
 void x16_simd512_cuda_hash_80(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
 
-//void x16_hamsi512_setBlock_80(void *pdata);
-//void x16_hamsi512_cuda_hash_80(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
-
-//void x16_fugue512_cpu_init(int thr_id, uint32_t threads);
-//void x16_fugue512_cpu_free(int thr_id);
-//void x16_fugue512_setBlock_80(void *pdata);
-//void x16_fugue512_cuda_hash_80(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
-
 void x16_whirlpool512_init(int thr_id, uint32_t threads);
 void x16_whirlpool512_setBlock_80(void* endiandata);
 void x16_whirlpool512_hash_80(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
-
-//void x16_sha512_setBlock_80(void *pdata);
-//void x16_sha512_cuda_hash_80(int thr_id, const uint32_t threads, const uint32_t startNonce, uint32_t *d_hash);
diff --git a/x16x/x16x.cu b/x16x/x16x.cu
@@ -443,8 +443,9 @@ extern "C" int scanhash_x16x(int thr_id, struct work* work, uint32_t max_nonce,
         x11_simd512_cpu_init(thr_id, throughput); // 64
         cuda_base_echo512_cpu_init(thr_id);
         cuda_base_hamsi512_cpu_init();
-        x15_whirlpool_cpu_init(thr_id, throughput, 0);
-        x16_whirlpool512_init(thr_id, throughput);
+        //x15_whirlpool_cpu_init(thr_id, throughput, 0);
+        cuda_base_whirlpool_cpu_init();
+        x16_whirlpool512_init(thr_id, throughput); // 80
         cuda_base_sha512_cpu_init();
 
         CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], (size_t)64 * throughput), 0);
@@ -776,8 +777,13 @@ extern "C" int scanhash_x16x(int thr_id, struct work* work, uint32_t max_nonce,
                     TRACE("shabal :");
                     break;
                 case WHIRLPOOL:
-                    x15_whirlpool_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], i);
-                    TRACE("shabal :");
+                    if (i == 15) {
+                        cudaHashFinalDone = true;
+                        cuda_base_whirlpool_cpu_hash_64f(throughput, d_hash[thr_id], pdata[19], d_x16ResNonce[thr_id], *(uint64_t*)&ptarget[6]);
+                    } else {
+                        cuda_base_whirlpool_cpu_hash_64(throughput, d_hash[thr_id]);
+                    }
+                    TRACE("whirlpool:");
                     break;
                 case SHA512:
                     if (i == 15) {
@@ -918,7 +924,7 @@ extern "C" void free_x16r(int thr_id) {
 
     quark_groestl512_cpu_free(thr_id);
     x11_simd512_cpu_free(thr_id);
-    x15_whirlpool_cpu_free(thr_id);
+    //x15_whirlpool_cpu_free(thr_id);
 
     cuda_check_cpu_free(thr_id);