correct kernels and merkle tree calculation

algorithm.c

@@ -9,6 +9,7 @@
 
 #include "algorithm.h"
 #include "sph/sph_sha2.h"
+#include "sph/sph_gost.h"
 #include "ocl.h"
 #include "ocl/build_kernel.h"
 
@@ -106,6 +107,13 @@ void gen_hash(const unsigned char *data, unsigned int len, unsigned char *hash)
   sph_sha256_close(&ctx_sha2, hash);
 }
 
+void gostcoin_gen_hash(const unsigned char *data, unsigned int len, unsigned char *hash)
+{
+	unsigned char h1[64];
+	sph_gost512(h1, (const void*)data, len);
+	sph_gost256(hash, (const void*)h1, 64);
+}
+
 void sha256d_midstate(struct work *work)
 {
   unsigned char data[64];
@@ -1249,7 +1257,8 @@ static algorithm_settings_t algos[] = {
 
   { "pascal", ALGO_PASCAL, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 0, 0, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, pascal_regenhash, pascal_midstate, NULL, queue_pascal_kernel, NULL, NULL },
 
-  { "gostcoin-mod", ALGO_GOSTCOIN, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 2, 4 * 8 * 4194304, 0, gostcoin_regenhash, NULL, NULL, queue_gostcoin_mod_kernel, NULL, NULL },	
+
+ { "gostcoin-mod", ALGO_GOSTCOIN, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 0, 4 * 8 * 4194304, 0, gostcoin_regenhash, NULL, NULL, queue_gostcoin_mod_kernel, gostcoin_gen_hash, NULL },	
 
   // Terminator (do not remove)
   { NULL, ALGO_UNK, "", 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, NULL, NULL, NULL, NULL, NULL }

kernel/gostcoin-mod.cl

@@ -930,6 +930,6 @@ __kernel void search(__global unsigned char* block, volatile __global uint* outp
     GOST_g_0(hash, hash1);		
     // result is first 32 bytes of hash
     if (SWAP8 (hash[0]) <= target)
-		output[output[0xFF]++] = nonce; 
+		output[output[0xFF]++] = SWAP4 (nonce); 
 }
 

sgminer.c

@@ -6145,13 +6145,17 @@ static void gen_stratum_work(struct pool *pool, struct work *work)
 
   /* Downgrade to a read lock to read off the pool variables */
   cg_dwlock(&pool->data_lock);
-
+	
   if (pool->algorithm.type != ALGO_DECRED && pool->algorithm.type != ALGO_SIA && pool->algorithm.type != ALGO_PASCAL) {
     /* Generate merkle root */
     pool->algorithm.gen_hash(pool->coinbase, pool->swork.cb_len, merkle_root);
     memcpy(merkle_sha, merkle_root, 32);
-    for (i = 0; i < pool->swork.merkles; i++) {
-      memcpy(merkle_sha + 32, pool->swork.merkle_bin[i], 32);
+    for (i = 0; i < pool->swork.merkles; i++) 
+	{
+	  if (pool->algorithm.type != ALGO_GOSTCOIN) // TODO: remove after hard fork 1	
+		memcpy(merkle_sha + 32, merkle_sha, 32); 	
+	  else		
+      	memcpy(merkle_sha + 32, pool->swork.merkle_bin[i], 32);
       gen_hash(merkle_sha, 64, merkle_root);
       memcpy(merkle_sha, merkle_root, 32);
     }