Krnlx Nvidia xevan miner - 3.3+ mh on 1070, ~6mh on 80ti FREE, OPENSOURCE

[palgin]

Didn't help cause this file are allready exist in source files. If i remoove it and add again - same errore:

Error   1   error LNK2001: unresolved external symbol "void __cdecl xevan_sha512_cpu_hash_64(int,int,unsigned int *)" (?xevan_sha512_cpu_hash_64@@YAXHHPEAI@Z)   D:\ccminer-xevan-master\x17.cu.obj

Nothing will help with Winbuild before any Linux user send hashing output using "debug" ver of x17.cu, so don't even bother. You can check yiimp bench, no ccminer-krnlx running on Win  

Think your error is caused by wrong Curl lib or Dev environment, you need VS2013 and Curl built with the same VS version, and don't forget about architecture (x86/x64).

[1715367604]

1715367604

[1715367604]

1715367604

[1715367604]

1715367604

[1715367604]

1715367604

[stas260385]

Didn't help cause this file are allready exist in source files. If i remoove it and add again - same errore:

Error   1   error LNK2001: unresolved external symbol "void __cdecl xevan_sha512_cpu_hash_64(int,int,unsigned int *)" (?xevan_sha512_cpu_hash_64@@YAXHHPEAI@Z)   D:\ccminer-xevan-master\x17.cu.obj

Nothing will help with Winbuild before any Linux user send hashing output using "debug" ver of x17.cu, so don't even bother. You can check yiimp bench, no ccminer-krnlx running on Win  

Think your error is caused by wrong Curl lib or Dev environment, you need VS2013 and Curl built with the same VS version, and don't forget about architecture (x86/x64).

All proper Curl libs, VS2013 and CUDA 7.5 are installed

[palgin]

All proper Curl libs, VS2013 and CUDA 7.5 are installed

cuda_sha512_lbry.cu - function you need is in this file, check if it's attached to your project.

[anorganix]

Finally managed to compile it for Windows x64 @ CUDA 7.5.
Of course it's not working, looking forward to the fix from the experienced ones.

[stas260385]

Finally managed to compile it for Windows x64 @ CUDA 7.5.
Of course it's not working, looking forward to the fix from the experienced ones.

https://i.imgur.com/Mgh3vBk.png

I've compiled too but the same issue...(

[integrale]

Great Job done Guys

have it succesfull compiled on Ubuntu 16.04 LTS Cuda 7.5  Nv 384.90

running on a poor small GTX 1050 @ 1.1Mh/s with  intensity of 19

[anorganix]

If anyone with Linux wants to help Win users, replace your x17.cu contents with this code:

Code: (x17.cu)

/**
 * X17 algorithm (X15 + sha512 + haval256)
 */

extern "C" {
#include "sph/sph_blake.h"
#include "sph/sph_bmw.h"
#include "sph/sph_groestl.h"
#include "sph/sph_skein.h"
#include "sph/sph_jh.h"
#include "sph/sph_keccak.h"

#include "sph/sph_luffa.h"
#include "sph/sph_cubehash.h"
#include "sph/sph_shavite.h"
#include "sph/sph_simd.h"
#include "sph/sph_echo.h"

#include "sph/sph_hamsi.h"
#include "sph/sph_fugue.h"

#include "sph/sph_shabal.h"
#include "sph/sph_whirlpool.h"

#include "sph/sph_sha2.h"
#include "sph/sph_haval.h"
}

#include "miner.h"
#include "cuda_helper.h"
#include "x11/cuda_x11.h"

#define NBN 2

// Memory for the hash functions
static uint32_t *d_hash[MAX_GPUS];
static uint32_t *d_resNonce[MAX_GPUS];
static uint32_t *h_resNonce[MAX_GPUS];

void print_hash(unsigned int *data, int size){
	for (int i = 0; i<size; i++)
		gpulog(LOG_WARNING, 0, "%x ", data[i]);
	//gpulog(LOG_WARNING, 0, "-------------");
}

extern void x13_hamsi_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x14_shabal512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x15_whirlpool_cpu_init(int thr_id, uint32_t threads, int mode);
extern void x15_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x15_whirlpool_cpu_free(int thr_id);

