Description: Lyra2v2 improvments from PR #56
 * remove unnecessary IFDEF from cuda_cubehash256.cu,
 * remove LROT from cuda_cubehash256.cu as cuda_helper.h defines ROTL32 for
   the exactly same purpose,
 * use __ldg in cuda_skein256.cu (perf.),
 * use uint64_t in cuda_sekin256.cu as it makes lyra2v2 ~1% faster on my GTX
   1080Ti (perf.),
 * remove unneeded old kernel from cuda_skein256.cu,
 * remove Round_8_512v35_final, which skips unnecessary computations, because
   compiler can trivially perform this optimization once the round function
   is inlined.
 .
 ccminer-tpruvot (2.2.4-3) unstable; urgency=medium
 .
   * Support CUDA 9.1 in debian/rules for Debian Buster and Ubuntu Bionic.
   * Enable libssl 1.1 build, as 1.0 is conflicting with libcurl4-openssl-dev
     on Ubuntu Bionic.
Author: Adam Cecile <acecile@le-vert.net>

---
The information above should follow the Patch Tagging Guidelines, please
checkout http://dep.debian.net/deps/dep3/ to learn about the format. Here
are templates for supplementary fields that you might want to add:

Origin: upstream, https://github.com/tpruvot/ccminer/pull/56
Forwarded: not-needed
Reviewed-By: Shunsuke Shimizu <grafi@grafi.jp>
Last-Update: 2018-03-27

--- ccminer-tpruvot-2.2.4.orig/Algo256/cuda_cubehash256.cu
+++ ccminer-tpruvot-2.2.4/Algo256/cuda_cubehash256.cu
@@ -8,19 +8,18 @@
 #define __CUDA_ARCH__ 520
 #endif
 
-#if __CUDA_ARCH__ < 350
-#define LROT(x,bits) ((x << bits) | (x >> (32 - bits)))
-#else
-#define LROT(x, bits) __funnelshift_l(x, x, bits)
-#endif
-
 #define TPB35 576
 #define TPB50 1024
 
-#define ROTATEUPWARDS7(a)  LROT(a,7)
-#define ROTATEUPWARDS11(a) LROT(a,11)
+static __device__ __forceinline__ uint32_t ROTATEUPWARDS7(uint32_t a) {
+	return ROTL32(a, 7);
+}
+static __device__ __forceinline__ uint32_t ROTATEUPWARDS11(uint32_t a) {
+	return ROTL32(a, 11);
+}
 
-__device__ __forceinline__ void rrounds(uint32_t x[2][2][2][2][2])
+static __device__ __forceinline__
+void rrounds(uint32_t x[2][2][2][2][2])
 {
 	int r;
 
@@ -246,7 +245,7 @@ __device__ __forceinline__ void rrounds(
 	}
 }
 
