Clusterfuck edit mode engaged: Yoinked some version of cbuchner's ccm…

…iner, grabbed some shavite stuff from djm's fork, slapped together an implementation for s3 hash.

JHA/cuda_jha_keccak512.cu

@@ -567,8 +567,6 @@ __host__ void jackpot_keccak512_cpu_hash(int thr_id, int threads, uint32_t start
     // Größe des dynamischen Shared Memory Bereichs
     size_t shared_size = 0;
 
-//    fprintf(stderr, "threads=%d, %d blocks, %d threads per block, %d bytes shared\n", threads, grid.x, block.x, shared_size);
-
     jackpot_keccak512_gpu_hash<<<grid, block, shared_size>>>(threads, startNounce, (uint64_t*)d_hash);
     MyStreamSynchronize(NULL, order, thr_id);
 }

JHA/jackpotcoin.cu

@@ -101,14 +101,12 @@ extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
 {
 	const uint32_t first_nonce = pdata[19];
 
-	// TODO: entfernen für eine Release! Ist nur zum Testen!
 	if (opt_benchmark)
 		((uint32_t*)ptarget)[7] = 0x0000ff;
 
 	const uint32_t Htarg = ptarget[7];
 
 	const int throughput = 256*4096*4; // 100;
-	//const int throughput = 256*256*2+100; // 100;
 
 	static bool init[8] = {0,0,0,0,0,0,0,0};
 	if (!init[thr_id])
@@ -167,16 +165,18 @@ extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
 			quark_jh512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++);
 		}
 
-		// Runde 2 (ohne Gröstl)
+		// Runde 3 (komplett)
 
 		// jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01)
 		jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id],
 			d_branch1Nonces[thr_id], &nrm1,
-			d_branch3Nonces[thr_id], &nrm3,
+			d_branch2Nonces[thr_id], &nrm2,
 			order++);
 
-		// verfolge den skein-pfad weiter
-		quark_skein512_cpu_hash_64(thr_id, nrm3, pdata[19], d_branch3Nonces[thr_id], d_hash[thr_id], order++);
+		if (nrm1+nrm2 == nrm3) {
+			quark_groestl512_cpu_hash_64(thr_id, nrm1, pdata[19], d_branch1Nonces[thr_id], d_hash[thr_id], order++);
+			quark_skein512_cpu_hash_64(thr_id, nrm2, pdata[19], d_branch2Nonces[thr_id], d_hash[thr_id], order++);
+		}
 
 		// jackpotNonces in branch1/2 aufsplitten gemäss if (hash[0] & 0x01)
 		jackpot_compactTest_cpu_hash_64(thr_id, nrm3, pdata[19], d_hash[thr_id], d_branch3Nonces[thr_id],
@@ -226,7 +226,7 @@ extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
 			if ((vhash64[7]<=Htarg) && fulltest(vhash64, ptarget)) {
 
 				pdata[19] = foundNonce;
-				*hashes_done = (foundNonce - first_nonce + 1)/4;
+				*hashes_done = (foundNonce - first_nonce + 1)/2;
 				//applog(LOG_INFO, "GPU #%d: result for nonce $%08X does validate on CPU (%d rounds)!", thr_id, foundNonce, rounds);
 				return 1;
 			} else {
@@ -238,6 +238,6 @@ extern "C" int scanhash_jackpot(int thr_id, uint32_t *pdata,
 
 	} while (pdata[19] < max_nonce && !work_restart[thr_id].restart);
 
-	*hashes_done = (pdata[19] - first_nonce + 1)/4;
+	*hashes_done = (pdata[19] - first_nonce + 1)/2;
 	return 0;
 }

Makefile.am

@@ -33,7 +33,7 @@ ccminer_SOURCES		= elist.h miner.h compat.h \
 			  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 quark/quarkcoin.cu quark/animecoin.cu \
 			  quark/cuda_quark_compactionTest.cu \
-			  cuda_nist5.cu \
+			  cuda_nist5.cu cuda_s3.cu \
 			  sph/cubehash.c sph/echo.c sph/luffa.c sph/shavite.c sph/simd.c \
 			  sph/hamsi.c sph/hamsi_helper.c sph/sph_hamsi.h \
 			  x13/x13.cu x13/cuda_x13_hamsi512.cu x13/cuda_x13_fugue512.cu \

README.txt

@@ -1,5 +1,5 @@
 
-ccMiner release 1.0 (May 10th 2014) - "Did anyone say X11?"
+ccMiner release 1.2 (June 15th 2014) - "Killer Groestl + X13"
 -------------------------------------------------------------
 
 ***************************************************************
@@ -30,13 +30,12 @@ FugueCoin
 GroestlCoin & Myriad-Groestl
 JackpotCoin
 QuarkCoin family & AnimeCoin
+TalkCoin
+DarkCoin and other X11 coins
 
 where some of these coins have a VERY NOTABLE nVidia advantage
 over competing AMD (OpenCL) implementations.
 
-X11 algo is being worked on. It will be released when we
-have achieved a nice nVidia advantage.
-
 We did not take a big effort on improving usability, so please set
 your parameters carefuly.
 
@@ -140,6 +139,15 @@ features.
 
 >>> RELEASE HISTORY <<<
 
+  June 15th 2014  add X13 and Diamond Groestl support.
+                  Thanks to tsiv and to Bombadil for the contributions!
+
+  June 14th 2014  released Killer Groestl quad version which I deem
+                  sufficiently hard to port over to AMD. It isn't
+                  the fastest option for Compute 3.5 and 5.0 cards,
+                  but it is still much faster than the table based
+                  versions.
+
   May 10th 2014   added X11, but without the bells & whistles
                   (no killer Groestl, SIMD hash quite slow still)