extern void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x17_haval256_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t* resNonce, uint64_t target);
extern void bmw256_cpu_hash_32_full(int thr_id, uint32_t threads, uint32_t *g_hash);
extern void quark_bmw512_cpu_hash_64x(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_hash);
extern void quark_groestl512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void groestl512_cpu_init(int thr_id, uint32_t threads);
extern void groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_skein512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void qubit_luffa512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_outputHash);
extern void x11_cubehash512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_init(int thr_id, uint32_t threads);
extern void xevan_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_hamsi512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_sha512_cpu_hash_64(int thr_id, int threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void xevan_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t *resNonce, uint64_t target);
extern void xevan_groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64_A(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void quark_blake512_cpu_hash_128(int thr_id, uint32_t threads, uint32_t *d_outputHash);
extern void quark_groestl512_cpu_hash_128(int thr_id, uint32_t threads,  uint32_t *d_hash);
extern void x11_luffa512_cpu_hash_128(int thr_id, uint32_t threads,uint32_t *d_hash);



// X17 CPU Hash (Validation)
extern "C" void x17hash(void *output, const void *input)
{
	uint32_t _ALIGN(64) hash[32]; // 128 bytes required
	uint32_t input_zero[20] = { 0 };
	const int dataLen = 128;

//return;
	sph_blake512_context     ctx_blake;
	sph_bmw512_context       ctx_bmw;
	sph_groestl512_context   ctx_groestl;
	sph_skein512_context     ctx_skein;
	sph_jh512_context        ctx_jh;
	sph_keccak512_context    ctx_keccak;
	sph_luffa512_context     ctx_luffa;
	sph_cubehash512_context  ctx_cubehash;
	sph_shavite512_context   ctx_shavite;
	sph_simd512_context      ctx_simd;
	sph_echo512_context      ctx_echo;
	sph_hamsi512_context     ctx_hamsi;
	sph_fugue512_context     ctx_fugue;
	sph_shabal512_context    ctx_shabal;
	sph_whirlpool_context    ctx_whirlpool;
	sph_sha512_context       ctx_sha512;
	sph_haval256_5_context   ctx_haval;

	print_hash(input_zero,20);
	gpulog(LOG_WARNING, 0, "--INPUT ZEROES--");

	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, input_zero, 80);
	sph_blake512_close(&ctx_blake, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BLAKE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	memset(&hash[16], 0, 64);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--AFTER MEMSET--");

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BMW512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--GROESTL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SKEIN512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--JH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--KECCAK512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--LUFFA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--CUBEHASH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHAVITE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BLAKE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--ECHO512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAMSI512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--FUGUE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHABAL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--WHIRLPOOL--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAVAL256--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	memset(&hash[8], 0, dataLen - 32);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--AFTER MEMSET--");

	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, hash, dataLen);
	sph_blake512_close(&ctx_blake, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BLAKE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BMW512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--GROESTL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SKEIN512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--JH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--KECCAK512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--LUFFA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--CUBEHASH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHAVITE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SIMD512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--ECHO512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAMSI512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--FUGUE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHABAL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--WHIRLPOOL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAVAL256--");

	for (int i = 0; i<32; i++) hash[i] = 0;

	memcpy(output, hash, 32);
}

static bool init[MAX_GPUS] = { 0 };





extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done){

	int dev_id = device_map[thr_id];

	uint32_t *pdata = work->data;
	uint32_t *ptarget = work->target;
	const uint32_t first_nonce = pdata[19];
	uint32_t default_throughput;
	if(device_sm[dev_id]<=500) default_throughput = 1<<20;
	else if(device_sm[dev_id]<=520) default_throughput = 1<<21;
	else if(device_sm[dev_id]>520) default_throughput = (1<<22) + (1<<21);
	
	if((strstr(device_name[dev_id], "1060")))default_throughput = 1<<20;
	if((strstr(device_name[dev_id], "1070")))default_throughput = 1<<20;
	if((strstr(device_name[dev_id], "1080")))default_throughput = 1<<20;
	
	uint32_t throughput = cuda_default_throughput(thr_id, default_throughput); // 19=256*256*8;
	if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);

	throughput&=0xFFFFFF70; //multiples of 128 due to simd_echo kernel

	if (opt_benchmark)
		((uint32_t*)ptarget)[7] = 0xff;

	gpulog(LOG_INFO,thr_id,"target %x %x %x",ptarget[5], ptarget[6], ptarget[7]);
        gpulog(LOG_INFO,thr_id,"target %llx",*(uint64_t*)&ptarget[6]);

	if (!init[thr_id])
	{
		cudaSetDevice(device_map[thr_id]);
		if (opt_cudaschedule == -1 && gpu_threads == 1) {
			cudaDeviceReset();
			// reduce cpu usage
			cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
		}
		gpulog(LOG_INFO,thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);

		x15_whirlpool_cpu_init(thr_id, throughput, 0);
		groestl512_cpu_init(thr_id, throughput);
		x11_simd512_cpu_init(thr_id, throughput);
		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 8 * sizeof(uint64_t) * throughput));
		CUDA_SAFE_CALL(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t)));
		h_resNonce[thr_id] = (uint32_t*) malloc(NBN  * 8 * sizeof(uint32_t));
		if(h_resNonce[thr_id] == NULL){
			gpulog(LOG_ERR,thr_id,"Host memory allocation failed");
			exit(EXIT_FAILURE);
		}		
		init[thr_id] = true;
	}

	uint32_t _ALIGN(64) endiandata[20];
	for (int k=0; k < 20; k++)
		be32enc(&endiandata[k], pdata[k]);

	quark_blake512_cpu_setBlock_80(thr_id, endiandata);
	cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));


	do {
		// Hash with CUDA

		quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]);//A

		quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);

		quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);

		quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);//A //fast

		x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A

		x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //A 256

		xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//P slow r2
		 
        x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);  //A slow r3
		 
		x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//A
		 
        x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
		x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast ++
		 
		//x13_hamsi_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
		xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //opt2
		 
		xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
		xevan_haval512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
        quark_blake512_cpu_hash_128(thr_id, throughput,  d_hash[thr_id]);//BAD
		 
        quark_bmw512_cpu_hash_64x(thr_id, throughput, NULL, d_hash[thr_id]);
		 
		quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);
		 
        quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);
		 
        quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);
		 
        x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A
		 
        x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//move to shared
		 
        x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); 
		 
        x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
		//x13_hamsi_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
        x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
		xevan_haval512_cpu_hash_64_final(thr_id, throughput, d_hash[thr_id],d_resNonce[thr_id],*(uint64_t*)&ptarget[6]);
		 

		cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost);

		if (h_resNonce[thr_id][0] != UINT32_MAX){
			const uint32_t Htarg = ptarget[7];
			const uint32_t startNounce = pdata[19];
			uint32_t vhash64[8];
			be32enc(&endiandata[19], startNounce + h_resNonce[thr_id][0]);
			x17hash(vhash64, endiandata);
//			*hashes_done = pdata[19] - first_nonce + throughput + 1;
//			pdata[19] = startNounce + h_resNonce[thr_id][0];
			gpulog(LOG_WARNING, 0,"NONCE FOUND ");
//			return 1;
			if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
				int res = 1;
				*hashes_done = pdata[19] - first_nonce + throughput + 1;
				work_set_target_ratio(work, vhash64);
				pdata[19] = startNounce + h_resNonce[thr_id][0];
				if (h_resNonce[thr_id][1] != UINT32_MAX) {
					pdata[21] = startNounce+h_resNonce[thr_id][1];
					if(!opt_quiet)
						gpulog(LOG_BLUE,dev_id,"Found 2nd nonce: %08x", pdata[21]);
					be32enc(&endiandata[19], pdata[21]);
					x17hash(vhash64, endiandata);
					if (bn_hash_target_ratio(vhash64, ptarget) > work->shareratio[0]){
						work_set_target_ratio(work, vhash64);
						xchg(pdata[19],pdata[21]);
					}
					res++;
				}
				return res;
			}
			else {
				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", h_resNonce[thr_id][0]);
				cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));				
			}
		}

		pdata[19] += throughput;
	} while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > (uint64_t)throughput + pdata[19]));

	*hashes_done = pdata[19] - first_nonce + 1;

	return 0;
}

// cleanup
extern "C" void free_x17(int thr_id)
{
	if (!init[thr_id])
		return;

	cudaDeviceSynchronize();

	free(h_resNonce[thr_id]);
	cudaFree(d_resNonce[thr_id]);
	cudaFree(d_hash[thr_id]);

	x11_simd_echo_512_cpu_free(thr_id);
	x15_whirlpool_cpu_free(thr_id);
	cudaDeviceSynchronize();
	init[thr_id] = false;
}

Then recompile and run for a few seconds, before first rejected share. Send me your output in PM or publish it here.

Thanks for everybody in advance!

Bump!
Any help much appreciated.

[palgin]

If anyone with Linux wants to help Win users, replace your x17.cu contents with this code:

Code: (x17.cu)

/**
 * X17 algorithm (X15 + sha512 + haval256)
 */

extern "C" {
#include "sph/sph_blake.h"
#include "sph/sph_bmw.h"
#include "sph/sph_groestl.h"
#include "sph/sph_skein.h"
#include "sph/sph_jh.h"
#include "sph/sph_keccak.h"

#include "sph/sph_luffa.h"
#include "sph/sph_cubehash.h"
#include "sph/sph_shavite.h"
#include "sph/sph_simd.h"
#include "sph/sph_echo.h"

#include "sph/sph_hamsi.h"
#include "sph/sph_fugue.h"

#include "sph/sph_shabal.h"
#include "sph/sph_whirlpool.h"

#include "sph/sph_sha2.h"
#include "sph/sph_haval.h"
}

#include "miner.h"
#include "cuda_helper.h"
#include "x11/cuda_x11.h"

#define NBN 2