-__device__ __forceinline__
+static __device__ __forceinline__
 void Final(uint32_t x[2][2][2][2][2], uint32_t *hashval)
 {
 	/* "the integer 1 is xored into the last state word x_11111" */
@@ -276,21 +275,12 @@ void cubehash256_gpu_hash_32(uint32_t th
 	uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
 	if (thread < threads)
 	{
-#if __CUDA_ARCH__ >= 500
 		uint2 Hash[4];
 
 		Hash[0] = __ldg(&g_hash[thread]);
 		Hash[1] = __ldg(&g_hash[thread + 1 * threads]);
 		Hash[2] = __ldg(&g_hash[thread + 2 * threads]);
 		Hash[3] = __ldg(&g_hash[thread + 3 * threads]);
-#else
-		uint32_t Hash[8];
-
-		LOHI(Hash[0], Hash[1], __ldg(&((uint64_t*)g_hash)[thread]));
-		LOHI(Hash[2], Hash[3], __ldg(&((uint64_t*)g_hash)[thread + 1 * threads]));
-		LOHI(Hash[4], Hash[5], __ldg(&((uint64_t*)g_hash)[thread + 2 * threads]));
-		LOHI(Hash[6], Hash[7], __ldg(&((uint64_t*)g_hash)[thread + 3 * threads]));
-#endif
 
 		uint32_t x[2][2][2][2][2] =
 		{
@@ -304,7 +294,6 @@ void cubehash256_gpu_hash_32(uint32_t th
 			0x15815AEB, 0x4AB6AAD6, 0x9CDAF8AF, 0xD6032C0A
 		};
 
-#if __CUDA_ARCH__ >= 500
 		x[0][0][0][0][0] ^= Hash[0].x;
 		x[0][0][0][0][1] ^= Hash[0].y;
 		x[0][0][0][1][0] ^= Hash[1].x;
@@ -313,35 +302,17 @@ void cubehash256_gpu_hash_32(uint32_t th
 		x[0][0][1][0][1] ^= Hash[2].y;
 		x[0][0][1][1][0] ^= Hash[3].x;
 		x[0][0][1][1][1] ^= Hash[3].y;
-#else
-		x[0][0][0][0][0] ^= Hash[0];
-		x[0][0][0][0][1] ^= Hash[1];
-		x[0][0][0][1][0] ^= Hash[2];
-		x[0][0][0][1][1] ^= Hash[3];
-		x[0][0][1][0][0] ^= Hash[4];
-		x[0][0][1][0][1] ^= Hash[5];
-		x[0][0][1][1][0] ^= Hash[6];
-		x[0][0][1][1][1] ^= Hash[7];
-#endif
+
 		rrounds(x);
 		x[0][0][0][0][0] ^= 0x80U;
 		rrounds(x);
 
-#if __CUDA_ARCH__ >= 500
 		Final(x, (uint32_t*)Hash);
 
 		g_hash[thread] = Hash[0];
 		g_hash[1 * threads + thread] = Hash[1];
 		g_hash[2 * threads + thread] = Hash[2];
 		g_hash[3 * threads + thread] = Hash[3];
-#else
-		Final(x, Hash);
-
-		((uint64_t*)g_hash)[thread] = ((uint64_t*)Hash)[0];
-		((uint64_t*)g_hash)[1 * threads + thread] = ((uint64_t*)Hash)[1];
-		((uint64_t*)g_hash)[2 * threads + thread] = ((uint64_t*)Hash)[2];
-		((uint64_t*)g_hash)[3 * threads + thread] = ((uint64_t*)Hash)[3];
-#endif
 	}
 }
 
--- ccminer-tpruvot-2.2.4.orig/Algo256/cuda_skein256.cu
+++ ccminer-tpruvot-2.2.4/Algo256/cuda_skein256.cu
@@ -2,24 +2,26 @@
 
 #include "cuda_helper.h"
 
-static __forceinline__ __device__
-void Round512v35(uint2 &p0, uint2 &p1, uint2 &p2, uint2 &p3, uint2 &p4, uint2 &p5, uint2 &p6, uint2 &p7,
-	const int ROT0, const int ROT1, const int ROT2, const int ROT3)
+static __device__ __forceinline__
+void Round512(uint64_t &p0, uint64_t &p1, uint64_t &p2, uint64_t &p3,
+	uint64_t &p4, uint64_t &p5, uint64_t &p6, uint64_t &p7,
+	int ROT0, int ROT1, int ROT2, int ROT3)
 {
-	p0 += p1; p1 = ROL2(p1, ROT0) ^ p0;
-	p2 += p3; p3 = ROL2(p3, ROT1) ^ p2;
-	p4 += p5; p5 = ROL2(p5, ROT2) ^ p4;
-	p6 += p7; p7 = ROL2(p7, ROT3) ^ p6;
+	p0 += p1; p1 = ROTL64(p1, ROT0) ^ p0;
+	p2 += p3; p3 = ROTL64(p3, ROT1) ^ p2;
+	p4 += p5; p5 = ROTL64(p5, ROT2) ^ p4;
+	p6 += p7; p7 = ROTL64(p7, ROT3) ^ p6;
 }
 
-__forceinline__ __device__
-void Round_8_512v35(const uint2 *const __restrict__ ks, const uint2 *const __restrict__ ts,
-	uint2 &p0, uint2 &p1, uint2 &p2, uint2 &p3, uint2 &p4, uint2 &p5, uint2 &p6, uint2 &p7, int R)
+static __device__ __forceinline__
+void Round_8_512(const uint64_t *__restrict__ ks, const uint64_t *__restrict__ ts,
+	uint64_t &p0, uint64_t &p1, uint64_t &p2, uint64_t &p3,
+	uint64_t &p4, uint64_t &p5, uint64_t &p6, uint64_t &p7, int R)
 {
-	Round512v35(p0, p1, p2, p3, p4, p5, p6, p7, 46, 36, 19, 37);
-	Round512v35(p2, p1, p4, p7, p6, p5, p0, p3, 33, 27, 14, 42);
-	Round512v35(p4, p1, p6, p3, p0, p5, p2, p7, 17, 49, 36, 39);
-	Round512v35(p6, p1, p0, p7, p2, p5, p4, p3, 44,  9, 54, 56);
+	Round512(p0, p1, p2, p3, p4, p5, p6, p7, 46, 36, 19, 37);
+	Round512(p2, p1, p4, p7, p6, p5, p0, p3, 33, 27, 14, 42);
+	Round512(p4, p1, p6, p3, p0, p5, p2, p7, 17, 49, 36, 39);
+	Round512(p6, p1, p0, p7, p2, p5, p4, p3, 44,  9, 54, 56);
 
 	p0 += ks[(R+0) % 9];
 	p1 += ks[(R+1) % 9];
@@ -28,12 +30,12 @@ void Round_8_512v35(const uint2 *const _
 	p4 += ks[(R+4) % 9];
 	p5 += ks[(R+5) % 9] + ts[(R+0) % 3];
 	p6 += ks[(R+6) % 9] + ts[(R+1) % 3];
-	p7 += ks[(R+7) % 9] + make_uint2(R, 0);
+	p7 += ks[(R+7) % 9] + R;
 
-	Round512v35(p0, p1, p2, p3, p4, p5, p6, p7, 39, 30, 34, 24);
-	Round512v35(p2, p1, p4, p7, p6, p5, p0, p3, 13, 50, 10, 17);
-	Round512v35(p4, p1, p6, p3, p0, p5, p2, p7, 25, 29, 39, 43);
-	Round512v35(p6, p1, p0, p7, p2, p5, p4, p3, 8,  35, 56, 22);
+	Round512(p0, p1, p2, p3, p4, p5, p6, p7, 39, 30, 34, 24);
+	Round512(p2, p1, p4, p7, p6, p5, p0, p3, 13, 50, 10, 17);
+	Round512(p4, p1, p6, p3, p0, p5, p2, p7, 25, 29, 39, 43);
+	Round512(p6, p1, p0, p7, p2, p5, p4, p3,  8, 35, 56, 22);
 
 	p0 += ks[(R+1) % 9];
 	p1 += ks[(R+2) % 9];
@@ -42,36 +44,7 @@ void Round_8_512v35(const uint2 *const _
 	p4 += ks[(R+5) % 9];
 	p5 += ks[(R+6) % 9] + ts[(R+1) % 3];
 	p6 += ks[(R+7) % 9] + ts[(R+2) % 3];
-	p7 += ks[(R+8) % 9] + make_uint2(R+1, 0);
-}
-
-__forceinline__ __device__
-void Round_8_512v35_final(const uint2 *const __restrict__ ks, const uint2 *const __restrict__ ts,
-	uint2 &p0, uint2 &p1, uint2 &p2, uint2 &p3, uint2 &p4, uint2 &p5, uint2 &p6, uint2 &p7)
-{
-	Round512v35(p0, p1, p2, p3, p4, p5, p6, p7, 46, 36, 19, 37);
-	Round512v35(p2, p1, p4, p7, p6, p5, p0, p3, 33, 27, 14, 42);
-	Round512v35(p4, p1, p6, p3, p0, p5, p2, p7, 17, 49, 36, 39);
-	Round512v35(p6, p1, p0, p7, p2, p5, p4, p3, 44, 9, 54, 56);
-
-	p0 += ks[8];
-	p1 += ks[0];
-	p2 += ks[1];
-	p3 += ks[2];
-	p4 += ks[3];
-	p5 += ks[4] + ts[2];
-	p6 += ks[5] + ts[0];
-	p7 += ks[6] + make_uint2(17, 0);
-
-	Round512v35(p0, p1, p2, p3, p4, p5, p6, p7, 39, 30, 34, 24);
-	Round512v35(p2, p1, p4, p7, p6, p5, p0, p3, 13, 50, 10, 17);
-	Round512v35(p4, p1, p6, p3, p0, p5, p2, p7, 25, 29, 39, 43);
-	Round512v35(p6, p1, p0, p7, p2, p5, p4, p3, 8,  35, 56, 22);
-
-	p0 += ks[0];
-	p1 += ks[1];
-	p2 += ks[2];
-	p3 += ks[3];
+	p7 += ks[(R+8) % 9] + R+1;
 }
 
 __global__ __launch_bounds__(256,3)
@@ -81,34 +54,27 @@ void skein256_gpu_hash_32(uint32_t threa
 
 	if (thread < threads)
 	{
-		const uint2 skein_ks_parity = { 0xA9FC1A22, 0x1BD11BDA };
-		const uint2 t12[6] = {
-			{ 0x20, 0 },
-			{ 0,    0xf0000000 },
-			{ 0x20, 0xf0000000 },
-			{ 0x08, 0 },
-			{ 0,    0xff000000 },
-			{ 0x08, 0xff000000 }
+		const uint64_t skein_ks_parity64 = 0x1BD11BDAA9FC1A22ull;
+
+		const uint64_t t12[6] = {
+			0x20, 0xf000000000000000ull, 0xf000000000000020ull,
+			0x08, 0xff00000000000000ull, 0xff00000000000008ull
 		};
 
-		uint2 h[9] = {
-			{ 0x2FDB3E13, 0xCCD044A1 },
-			{ 0x1A79A9EB, 0xE8359030 },
-			{ 0x4F816E6F, 0x55AEA061 },
-			{ 0xAE9B94DB, 0x2A2767A4 },
-			{ 0x74DD7683, 0xEC06025E },
-			{ 0xC4746251, 0xE7A436CD },
-			{ 0x393AD185, 0xC36FBAF9 },
-			{ 0x33EDFC13, 0x3EEDBA18 },
-			{ 0xC73A4E2A, 0xB69D3CFC }
+		uint64_t h[9] = {
+			0xCCD044A12FDB3E13ull, 0xE83590301A79A9EBull,
+			0x55AEA0614F816E6Full, 0x2A2767A4AE9B94DBull,
+			0xEC06025E74DD7683ull, 0xE7A436CDC4746251ull,
+			0xC36FBAF9393AD185ull, 0x3EEDBA1833EDFC13ull,
+			0xB69D3CFCC73A4E2Aull // skein_ks_parity64 ^ h[0..7]
 		};
-		uint2 dt0,dt1,dt2,dt3;
-		uint2 p0, p1, p2, p3, p4, p5, p6, p7;
+		uint64_t dt0, dt1, dt2, dt3;
+		uint64_t p0, p1, p2, p3, p4, p5, p6, p7;
 
-		LOHI(dt0.x,dt0.y,outputHash[thread]);
-		LOHI(dt1.x,dt1.y,outputHash[threads+thread]);
-		LOHI(dt2.x,dt2.y,outputHash[2*threads+thread]);
-		LOHI(dt3.x,dt3.y,outputHash[3*threads+thread]);
+		dt0 = __ldg(&outputHash[0 * threads + thread]);
+		dt1 = __ldg(&outputHash[1 * threads + thread]);
+		dt2 = __ldg(&outputHash[2 * threads + thread]);
+		dt3 = __ldg(&outputHash[3 * threads + thread]);
 
 		p0 = h[0] + dt0;
 		p1 = h[1] + dt1;
@@ -120,15 +86,15 @@ void skein256_gpu_hash_32(uint32_t threa
 		p7 = h[7];
 
 		// forced unroll required
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 1);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 3);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 5);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 7);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 9);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 11);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 13);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 15);
-		Round_8_512v35(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 17);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 1);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 3);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 5);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 7);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 9);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 11);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 13);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 15);
+		Round_8_512(h, t12, p0, p1, p2, p3, p4, p5, p6, p7, 17);
 
 		p0 ^= dt0;
 		p1 ^= dt1;