// Memory for the hash functions
static uint32_t *d_hash[MAX_GPUS];
static uint32_t *d_resNonce[MAX_GPUS];
static uint32_t *h_resNonce[MAX_GPUS];

void print_hash(unsigned int *data, int size){
	for (int i = 0; i<size; i++)
		gpulog(LOG_WARNING, 0, "%x ", data[i]);
	//gpulog(LOG_WARNING, 0, "-------------");
}

extern void x13_hamsi_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x14_shabal512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x15_whirlpool_cpu_init(int thr_id, uint32_t threads, int mode);
extern void x15_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x15_whirlpool_cpu_free(int thr_id);

extern void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x17_haval256_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t* resNonce, uint64_t target);
extern void bmw256_cpu_hash_32_full(int thr_id, uint32_t threads, uint32_t *g_hash);
extern void quark_bmw512_cpu_hash_64x(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_hash);
extern void quark_groestl512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void groestl512_cpu_init(int thr_id, uint32_t threads);
extern void groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_skein512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void qubit_luffa512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_outputHash);
extern void x11_cubehash512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_init(int thr_id, uint32_t threads);
extern void xevan_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_hamsi512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_sha512_cpu_hash_64(int thr_id, int threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void xevan_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t *resNonce, uint64_t target);
extern void xevan_groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64_A(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void quark_blake512_cpu_hash_128(int thr_id, uint32_t threads, uint32_t *d_outputHash);
extern void quark_groestl512_cpu_hash_128(int thr_id, uint32_t threads,  uint32_t *d_hash);
extern void x11_luffa512_cpu_hash_128(int thr_id, uint32_t threads,uint32_t *d_hash);



// X17 CPU Hash (Validation)
extern "C" void x17hash(void *output, const void *input)
{
	uint32_t _ALIGN(64) hash[32]; // 128 bytes required
	uint32_t input_zero[20] = { 0 };
	const int dataLen = 128;

//return;
	sph_blake512_context     ctx_blake;
	sph_bmw512_context       ctx_bmw;
	sph_groestl512_context   ctx_groestl;
	sph_skein512_context     ctx_skein;
	sph_jh512_context        ctx_jh;
	sph_keccak512_context    ctx_keccak;
	sph_luffa512_context     ctx_luffa;
	sph_cubehash512_context  ctx_cubehash;
	sph_shavite512_context   ctx_shavite;
	sph_simd512_context      ctx_simd;
	sph_echo512_context      ctx_echo;
	sph_hamsi512_context     ctx_hamsi;
	sph_fugue512_context     ctx_fugue;
	sph_shabal512_context    ctx_shabal;
	sph_whirlpool_context    ctx_whirlpool;
	sph_sha512_context       ctx_sha512;
	sph_haval256_5_context   ctx_haval;

	print_hash(input_zero,20);
	gpulog(LOG_WARNING, 0, "--INPUT ZEROES--");

	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, input_zero, 80);
	sph_blake512_close(&ctx_blake, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BLAKE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	memset(&hash[16], 0, 64);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--AFTER MEMSET--");

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BMW512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--GROESTL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SKEIN512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--JH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--KECCAK512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--LUFFA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--CUBEHASH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHAVITE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BLAKE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--ECHO512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAMSI512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--FUGUE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHABAL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--WHIRLPOOL--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAVAL256--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	memset(&hash[8], 0, dataLen - 32);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--AFTER MEMSET--");

	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, hash, dataLen);
	sph_blake512_close(&ctx_blake, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BLAKE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--BMW512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--GROESTL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SKEIN512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--JH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--KECCAK512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--LUFFA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--CUBEHASH512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHAVITE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SIMD512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--ECHO512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAMSI512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--FUGUE512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHABAL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--WHIRLPOOL512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--SHA512--");
	//for (int i = 0; i<32; i++) hash[i] = 0;

	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);

	print_hash(hash, 32);
	gpulog(LOG_WARNING, 0, "--HAVAL256--");

	for (int i = 0; i<32; i++) hash[i] = 0;

	memcpy(output, hash, 32);
}

static bool init[MAX_GPUS] = { 0 };





extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done){

	int dev_id = device_map[thr_id];

	uint32_t *pdata = work->data;
	uint32_t *ptarget = work->target;
	const uint32_t first_nonce = pdata[19];
	uint32_t default_throughput;
	if(device_sm[dev_id]<=500) default_throughput = 1<<20;
	else if(device_sm[dev_id]<=520) default_throughput = 1<<21;
	else if(device_sm[dev_id]>520) default_throughput = (1<<22) + (1<<21);
	
	if((strstr(device_name[dev_id], "1060")))default_throughput = 1<<20;
	if((strstr(device_name[dev_id], "1070")))default_throughput = 1<<20;
	if((strstr(device_name[dev_id], "1080")))default_throughput = 1<<20;
	
	uint32_t throughput = cuda_default_throughput(thr_id, default_throughput); // 19=256*256*8;
	if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);

	throughput&=0xFFFFFF70; //multiples of 128 due to simd_echo kernel

	if (opt_benchmark)
		((uint32_t*)ptarget)[7] = 0xff;

	gpulog(LOG_INFO,thr_id,"target %x %x %x",ptarget[5], ptarget[6], ptarget[7]);
        gpulog(LOG_INFO,thr_id,"target %llx",*(uint64_t*)&ptarget[6]);

	if (!init[thr_id])
	{
		cudaSetDevice(device_map[thr_id]);
		if (opt_cudaschedule == -1 && gpu_threads == 1) {
			cudaDeviceReset();
			// reduce cpu usage
			cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
		}
		gpulog(LOG_INFO,thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);

		x15_whirlpool_cpu_init(thr_id, throughput, 0);
		groestl512_cpu_init(thr_id, throughput);
		x11_simd512_cpu_init(thr_id, throughput);
		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 8 * sizeof(uint64_t) * throughput));
		CUDA_SAFE_CALL(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t)));
		h_resNonce[thr_id] = (uint32_t*) malloc(NBN  * 8 * sizeof(uint32_t));
		if(h_resNonce[thr_id] == NULL){
			gpulog(LOG_ERR,thr_id,"Host memory allocation failed");
			exit(EXIT_FAILURE);
		}		
		init[thr_id] = true;
	}

	uint32_t _ALIGN(64) endiandata[20];
	for (int k=0; k < 20; k++)
		be32enc(&endiandata[k], pdata[k]);

	quark_blake512_cpu_setBlock_80(thr_id, endiandata);
	cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));


	do {
		// Hash with CUDA

		quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]);//A

		quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);

		quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);

		quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);//A //fast

		x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A

		x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //A 256

		xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//P slow r2
		 
        x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);  //A slow r3
		 
		x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//A
		 
        x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
		x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast ++
		 
		//x13_hamsi_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
		xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //opt2
		 
		xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
		xevan_haval512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		 
        quark_blake512_cpu_hash_128(thr_id, throughput,  d_hash[thr_id]);//BAD
		 
        quark_bmw512_cpu_hash_64x(thr_id, throughput, NULL, d_hash[thr_id]);
		 
		quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);
		 
        quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);
		 
        quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);
		 
        x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A
		 
        x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//move to shared
		 
        x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); 
		 
        x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
		//x13_hamsi_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
        x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
        xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
		 
		xevan_haval512_cpu_hash_64_final(thr_id, throughput, d_hash[thr_id],d_resNonce[thr_id],*(uint64_t*)&ptarget[6]);
		 

		cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost);

		if (h_resNonce[thr_id][0] != UINT32_MAX){
			const uint32_t Htarg = ptarget[7];
			const uint32_t startNounce = pdata[19];
			uint32_t vhash64[8];
			be32enc(&endiandata[19], startNounce + h_resNonce[thr_id][0]);
			x17hash(vhash64, endiandata);
//			*hashes_done = pdata[19] - first_nonce + throughput + 1;
//			pdata[19] = startNounce + h_resNonce[thr_id][0];
			gpulog(LOG_WARNING, 0,"NONCE FOUND ");
//			return 1;
			if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
				int res = 1;
				*hashes_done = pdata[19] - first_nonce + throughput + 1;
				work_set_target_ratio(work, vhash64);
				pdata[19] = startNounce + h_resNonce[thr_id][0];
				if (h_resNonce[thr_id][1] != UINT32_MAX) {
					pdata[21] = startNounce+h_resNonce[thr_id][1];
					if(!opt_quiet)
						gpulog(LOG_BLUE,dev_id,"Found 2nd nonce: %08x", pdata[21]);
					be32enc(&endiandata[19], pdata[21]);
					x17hash(vhash64, endiandata);
					if (bn_hash_target_ratio(vhash64, ptarget) > work->shareratio[0]){
						work_set_target_ratio(work, vhash64);
						xchg(pdata[19],pdata[21]);
					}
					res++;
				}
				return res;
			}
			else {
				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", h_resNonce[thr_id][0]);
				cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));				
			}
		}

		pdata[19] += throughput;
	} while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > (uint64_t)throughput + pdata[19]));

	*hashes_done = pdata[19] - first_nonce + 1;

	return 0;
}

// cleanup
extern "C" void free_x17(int thr_id)
{
	if (!init[thr_id])
		return;

	cudaDeviceSynchronize();

	free(h_resNonce[thr_id]);
	cudaFree(d_resNonce[thr_id]);
	cudaFree(d_hash[thr_id]);

	x11_simd_echo_512_cpu_free(thr_id);
	x15_whirlpool_cpu_free(thr_id);
	cudaDeviceSynchronize();
	init[thr_id] = false;
}

Then recompile and run for a few seconds, before first rejected share. Send me your output in PM or publish it here.

Thanks for everybody in advance!

Bump!
Any help much appreciated.

Seems no *nix user wants to help MS guys 

Will have to wait for krnlx to run it on his machine.

[integrale]

im running on "nix"   if i can help ?  im not a pro   but can try it

[stas260385]

It's from Z-pool bench:


xevan   26h   1070   GeForce GTX 1070   10de:119d   SM 6.1   2.27 MH   21   1822   109   20.791 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 382.53
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.26 MH   21   1797   101   22.363 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   16h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.26 MH   21   1797   108   20.905 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.26 MH   21   1797   107   21.088 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.25 MH   21   1797   107   21.074 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   26h   1070   GeForce GTX 1070   10de:119d   SM 6.1   2.25 MH   21   1822   110   20.479 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 382.53
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.25 MH   21   1797   104   21.658 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   16h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.25 MH   21   1797   140   16.056 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.25 MH   21   1797   91   24.674 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.24 MH   21   1797   108   20.739 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.24 MH   21   1797   105   21.308 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   17h   1070   EVGA GTX 1070 FTW   3842:6276   SM 6.1   2.24 MH   21   1797   96   23.289 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 385.28
xevan   2d   1070   GeForce GTX 1070   10de:119d   SM 6.1   2 MH   21   1683   79   25.305 k   ccminer/8.13-KlausT   linux   CUDA 8.0 375.66
xevan   2d   1080 Ti   GeForce GTX 1080 Ti   10de:120f   SM 6.1   1.97 MH   21   1721   81   24.36 k   cpuminer-opt/6.6.6   win64   CUDA 8.0 384.94
xevan   2d   1080 Ti   GeForce GTX 1080 Ti   10de:120f   SM 6.1   1.9 MH   21   1721   186   10.209 k   ccminer/8.13-dj   win64   CUDA 8.0 384.94
xevan   2d   1080 Ti   GeForce GTX 1080 Ti   10de:120f   SM 6.1   1.71 MH   21   1721   96   17.835 k   ccminer/8.13-dj   win64   CUDA 8.0 384.94
xevan   2d   1080 Ti   GeForce GTX 1080 Ti   10de:120f   SM 6.1   1.61 MH   21   1721   153   10.535 k   ccminer/8.13-dj   win64   CUDA 8.0 384.94
xevan   2d   1080 Ti   GeForce GTX 1080 Ti   10de:120f   SM 6.1   1.54 MH   21   1721   117   13.157 k   ccminer/8.13-dj   win64   CUDA 8.0 384.94
xevan   2d   1080 Ti   GeForce GTX 1080 Ti   10de:120f   SM 6.1   1.52 MH   21   1721   106   14.386 k   ccminer/8.13-dj   win64   CUDA 8.0 384.94
xevan   2d   1080 Ti   GeForce GTX 1080 Ti   10de:120f   SM 6.1   1.48 MH   21   1721   124   11.964 k   ccminer/8.13-dj   win64   CUDA 8.0 384.94

[sickofscamcoins]

http://paste.ubuntu.com/25680610/

this is results of output, replacing x17.cu and recompile with one posted palgin.

[palgin]

http://paste.ubuntu.com/25680610/

this is results of output, replacing x17.cu and recompile with one posted palgin.

Thank you, will check it!

[integrale]

copied from my Linux build   x17.cu  hopefully it helps

Code:

/**
 * X17 algorithm (X15 + sha512 + haval256)
 */

extern "C" {
#include "sph/sph_blake.h"
#include "sph/sph_bmw.h"
#include "sph/sph_groestl.h"
#include "sph/sph_skein.h"
#include "sph/sph_jh.h"
#include "sph/sph_keccak.h"

#include "sph/sph_luffa.h"
#include "sph/sph_cubehash.h"
#include "sph/sph_shavite.h"
#include "sph/sph_simd.h"
#include "sph/sph_echo.h"

#include "sph/sph_hamsi.h"
#include "sph/sph_fugue.h"

#include "sph/sph_shabal.h"
#include "sph/sph_whirlpool.h"

#include "sph/sph_sha2.h"
#include "sph/sph_haval.h"
}

#include "miner.h"
#include "cuda_helper.h"
#include "x11/cuda_x11.h"

#define NBN 2

// Memory for the hash functions
static uint32_t *d_hash[MAX_GPUS];
static uint32_t *d_resNonce[MAX_GPUS];
static uint32_t *h_resNonce[MAX_GPUS];