@@ -143,150 +109,30 @@ void skein256_gpu_hash_32(uint32_t threa
 		h[5] = p5;
 		h[6] = p6;
 		h[7] = p7;
-		h[8] = skein_ks_parity ^ h[0] ^ h[1] ^ h[2] ^ h[3] ^ h[4] ^ h[5] ^ h[6] ^ h[7];
+		h[8] = skein_ks_parity64 ^ h[0] ^ h[1] ^ h[2] ^ h[3] ^ h[4] ^ h[5] ^ h[6] ^ h[7];
 
-		const uint2 *t = t12+3;
+		const uint64_t *t = t12+3;
 		p5 += t12[3];  //p5 already equal h[5]
 		p6 += t12[4];
 
 		// forced unroll
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 1);
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 3);
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 5);
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 7);
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 9);
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 11);
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 13);
-		Round_8_512v35(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 15);
-		Round_8_512v35_final(h, t, p0, p1, p2, p3, p4, p5, p6, p7);
-
-		outputHash[thread]           = devectorize(p0);
-		outputHash[threads+thread]   = devectorize(p1);
-		outputHash[2*threads+thread] = devectorize(p2);
-		outputHash[3*threads+thread] = devectorize(p3);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 1);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 3);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 5);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 7);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 9);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 11);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 13);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 15);
+		Round_8_512(h, t, p0, p1, p2, p3, p4, p5, p6, p7, 17);
+
+		outputHash[0 * threads + thread] = p0;
+		outputHash[1 * threads + thread] = p1;
+		outputHash[2 * threads + thread] = p2;
+		outputHash[3 * threads + thread] = p3;
 	}
 }
 