extern void x13_hamsi_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x14_shabal512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x15_whirlpool_cpu_init(int thr_id, uint32_t threads, int mode);
extern void x15_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x15_whirlpool_cpu_free(int thr_id);

extern void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x17_haval256_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t* resNonce, uint64_t target);
extern void bmw256_cpu_hash_32_full(int thr_id, uint32_t threads, uint32_t *g_hash);
extern void quark_bmw512_cpu_hash_64x(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_hash);
extern void quark_groestl512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void groestl512_cpu_init(int thr_id, uint32_t threads);
extern void groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_skein512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void qubit_luffa512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_outputHash);
extern void x11_cubehash512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_init(int thr_id, uint32_t threads);
extern void xevan_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_hamsi512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_sha512_cpu_hash_64(int thr_id, int threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void xevan_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t *resNonce, uint64_t target);
extern void xevan_groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64_A(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void quark_blake512_cpu_hash_128(int thr_id, uint32_t threads, uint32_t *d_outputHash);
extern void quark_groestl512_cpu_hash_128(int thr_id, uint32_t threads,  uint32_t *d_hash);
extern void x11_luffa512_cpu_hash_128(int thr_id, uint32_t threads,uint32_t *d_hash);



// X17 CPU Hash (Validation)
extern "C" void x17hash(void *output, const void *input)
{
	uint32_t _ALIGN(64) hash[32]; // 128 bytes required
	const int dataLen = 128;
//return;
	sph_blake512_context     ctx_blake;
	sph_bmw512_context       ctx_bmw;
	sph_groestl512_context   ctx_groestl;
	sph_skein512_context     ctx_skein;
	sph_jh512_context        ctx_jh;
	sph_keccak512_context    ctx_keccak;
	sph_luffa512_context     ctx_luffa;
	sph_cubehash512_context  ctx_cubehash;
	sph_shavite512_context   ctx_shavite;
	sph_simd512_context      ctx_simd;
	sph_echo512_context      ctx_echo;
	sph_hamsi512_context     ctx_hamsi;
	sph_fugue512_context     ctx_fugue;
	sph_shabal512_context    ctx_shabal;
	sph_whirlpool_context    ctx_whirlpool;
	sph_sha512_context       ctx_sha512;
	sph_haval256_5_context   ctx_haval;

//print_hash(input,20);
	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, input, 80);
	sph_blake512_close(&ctx_blake, hash);
//print_hash(hash,32);
	memset(&hash[16], 0, 64);

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);
//print_hash(hash,32);

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);
//print_hash(hash,32);
	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);
//print_hash(hash,32);
	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);
//print_hash(hash,32);
	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);
//print_hash(hash,32);
	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);
//print_hash(hash,32);
	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);
//print_hash(hash,32);
	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);
//print_hash(hash,32);

	memset(&hash[8], 0, dataLen - 32);

	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, hash, dataLen);
	sph_blake512_close(&ctx_blake, hash);

//print_hash(hash,32);

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);

	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);

	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);

	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);

	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);

	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);

	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);

	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);

	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);

	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);

	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);

	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);

//print_hash(hash,32);
	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);
//print_hash(hash,8);
	memcpy(output, hash, 32);
}

static bool init[MAX_GPUS] = { 0 };


void print_hash(unsigned int *data,int size){
for(int i=0;i<size;i++)
        gpulog(LOG_WARNING, 0,"%x ",data[i]);
gpulog(LOG_WARNING, 0,"-------------");
}


extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done){

	int dev_id = device_map[thr_id];

	uint32_t *pdata = work->data;
	uint32_t *ptarget = work->target;
	const uint32_t first_nonce = pdata[19];
/*
	uint32_t default_throughput = 1<<20;
	
	if (strstr(device_name[dev_id], "GTX 970")) default_throughput+=256*256*6;
	if (strstr(device_name[dev_id], "GTX 980")) default_throughput =1<<22;
	
	uint32_t throughput = cuda_default_throughput(thr_id, default_throughput); // 19=256*256*8;
*/
	uint32_t default_throughput;
	if(device_sm[dev_id]<=500) default_throughput = 1<<20;
	else if(device_sm[dev_id]<=520) default_throughput = 1<<21;
	else if(device_sm[dev_id]>520) default_throughput = (1<<22) + (1<<21);
	
	if((strstr(device_name[dev_id], "1070")))default_throughput = 1<<20;
	if((strstr(device_name[dev_id], "1080")))default_throughput = 1<<20;
	
	uint32_t throughput = cuda_default_throughput(thr_id, default_throughput); // 19=256*256*8;
	if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);

	throughput&=0xFFFFFF70; //multiples of 128 due to simd_echo kernel

	if (opt_benchmark)
		((uint32_t*)ptarget)[7] = 0xff;

	gpulog(LOG_INFO,thr_id,"target %x %x %x",ptarget[5], ptarget[6], ptarget[7]);
        gpulog(LOG_INFO,thr_id,"target %llx",*(uint64_t*)&ptarget[6]);

	if (!init[thr_id])
	{
		cudaSetDevice(device_map[thr_id]);
		if (opt_cudaschedule == -1 && gpu_threads == 1) {
			cudaDeviceReset();
			// reduce cpu usage
			cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
//			cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);
		}
		gpulog(LOG_INFO,thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);

//		x11_simd_echo_512_cpu_init(thr_id, throughput);
		x15_whirlpool_cpu_init(thr_id, throughput, 0);
		groestl512_cpu_init(thr_id, throughput);
		x11_simd512_cpu_init(thr_id, throughput);
//for(;;);
		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 8 * sizeof(uint64_t) * throughput));
		CUDA_SAFE_CALL(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t)));
		h_resNonce[thr_id] = (uint32_t*) malloc(NBN  * 8 * sizeof(uint32_t));
		if(h_resNonce[thr_id] == NULL){
			gpulog(LOG_ERR,thr_id,"Host memory allocation failed");
			exit(EXIT_FAILURE);
		}		
		init[thr_id] = true;
	}

	uint32_t _ALIGN(64) endiandata[20];
	for (int k=0; k < 20; k++)
		be32enc(&endiandata[k], pdata[k]);
//		endiandata[k]=0;
//	print_hash(endiandata,20);
	quark_blake512_cpu_setBlock_80(thr_id, endiandata);
	cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));
//	x11_simd512_cpu_init(thr_id, throughput);
//	for(;;);
	do {
		// Hash with CUDA


		quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]);//A
		quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);

		quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);
		quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);//A //fast
//		keccak_xevan_cpu_hash_64_A(thr_id, throughput,  d_hash[thr_id]);//A

//cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));
//		x11_luffa512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //P
//cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
//print_hash(h_resNonce[thr_id],16);
//cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));

		x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A
//cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
//print_hash(h_resNonce[thr_id],16);
//for(;;);

		x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //A 256
		xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//P slow r2
                x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);  //A slow r3

//                cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));


//		xevan_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //slow r1

//                cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
//		print_hash(h_resNonce[thr_id],16);


  //              cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));

		x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//A


//                cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
 //               print_hash(h_resNonce[thr_id],16);

//for(;;);

                x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast ++
		x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //opt2
		xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast
		xevan_haval512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast


//		xevan_blake512_cpu_hash_64(thr_id, throughput,  d_hash[thr_id]);//BAD
quark_blake512_cpu_hash_128(thr_id, throughput,  d_hash[thr_id]);//BAD

//
                quark_bmw512_cpu_hash_64x(thr_id, throughput, NULL, d_hash[thr_id]);
//                xevan_groestl512_cpu_hash(thr_id, throughput, d_hash[thr_id]);
quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);

//                xevan_skein512(thr_id, throughput, d_hash[thr_id]);
                quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);

                quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);
//                keccak_xevan_cpu_hash_64_A(thr_id, throughput,  d_hash[thr_id]);
//                x11_luffa512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A

                x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//move to shared
                x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); 

//                xevan_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

                x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);







/*
for(int i = 10000;i< 10016;i++){
                cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][16*i], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
print_hash(h_resNonce[thr_id],8);
}
		for(;;);

*/
		xevan_haval512_cpu_hash_64_final(thr_id, throughput, d_hash[thr_id],d_resNonce[thr_id],*(uint64_t*)&ptarget[6]);

		cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost);

		if (h_resNonce[thr_id][0] != UINT32_MAX){
			const uint32_t Htarg = ptarget[7];
			const uint32_t startNounce = pdata[19];
			uint32_t vhash64[8];
			be32enc(&endiandata[19], startNounce + h_resNonce[thr_id][0]);
			x17hash(vhash64, endiandata);
//			*hashes_done = pdata[19] - first_nonce + throughput + 1;
//			pdata[19] = startNounce + h_resNonce[thr_id][0];
			gpulog(LOG_WARNING, 0,"NONCE FOUND ");
//			return 1;
			if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
				int res = 1;
				*hashes_done = pdata[19] - first_nonce + throughput + 1;
				work_set_target_ratio(work, vhash64);
				pdata[19] = startNounce + h_resNonce[thr_id][0];
				if (h_resNonce[thr_id][1] != UINT32_MAX) {
					pdata[21] = startNounce+h_resNonce[thr_id][1];
					if(!opt_quiet)
						gpulog(LOG_BLUE,dev_id,"Found 2nd nonce: %08x", pdata[21]);
					be32enc(&endiandata[19], pdata[21]);
					x17hash(vhash64, endiandata);
					if (bn_hash_target_ratio(vhash64, ptarget) > work->shareratio[0]){
						work_set_target_ratio(work, vhash64);
						xchg(pdata[19],pdata[21]);
					}
					res++;
				}
				return res;
			}
			else {
				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", h_resNonce[thr_id][0]);
				cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));				
			}
		}

		pdata[19] += throughput;
	} while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > (uint64_t)throughput + pdata[19]));

	*hashes_done = pdata[19] - first_nonce + 1;

	return 0;
}

// cleanup
extern "C" void free_x17(int thr_id)
{
	if (!init[thr_id])
		return;

	cudaDeviceSynchronize();

	free(h_resNonce[thr_id]);
	cudaFree(d_resNonce[thr_id]);
	cudaFree(d_hash[thr_id]);

	x11_simd_echo_512_cpu_free(thr_id);
	x15_whirlpool_cpu_free(thr_id);
	cudaDeviceSynchronize();
	init[thr_id] = false;
}

[sickofscamcoins]

http://paste.ubuntu.com/25680610/

this is results of output, replacing x17.cu and recompile with one posted palgin.

Thank you, will check it!

http://paste.ubuntu.com/25680643/

no escape characters this time

[palgin]

UPDATE: Hashing is exactly the same as on Win, so it comes out to be very weird problem. The one thing I suspect is ulong and uint compiler handling difference on Win and on *nix, for example, ulong will be equal to uint64_t on 64-bit build, but different on 32-bit build. Will look into this way, maybe will come up with something.

[sickofscamcoins]

welp i tried for my cookie anyways.  Anything else I can do to help lmk

[Nesp]

Any binaries?

[anorganix]

Any binaries?

Yes, we have 0 and 1.

[palgin]

Next part of check, now GPU part. All the same, x17.cu is following:

Code:

/**
 * X17 algorithm (X15 + sha512 + haval256)
 */

extern "C" {
#include "sph/sph_blake.h"
#include "sph/sph_bmw.h"
#include "sph/sph_groestl.h"
#include "sph/sph_skein.h"
#include "sph/sph_jh.h"
#include "sph/sph_keccak.h"

#include "sph/sph_luffa.h"
#include "sph/sph_cubehash.h"
#include "sph/sph_shavite.h"
#include "sph/sph_simd.h"
#include "sph/sph_echo.h"

#include "sph/sph_hamsi.h"
#include "sph/sph_fugue.h"

#include "sph/sph_shabal.h"
#include "sph/sph_whirlpool.h"

#include "sph/sph_sha2.h"
#include "sph/sph_haval.h"
}

#include "miner.h"
#include "cuda_helper.h"
#include "x11/cuda_x11.h"

#define NBN 2

// Memory for the hash functions
static uint32_t *d_hash[MAX_GPUS];
static uint32_t *d_resNonce[MAX_GPUS];
static uint32_t *h_resNonce[MAX_GPUS];

extern void x13_hamsi_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x14_shabal512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x15_whirlpool_cpu_init(int thr_id, uint32_t threads, int mode);
extern void x15_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x15_whirlpool_cpu_free(int thr_id);

extern void x17_sha512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);

extern void x17_haval256_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t* resNonce, uint64_t target);
extern void bmw256_cpu_hash_32_full(int thr_id, uint32_t threads, uint32_t *g_hash);
extern void quark_bmw512_cpu_hash_64x(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_hash);
extern void quark_groestl512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void groestl512_cpu_init(int thr_id, uint32_t threads);
extern void groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_skein512(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void qubit_luffa512_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_outputHash);
extern void x11_cubehash512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_shavite512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_echo512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x11_simd512_cpu_init(int thr_id, uint32_t threads);
extern void xevan_simd512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_hamsi512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void x13_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_whirlpool_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_sha512_cpu_hash_64(int thr_id, int threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void xevan_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void xevan_haval512_cpu_hash_64_final(int thr_id, uint32_t threads, uint32_t *d_hash, uint32_t *resNonce, uint64_t target);
extern void xevan_groestl512_cpu_hash(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void keccak_xevan_cpu_hash_64_A(int thr_id, uint32_t threads, uint32_t *d_hash);
extern void quark_blake512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t *d_nonceVector, uint32_t *d_outputHash);
extern void quark_blake512_cpu_hash_128(int thr_id, uint32_t threads, uint32_t *d_outputHash);
extern void quark_groestl512_cpu_hash_128(int thr_id, uint32_t threads,  uint32_t *d_hash);
extern void x11_luffa512_cpu_hash_128(int thr_id, uint32_t threads,uint32_t *d_hash);



// X17 CPU Hash (Validation)
extern "C" void x17hash(void *output, const void *input)
{
	uint32_t _ALIGN(64) hash[32]; // 128 bytes required
	const int dataLen = 128;
//return;
	sph_blake512_context     ctx_blake;
	sph_bmw512_context       ctx_bmw;
	sph_groestl512_context   ctx_groestl;
	sph_skein512_context     ctx_skein;
	sph_jh512_context        ctx_jh;
	sph_keccak512_context    ctx_keccak;
	sph_luffa512_context     ctx_luffa;
	sph_cubehash512_context  ctx_cubehash;
	sph_shavite512_context   ctx_shavite;
	sph_simd512_context      ctx_simd;
	sph_echo512_context      ctx_echo;
	sph_hamsi512_context     ctx_hamsi;
	sph_fugue512_context     ctx_fugue;
	sph_shabal512_context    ctx_shabal;
	sph_whirlpool_context    ctx_whirlpool;
	sph_sha512_context       ctx_sha512;
	sph_haval256_5_context   ctx_haval;

//print_hash(input,20);
	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, input, 80);
	sph_blake512_close(&ctx_blake, hash);
//print_hash(hash,32);
	memset(&hash[16], 0, 64);

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);
//print_hash(hash,32);

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);
//print_hash(hash,32);
	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);