-static __forceinline__ __device__
-void Round512v30(uint64_t &p0, uint64_t &p1, uint64_t &p2, uint64_t &p3,
-	uint64_t &p4, uint64_t &p5, uint64_t &p6, uint64_t &p7,
-	const int ROT0, const int ROT1, const int ROT2, const int ROT3)
-{
-	p0 += p1; p1 = ROTL64(p1, ROT0) ^ p0;
-	p2 += p3; p3 = ROTL64(p3, ROT1) ^ p2;
-	p4 += p5; p5 = ROTL64(p5, ROT2) ^ p4;
-	p6 += p7; p7 = ROTL64(p7, ROT3) ^ p6;
-}
-
-static __forceinline__ __device__
-void Round_8_512v30(uint64_t *ks, uint64_t *ts,
-	uint64_t &p0, uint64_t &p1, uint64_t &p2, uint64_t &p3,
-	uint64_t &p4, uint64_t &p5, uint64_t &p6, uint64_t &p7, int R)
-{
-	Round512v30(p0, p1, p2, p3, p4, p5, p6, p7, 46, 36, 19, 37);
-	Round512v30(p2, p1, p4, p7, p6, p5, p0, p3, 33, 27, 14, 42);
-	Round512v30(p4, p1, p6, p3, p0, p5, p2, p7, 17, 49, 36, 39);
-	Round512v30(p6, p1, p0, p7, p2, p5, p4, p3, 44,  9, 54, 56);
-
-	p0 += ks[(R+0) % 9];
-	p1 += ks[(R+1) % 9];
-	p2 += ks[(R+2) % 9];
-	p3 += ks[(R+3) % 9];
-	p4 += ks[(R+4) % 9];
-	p5 += ks[(R+5) % 9] + ts[(R+0) % 3];
-	p6 += ks[(R+6) % 9] + ts[(R+1) % 3];
-	p7 += ks[(R+7) % 9] + R;
-
-	Round512v30(p0, p1, p2, p3, p4, p5, p6, p7, 39, 30, 34, 24);
-	Round512v30(p2, p1, p4, p7, p6, p5, p0, p3, 13, 50, 10, 17);
-	Round512v30(p4, p1, p6, p3, p0, p5, p2, p7, 25, 29, 39, 43);
-	Round512v30(p6, p1, p0, p7, p2, p5, p4, p3, 8,  35, 56, 22);
-
-	p0 += ks[(R+1) % 9];
-	p1 += ks[(R+2) % 9];
-	p2 += ks[(R+3) % 9];
-	p3 += ks[(R+4) % 9];
-	p4 += ks[(R+5) % 9];
-	p5 += ks[(R+6) % 9] + ts[(R+1) % 3];
-	p6 += ks[(R+7) % 9] + ts[(R+2) % 3];
-	p7 += ks[(R+8) % 9] + R+1;
-}
-
-#define skein_ks_parity64 0x1BD11BDAA9FC1A22ull
-#include <stdio.h>
-__global__  __launch_bounds__(256, 3)
-void skein256_gpu_hash_32_v30(uint32_t threads, uint32_t startNounce, uint64_t *outputHash)
-{
-	const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
-	if (thread < threads)
-	{
-		uint64_t h[12] = { // SKEIN_IV512_256
-			0xCCD044A12FDB3E13, 0xE83590301A79A9EB,
-			0x55AEA0614F816E6F, 0x2A2767A4AE9B94DB,
-			0xEC06025E74DD7683, 0xE7A436CDC4746251,
-			0xC36FBAF9393AD185, 0x3EEDBA1833EDFC13,
-			0xb69d3cfcc73a4e2a, // skein_ks_parity64 ^ h[0..7]
-			0x20, 0xf000000000000000, 0xf000000000000020 // t0..2
-		};
-		uint64_t dt0 = outputHash[thread];
-		uint64_t dt1 = outputHash[threads   + thread];
-		uint64_t dt2 = outputHash[threads*2 + thread];
-		uint64_t dt3 = outputHash[threads*3 + thread];
-
-		uint64_t *t = &h[9];
-		uint64_t p0 = h[0] + dt0;
-		uint64_t p1 = h[1] + dt1;
-		uint64_t p2 = h[2] + dt2;
-		uint64_t p3 = h[3] + dt3;
-		uint64_t p4 = h[4];
-		uint64_t p5 = h[5] + t[0];
-		uint64_t p6 = h[6] + t[1];
-		uint64_t p7 = h[7];
-
-		#pragma unroll 9
-		for (int i = 1; i<19; i += 2) {
-			Round_8_512v30(h, t, p0, p1, p2, p3, p4, p5, p6, p7, i);
-		}
-
-		p0 ^= dt0;
-		p1 ^= dt1;
-		p2 ^= dt2;
-		p3 ^= dt3;
-
-		h[0] = p0;
-		h[1] = p1;
-		h[2] = p2;
-		h[3] = p3;
-		h[4] = p4;
-		h[5] = p5;
-		h[6] = p6;
-		h[7] = p7;
-		h[8] = skein_ks_parity64;
-
-		#pragma unroll 8
-		for (int i = 0; i<8; i++) {
-			h[8] ^= h[i];
-		}
-
-		t[0] = 0x08;
-		t[1] = 0xff00000000000000;
-		t[2] = 0xff00000000000008;
-
-		p5 += t[0];  //p5 already equal h[5]
-		p6 += t[1];
-
-		#pragma unroll 9
-		for (int i = 1; i<19; i += 2) {
-			Round_8_512v30(h, t, p0, p1, p2, p3, p4, p5, p6, p7, i);
-		}
-
-		outputHash[thread] = p0;
-		outputHash[threads   + thread] = p1;
-		outputHash[threads*2 + thread] = p2;
-		outputHash[threads*3 + thread] = p3;
-	} //thread
-}
-
 __host__
 void skein256_cpu_init(int thr_id, uint32_t threads)
 {
@@ -302,11 +148,7 @@ void skein256_cpu_hash_32(int thr_id, ui
 	dim3 grid((threads + threadsperblock - 1) / threadsperblock);
 	dim3 block(threadsperblock);
 
-	// only 1kH/s perf change between kernels on a 960...
-	if (device_sm[dev_id] > 300 && cuda_arch[dev_id] > 300)
-		skein256_gpu_hash_32<<<grid, block>>>(threads, startNounce, d_outputHash);
-	else
-		skein256_gpu_hash_32_v30<<<grid, block>>>(threads, startNounce, d_outputHash);
+	skein256_gpu_hash_32<<<grid, block>>>(threads, startNounce, d_outputHash);
 
 	MyStreamSynchronize(NULL, order, thr_id);
 }