//print_hash(hash,32);
	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);
//print_hash(hash,32);
	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);
//print_hash(hash,32);
	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);
//print_hash(hash,32);
	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);
//print_hash(hash,32);
//for(int i=0;i<32;i++)hash[i]=0;
	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);
//print_hash(hash,32);
	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);
//print_hash(hash,32);

	memset(&hash[8], 0, dataLen - 32);

	sph_blake512_init(&ctx_blake);
	sph_blake512(&ctx_blake, hash, dataLen);
	sph_blake512_close(&ctx_blake, hash);

//print_hash(hash,32);

	sph_bmw512_init(&ctx_bmw);
	sph_bmw512(&ctx_bmw, hash, dataLen);
	sph_bmw512_close(&ctx_bmw, hash);

	sph_groestl512_init(&ctx_groestl);
	sph_groestl512(&ctx_groestl, hash, dataLen);
	sph_groestl512_close(&ctx_groestl, hash);

	sph_skein512_init(&ctx_skein);
	sph_skein512(&ctx_skein, hash, dataLen);
	sph_skein512_close(&ctx_skein, hash);

	sph_jh512_init(&ctx_jh);
	sph_jh512(&ctx_jh, hash, dataLen);
	sph_jh512_close(&ctx_jh, hash);

	sph_keccak512_init(&ctx_keccak);
	sph_keccak512(&ctx_keccak, hash, dataLen);
	sph_keccak512_close(&ctx_keccak, hash);

	sph_luffa512_init(&ctx_luffa);
	sph_luffa512(&ctx_luffa, hash, dataLen);
	sph_luffa512_close(&ctx_luffa, hash);

	sph_cubehash512_init(&ctx_cubehash);
	sph_cubehash512(&ctx_cubehash, hash, dataLen);
	sph_cubehash512_close(&ctx_cubehash, hash);

	sph_shavite512_init(&ctx_shavite);
	sph_shavite512(&ctx_shavite, hash, dataLen);
	sph_shavite512_close(&ctx_shavite, hash);

	sph_simd512_init(&ctx_simd);
	sph_simd512(&ctx_simd, hash, dataLen);
	sph_simd512_close(&ctx_simd, hash);

	sph_echo512_init(&ctx_echo);
	sph_echo512(&ctx_echo, hash, dataLen);
	sph_echo512_close(&ctx_echo, hash);

	sph_hamsi512_init(&ctx_hamsi);
	sph_hamsi512(&ctx_hamsi, hash, dataLen);
	sph_hamsi512_close(&ctx_hamsi, hash);

	sph_fugue512_init(&ctx_fugue);
	sph_fugue512(&ctx_fugue, hash, dataLen);
	sph_fugue512_close(&ctx_fugue, hash);

	sph_shabal512_init(&ctx_shabal);
	sph_shabal512(&ctx_shabal, hash, dataLen);
	sph_shabal512_close(&ctx_shabal, hash);

	sph_whirlpool_init(&ctx_whirlpool);
	sph_whirlpool(&ctx_whirlpool, hash, dataLen);
	sph_whirlpool_close(&ctx_whirlpool, hash);

	sph_sha512_init(&ctx_sha512);
	sph_sha512(&ctx_sha512,(const void*) hash, dataLen);
	sph_sha512_close(&ctx_sha512,(void*) hash);

//print_hash(hash,32);
	sph_haval256_5_init(&ctx_haval);
	sph_haval256_5(&ctx_haval,(const void*) hash, dataLen);
	sph_haval256_5_close(&ctx_haval, hash);
//print_hash(hash,8);
	memcpy(output, hash, 32);
}

static bool init[MAX_GPUS] = { 0 };


void print_hash(unsigned int *data,int size){
for(int i=0;i<size;i++)
        gpulog(LOG_WARNING, 0,"%x ",data[i]);
gpulog(LOG_WARNING, 0,"-------------");
}


extern "C" int scanhash_x17(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done){

	int dev_id = device_map[thr_id];

	uint32_t *pdata = work->data;
	uint32_t *ptarget = work->target;
	const uint32_t first_nonce = pdata[19];
/*
	uint32_t default_throughput = 1<<20;
	
	if (strstr(device_name[dev_id], "GTX 970")) default_throughput+=256*256*6;
	if (strstr(device_name[dev_id], "GTX 980")) default_throughput =1<<22;
	
	uint32_t throughput = cuda_default_throughput(thr_id, default_throughput); // 19=256*256*8;
*/
	uint32_t default_throughput;
	if(device_sm[dev_id]<=500) default_throughput = 1<<20;
	else if(device_sm[dev_id]<=520) default_throughput = 1<<21;
	else if(device_sm[dev_id]>520) default_throughput = (1<<22) + (1<<21);
	
	if((strstr(device_name[dev_id], "1070")))default_throughput = 1<<20;
	if((strstr(device_name[dev_id], "1080")))default_throughput = 1<<20;
	
	uint32_t throughput = cuda_default_throughput(thr_id, default_throughput); // 19=256*256*8;
	if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);

	throughput&=0xFFFFFF70; //multiples of 128 due to simd_echo kernel

	if (opt_benchmark)
		((uint32_t*)ptarget)[7] = 0xff;

	gpulog(LOG_INFO,thr_id,"target %x %x %x",ptarget[5], ptarget[6], ptarget[7]);
        gpulog(LOG_INFO,thr_id,"target %llx",*(uint64_t*)&ptarget[6]);

	if (!init[thr_id])
	{
		cudaSetDevice(device_map[thr_id]);
		if (opt_cudaschedule == -1 && gpu_threads == 1) {
			cudaDeviceReset();
			// reduce cpu usage
			cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
//			cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);
		}
		gpulog(LOG_INFO,thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);

//		x11_simd_echo_512_cpu_init(thr_id, throughput);
		x15_whirlpool_cpu_init(thr_id, throughput, 0);
		groestl512_cpu_init(thr_id, throughput);
		x11_simd512_cpu_init(thr_id, throughput);
//for(;;);
		CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], 8 * sizeof(uint64_t) * throughput));
		CUDA_SAFE_CALL(cudaMalloc(&d_resNonce[thr_id], NBN * sizeof(uint32_t)));
		h_resNonce[thr_id] = (uint32_t*) malloc(NBN  * 8 * sizeof(uint32_t));
		if(h_resNonce[thr_id] == NULL){
			gpulog(LOG_ERR,thr_id,"Host memory allocation failed");
			exit(EXIT_FAILURE);
		}		
		init[thr_id] = true;
	}

	uint32_t _ALIGN(64) endiandata[20];
	for (int k=0; k < 20; k++)
		be32enc(&endiandata[k], pdata[k]);
//		endiandata[k]=0;
//	print_hash(endiandata,20);
	quark_blake512_cpu_setBlock_80(thr_id, endiandata);
	cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));
//	x11_simd512_cpu_init(thr_id, throughput);
//	for(;;);
	do {
		// Hash with CUDA


		quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]);//A

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id],16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);//A //fast

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));
//		keccak_xevan_cpu_hash_64_A(thr_id, throughput,  d_hash[thr_id]);//A

//cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));
//		x11_luffa512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //P
//cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
//print_hash(h_resNonce[thr_id],16);
//cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));

		x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

//cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
//print_hash(h_resNonce[thr_id],16);
//for(;;);

		x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //A 256

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//P slow r2

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

                x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);  //A slow r3

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));


//                cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));


//		xevan_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //slow r1

//                cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
//		print_hash(h_resNonce[thr_id],16);


  //              cudaMemset(d_hash[thr_id], 0x00, 16*sizeof(uint32_t));

		x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//A

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));



//                cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
 //               print_hash(h_resNonce[thr_id],16);

//for(;;);

                x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast ++

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //opt2

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

		xevan_haval512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); //fast

		cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
		print_hash(h_resNonce[thr_id], 16);
		cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));



//		xevan_blake512_cpu_hash_64(thr_id, throughput,  d_hash[thr_id]);//BAD
quark_blake512_cpu_hash_128(thr_id, throughput,  d_hash[thr_id]);//BAD

cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
print_hash(h_resNonce[thr_id], 16);
cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));


//
                quark_bmw512_cpu_hash_64x(thr_id, throughput, NULL, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

//                xevan_groestl512_cpu_hash(thr_id, throughput, d_hash[thr_id]);
quark_groestl512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);

cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
print_hash(h_resNonce[thr_id], 16);


//                xevan_skein512(thr_id, throughput, d_hash[thr_id]);
                quark_skein512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));


                quark_jh512_cpu_hash_64(thr_id, throughput, NULL, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

//                keccak_xevan_cpu_hash_64_A(thr_id, throughput,  d_hash[thr_id]);
//                x11_luffa512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                x11_luffa512_cpu_hash_128(thr_id, throughput, d_hash[thr_id]);//A

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));


                x11_cubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

                xevan_shavite512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);//move to shared

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

                x11_simd512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]); 

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));


//                xevan_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);
                x11_echo512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));


                x13_hamsi512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

                x13_fugue512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

                x14_shabal512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

                xevan_whirlpool_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));

                xevan_sha512_cpu_hash_64(thr_id, throughput, d_hash[thr_id]);

				cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][0], 16 * sizeof(uint32_t), cudaMemcpyDeviceToHost);
				print_hash(h_resNonce[thr_id], 16);
				cudaMemset(d_hash[thr_id], 0x00, 16 * sizeof(uint32_t));








/*
for(int i = 10000;i< 10016;i++){
                cudaMemcpy(h_resNonce[thr_id], &d_hash[thr_id][16*i], 16*sizeof(uint32_t), cudaMemcpyDeviceToHost);
print_hash(h_resNonce[thr_id],8);
}
		for(;;);

*/
		xevan_haval512_cpu_hash_64_final(thr_id, throughput, d_hash[thr_id],d_resNonce[thr_id],*(uint64_t*)&ptarget[6]);

		cudaMemcpy(h_resNonce[thr_id], d_resNonce[thr_id], NBN*sizeof(uint32_t), cudaMemcpyDeviceToHost);

		print_hash(h_resNonce[thr_id], 16);

		if (h_resNonce[thr_id][0] != UINT32_MAX){
			const uint32_t Htarg = ptarget[7];
			const uint32_t startNounce = pdata[19];
			uint32_t vhash64[8];
			be32enc(&endiandata[19], startNounce + h_resNonce[thr_id][0]);
			x17hash(vhash64, endiandata);
//			*hashes_done = pdata[19] - first_nonce + throughput + 1;
//			pdata[19] = startNounce + h_resNonce[thr_id][0];
			gpulog(LOG_WARNING, 0,"NONCE FOUND ");
//			return 1;
			if (vhash64[7] <= Htarg && fulltest(vhash64, ptarget)) {
				int res = 1;
				*hashes_done = pdata[19] - first_nonce + throughput + 1;
				work_set_target_ratio(work, vhash64);
				pdata[19] = startNounce + h_resNonce[thr_id][0];
				if (h_resNonce[thr_id][1] != UINT32_MAX) {
					pdata[21] = startNounce+h_resNonce[thr_id][1];
					if(!opt_quiet)
						gpulog(LOG_BLUE,dev_id,"Found 2nd nonce: %08x", pdata[21]);
					be32enc(&endiandata[19], pdata[21]);
					x17hash(vhash64, endiandata);
					if (bn_hash_target_ratio(vhash64, ptarget) > work->shareratio[0]){
						work_set_target_ratio(work, vhash64);
						xchg(pdata[19],pdata[21]);
					}
					res++;
				}
				return res;
			}
			else {
				gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", h_resNonce[thr_id][0]);
				cudaMemset(d_resNonce[thr_id], 0xff, NBN*sizeof(uint32_t));				
			}
		}

		pdata[19] += throughput;
	} while (!work_restart[thr_id].restart && ((uint64_t)max_nonce > (uint64_t)throughput + pdata[19]));

	*hashes_done = pdata[19] - first_nonce + 1;

	return 0;
}

// cleanup
extern "C" void free_x17(int thr_id)
{
	if (!init[thr_id])
		return;

	cudaDeviceSynchronize();

	free(h_resNonce[thr_id]);
	cudaFree(d_resNonce[thr_id]);
	cudaFree(d_hash[thr_id]);

	x11_simd_echo_512_cpu_free(thr_id);
	x15_whirlpool_cpu_free(thr_id);
	cudaDeviceSynchronize();
	init[thr_id] = false;
}

Thanks for everybody who helps in investigation.

UPDATE: think I've found the bug, it's in xevan_haval512_cpu_hash_64_final function, so h_resNonce[0] and h_resNonce[1] are always random on Win with exactly the same input data... Need confirmation that behaviour is different on *nix, of course. CUDA memcpy bug?

[Nesp]

Any binaries?

Yes, we have 0 and 1.

Hahah, nice one