1.5.24

backend updates to bfgminer and cgminer
fabulouspanda · Feb 16, 2015 · 981672b · 981672b
1 parent cde6246
commit 981672b
Show file tree

Hide file tree

Showing 49 changed files with 567 additions and 61 deletions.
diff --git a/MacMiner/MacMiner-Info.plist b/MacMiner/MacMiner-Info.plist
@@ -17,11 +17,11 @@
 	<key>CFBundlePackageType</key>
 	<string>APPL</string>
 	<key>CFBundleShortVersionString</key>
-	<string>1.5.22</string>
+	<string>1.5.24</string>
 	<key>CFBundleSignature</key>
 	<string>????</string>
 	<key>CFBundleVersion</key>
-	<string>1522</string>
+	<string>1524</string>
 	<key>LSApplicationCategoryType</key>
 	<string>public.app-category.utilities</string>
 	<key>LSMinimumSystemVersion</key>

diff --git a/bfgminer/bin/bfgminer b/bfgminer/bin/bfgminer
diff --git a/bfgminer/bin/bfgminer-rpc b/bfgminer/bin/bfgminer-rpc
diff --git a/bfgminer/bin/bitforce-firmware-flash b/bfgminer/bin/bitforce-firmware-flash
diff --git a/bfgminer/lib/libbase58.0.dylib b/bfgminer/lib/libbase58.0.dylib
diff --git a/bfgminer/lib/libbase58.la b/bfgminer/lib/libbase58.la
@@ -25,7 +25,7 @@ weak_library_names=''
 # Version information for libbase58.
 current=0
 age=0
-revision=0
+revision=1
 
 # Is this an already installed library?
 installed=yes

diff --git a/bfgminer/lib/libblkmaker-0.1.0.dylib b/bfgminer/lib/libblkmaker-0.1.0.dylib
diff --git a/bfgminer/lib/libblkmaker-0.1.6.dylib b/bfgminer/lib/libblkmaker-0.1.6.dylib
diff --git a/bfgminer/lib/libblkmaker-0.1.la b/bfgminer/lib/libblkmaker-0.1.la
@@ -17,7 +17,7 @@ old_library=''
 inherited_linker_flags=' '
 
 # Libraries that this one depends upon.
-dependency_libs=' -L/Users/admin/Downloads/bfgminer-4.10.0/libbase58/.libs /Applications/MacMiner.app/Contents/Resources/bfgminer/lib/libbase58.la'
+dependency_libs=' -L/Users/admin/Downloads/bfgminer-5.0.0/libbase58/.libs /Applications/MacMiner.app/Contents/Resources/bfgminer/lib/libbase58.la'
 
 # Names of additional weak libraries provided by this library
 weak_library_names=''

diff --git a/bfgminer/lib/libblkmaker_jansson-0.1.0.dylib b/bfgminer/lib/libblkmaker_jansson-0.1.0.dylib
diff --git a/bfgminer/lib/libblkmaker_jansson-0.1.6.dylib b/bfgminer/lib/libblkmaker_jansson-0.1.6.dylib
diff --git a/bfgminer/lib/libblkmaker_jansson-0.1.la b/bfgminer/lib/libblkmaker_jansson-0.1.la
@@ -17,7 +17,7 @@ old_library=''
 inherited_linker_flags=' '
 
 # Libraries that this one depends upon.
-dependency_libs=' -L.libs /Applications/MacMiner.app/Contents/Resources/bfgminer/lib/libblkmaker-0.1.la -L/Users/admin/Downloads/bfgminer-4.10.0/libbase58/.libs /Applications/MacMiner.app/Contents/Resources/bfgminer/lib/libbase58.la -L/Applications/MacMiner.app/Contents/Resources/jansson/lib /Applications/MacMiner.app/Contents/Resources/jansson/lib/libjansson.la'
+dependency_libs=' -L.libs /Applications/MacMiner.app/Contents/Resources/bfgminer/lib/libblkmaker-0.1.la -L/Users/admin/Downloads/bfgminer-5.0.0/libbase58/.libs /Applications/MacMiner.app/Contents/Resources/bfgminer/lib/libbase58.la -L/Applications/MacMiner.app/Contents/Resources/jansson/lib /Applications/MacMiner.app/Contents/Resources/jansson/lib/libjansson.la'
 
 # Names of additional weak libraries provided by this library
 weak_library_names=''

diff --git a/bfgminer/lib/libevent-2.0.5.dylib b/bfgminer/lib/libevent-2.0.5.dylib
diff --git a/bfgminer/lib/libevent.a b/bfgminer/lib/libevent.a
diff --git a/bfgminer/lib/libevent_core-2.0.5.dylib b/bfgminer/lib/libevent_core-2.0.5.dylib
diff --git a/bfgminer/lib/libevent_core.a b/bfgminer/lib/libevent_core.a
diff --git a/bfgminer/lib/libevent_extra-2.0.5.dylib b/bfgminer/lib/libevent_extra-2.0.5.dylib
diff --git a/bfgminer/lib/libevent_extra.a b/bfgminer/lib/libevent_extra.a
diff --git a/bfgminer/lib/libevent_openssl-2.0.5.dylib b/bfgminer/lib/libevent_openssl-2.0.5.dylib
diff --git a/bfgminer/lib/libevent_openssl.a b/bfgminer/lib/libevent_openssl.a
diff --git a/bfgminer/lib/libevent_pthreads-2.0.5.dylib b/bfgminer/lib/libevent_pthreads-2.0.5.dylib
diff --git a/bfgminer/lib/libevent_pthreads.a b/bfgminer/lib/libevent_pthreads.a
diff --git a/bfgminer/lib/libjansson.4.dylib b/bfgminer/lib/libjansson.4.dylib
diff --git a/bfgminer/lib/libjansson.a b/bfgminer/lib/libjansson.a
diff --git a/bfgminer/lib/libmicrohttpd.10.dylib b/bfgminer/lib/libmicrohttpd.10.dylib
diff --git a/bfgminer/lib/libmicrohttpd.a b/bfgminer/lib/libmicrohttpd.a
diff --git a/bfgminer/lib/libusb-1.0.2.dylib b/bfgminer/lib/libusb-1.0.2.dylib
diff --git a/bfgminer/lib/libusb-1.0.a b/bfgminer/lib/libusb-1.0.a
diff --git a/bfgminer/lib/pkgconfig/libbase58.pc b/bfgminer/lib/pkgconfig/libbase58.pc
@@ -5,6 +5,6 @@ includedir=${prefix}/include
 
 Name: libbase58
 Description: Library for Bitcoin's base58 encoding.
-Version: 0.1
+Version: 0.1.3
 Libs: -L${libdir} -lbase58
 Cflags: -I${includedir}
diff --git a/bfgminer/share/bfgminer/opencl/diablo.cl b/bfgminer/share/bfgminer/opencl/diablo.cl
@@ -17,7 +17,7 @@
  *  along with this program.  If not, see <http://www.gnu.org/licenses/>.
  */
 
-// kernel-interface: diablo
+// kernel-interface: diablo SHA256d
 
 #ifdef VECTORS4
 	typedef uint4 z;

diff --git a/bfgminer/share/bfgminer/opencl/diakgcn.cl b/bfgminer/share/bfgminer/opencl/diakgcn.cl
@@ -3,7 +3,7 @@
 // Parts and / or ideas for this kernel are based upon the public-domain poclbm project, the phatk kernel by Phateus and the DiabloMiner kernel by DiabloD3.
 // The kernel was rewritten by me (Diapolo) and is still public-domain!
 
-// kernel-interface: diakgcn
+// kernel-interface: diakgcn SHA256d
 
 #ifdef VECTORS4
 	typedef uint4 u;

diff --git a/bfgminer/share/bfgminer/opencl/keccak.cl b/bfgminer/share/bfgminer/opencl/keccak.cl
@@ -0,0 +1,133 @@
+/*
+ * Scrypt-jane public domain, OpenCL implementation of scrypt(keccak,chacha,SCRYPTN,1,1) 2013 mtrlt
+ */
+
+// kernel-interface: fullheader Keccak
+
+#define ARGS_25(x) x ## 0, x ## 1, x ## 2, x ## 3, x ## 4, x ## 5, x ## 6, x ## 7, x ## 8, x ## 9, x ## 10, x ## 11, x ## 12, x ## 13, x ## 14, x ## 15, x ## 16, x ## 17, x ## 18, x ## 19, x ## 20, x ## 21, x ## 22, x ## 23, x ## 24
+
+__constant uint2 keccak_round_constants[24] = {
+	(uint2)(0x00000001,0x00000000), (uint2)(0x00008082,0x00000000),
+	(uint2)(0x0000808a,0x80000000), (uint2)(0x80008000,0x80000000),
+	(uint2)(0x0000808b,0x00000000), (uint2)(0x80000001,0x00000000),
+	(uint2)(0x80008081,0x80000000), (uint2)(0x00008009,0x80000000),
+	(uint2)(0x0000008a,0x00000000), (uint2)(0x00000088,0x00000000),
+	(uint2)(0x80008009,0x00000000), (uint2)(0x8000000a,0x00000000),
+	(uint2)(0x8000808b,0x00000000), (uint2)(0x0000008b,0x80000000),
+	(uint2)(0x00008089,0x80000000), (uint2)(0x00008003,0x80000000),
+	(uint2)(0x00008002,0x80000000), (uint2)(0x00000080,0x80000000),
+	(uint2)(0x0000800a,0x00000000), (uint2)(0x8000000a,0x80000000),
+	(uint2)(0x80008081,0x80000000), (uint2)(0x00008080,0x80000000),
+	(uint2)(0x80000001,0x00000000), (uint2)(0x80008008,0x80000000)
+};
+
+uint2 ROTL64_1(const uint2 x, const uint y)
+{
+	return (uint2)((x.x<<y)^(x.y>>(32-y)),(x.y<<y)^(x.x>>(32-y)));
+}
+uint2 ROTL64_2(const uint2 x, const uint y)
+{
+	return (uint2)((x.y<<y)^(x.x>>(32-y)),(x.x<<y)^(x.y>>(32-y)));
+}
+
+#define RND(i) \
+do{  \
+		m0 = *s0 ^ *s5 ^ *s10 ^ *s15 ^ *s20 ^ ROTL64_1(*s2 ^ *s7 ^ *s12 ^ *s17 ^ *s22, 1);\
+		m1 = *s1 ^ *s6 ^ *s11 ^ *s16 ^ *s21 ^ ROTL64_1(*s3 ^ *s8 ^ *s13 ^ *s18 ^ *s23, 1);\
+		m2 = *s2 ^ *s7 ^ *s12 ^ *s17 ^ *s22 ^ ROTL64_1(*s4 ^ *s9 ^ *s14 ^ *s19 ^ *s24, 1);\
+		m3 = *s3 ^ *s8 ^ *s13 ^ *s18 ^ *s23 ^ ROTL64_1(*s0 ^ *s5 ^ *s10 ^ *s15 ^ *s20, 1);\
+		m4 = *s4 ^ *s9 ^ *s14 ^ *s19 ^ *s24 ^ ROTL64_1(*s1 ^ *s6 ^ *s11 ^ *s16 ^ *s21, 1);\
+\
+		m5 = *s1^m0;\
+\
+		*s0 ^= m4;\
+		*s1 = ROTL64_2(*s6^m0, 12);\
+		*s6 = ROTL64_1(*s9^m3, 20);\
+		*s9 = ROTL64_2(*s22^m1, 29);\
+		*s22 = ROTL64_2(*s14^m3, 7);\
+		*s14 = ROTL64_1(*s20^m4, 18);\
+		*s20 = ROTL64_2(*s2^m1, 30);\
+		*s2 = ROTL64_2(*s12^m1, 11);\
+		*s12 = ROTL64_1(*s13^m2, 25);\
+		*s13 = ROTL64_1(*s19^m3,  8);\
+		*s19 = ROTL64_2(*s23^m2, 24);\
+		*s23 = ROTL64_2(*s15^m4, 9);\
+		*s15 = ROTL64_1(*s4^m3, 27);\
+		*s4 = ROTL64_1(*s24^m3, 14);\
+		*s24 = ROTL64_1(*s21^m0,  2);\
+		*s21 = ROTL64_2(*s8^m2, 23);\
+		*s8 = ROTL64_2(*s16^m0, 13);\
+		*s16 = ROTL64_2(*s5^m4, 4);\
+		*s5 = ROTL64_1(*s3^m2, 28);\
+		*s3 = ROTL64_1(*s18^m2, 21);\
+		*s18 = ROTL64_1(*s17^m1, 15);\
+		*s17 = ROTL64_1(*s11^m0, 10);\
+		*s11 = ROTL64_1(*s7^m1,  6);\
+		*s7 = ROTL64_1(*s10^m4,  3);\
+		*s10 = ROTL64_1(      m5,  1);\
+		\
+		m5 = *s0; m6 = *s1; *s0 = bitselect(*s0^*s2,*s0,*s1); *s1 = bitselect(*s1^*s3,*s1,*s2); *s2 = bitselect(*s2^*s4,*s2,*s3); *s3 = bitselect(*s3^m5,*s3,*s4); *s4 = bitselect(*s4^m6,*s4,m5);\
+		m5 = *s5; m6 = *s6; *s5 = bitselect(*s5^*s7,*s5,*s6); *s6 = bitselect(*s6^*s8,*s6,*s7); *s7 = bitselect(*s7^*s9,*s7,*s8); *s8 = bitselect(*s8^m5,*s8,*s9); *s9 = bitselect(*s9^m6,*s9,m5);\
+		m5 = *s10; m6 = *s11; *s10 = bitselect(*s10^*s12,*s10,*s11); *s11 = bitselect(*s11^*s13,*s11,*s12); *s12 = bitselect(*s12^*s14,*s12,*s13); *s13 = bitselect(*s13^m5,*s13,*s14); *s14 = bitselect(*s14^m6,*s14,m5);\
+		m5 = *s15; m6 = *s16; *s15 = bitselect(*s15^*s17,*s15,*s16); *s16 = bitselect(*s16^*s18,*s16,*s17); *s17 = bitselect(*s17^*s19,*s17,*s18); *s18 = bitselect(*s18^m5,*s18,*s19); *s19 = bitselect(*s19^m6,*s19,m5);\
+		m5 = *s20; m6 = *s21; *s20 = bitselect(*s20^*s22,*s20,*s21); *s21 = bitselect(*s21^*s23,*s21,*s22); *s22 = bitselect(*s22^*s24,*s22,*s23); *s23 = bitselect(*s23^m5,*s23,*s24); *s24 = bitselect(*s24^m6,*s24,m5);\
+\
+		*s0 ^= keccak_round_constants[i];  \
+}while(0)
+
+void keccak_block_noabsorb(ARGS_25(uint2* s))
+{
+	uint2 m0,m1,m2,m3,m4,m5,m6;
+#pragma unroll
+	for (uint i = 0; i < 24; ++i)
+		RND(i);
+}
+
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
+__kernel void search(
+#ifndef GOFFSET
+	const uint base,
+#endif
+	__global const uint2*restrict in, __global uint*restrict output)
+{
+#ifdef GOFFSET
+	const uint base = 0;
+#endif
+	uint2 ARGS_25(state);
+
+	state0 = in[0];
+	state1 = in[1];
+	state2 = in[2];
+	state3 = in[3];
+	state4 = in[4];
+	state5 = in[5];
+	state6 = in[6];
+	state7 = in[7];
+	state8 = in[8];
+	state9 = (uint2)(in[9].x, base + get_global_id(0));
+	state10 = (uint2)(1,0);
+	state11 = 0;
+	state12 = 0;
+	state13 = 0;
+	state14 = 0;
+	state15 = 0;
+	state16 = (uint2)(0,0x80000000U);
+	state17 = 0;
+	state18 = 0;
+	state19 = 0;
+	state20 = 0;
+	state21 = 0;
+	state22 = 0;
+	state23 = 0;
+	state24 = 0;
+
+	keccak_block_noabsorb(ARGS_25(&state));
+
+#define FOUND (0x0F)
+#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
+
+	if ((state3.y & 0xFFFFFFF0U) == 0)
+	{
+		SETFOUND(base + get_global_id(0));
+	}
+}
diff --git a/bfgminer/share/bfgminer/opencl/phatk.cl b/bfgminer/share/bfgminer/opencl/phatk.cl
@@ -2,7 +2,7 @@
 // I have therefore decided to keep it public-domain.
 // Modified version copyright 2011-2012 Con Kolivas
 
-// kernel-interface: phatk
+// kernel-interface: phatk SHA256d
 
 #ifdef VECTORS4
 	typedef uint4 u;

diff --git a/bfgminer/share/bfgminer/opencl/poclbm.cl b/bfgminer/share/bfgminer/opencl/poclbm.cl
@@ -5,7 +5,7 @@
 // This file is taken and modified from the public-domain poclbm project, and
 // we have therefore decided to keep it public-domain in Phoenix.
 
-// kernel-interface: poclbm
+// kernel-interface: poclbm SHA256d
 
 #ifdef VECTORS4
 	typedef uint4 u;

diff --git a/bfgminer/share/bfgminer/opencl/psw.cl b/bfgminer/share/bfgminer/opencl/psw.cl
@@ -29,7 +29,7 @@
  * online backup system.
  */
 
-// kernel-interface: scrypt
+// kernel-interface: scrypt scrypt
 
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
@@ -697,13 +697,13 @@ void salsa(uint4 B[8])
 #define Coord(x,y,z) x+y*(x ## SIZE)+z*(y ## SIZE)*(x ## SIZE)
 #define CO Coord(z,x,y)
 
-void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
+void scrypt_core(const uint gid, uint4 X[8], __global uint4 * restrict lookup)
 {
 	shittify(X);
 	const uint zSIZE = 8;
 	const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
 	const uint xSIZE = CONCURRENT_THREADS;
-	uint x = get_global_id(0)%xSIZE;
+	const uint x = gid % xSIZE;
 
 	for(uint y=0; y<1024/LOOKUP_GAP; ++y)
 	{
@@ -754,11 +754,19 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 #define SETFOUND(Xnonce) output[output[SCRYPT_FOUND]++] = Xnonce
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void search(__global const uint4 * restrict input,
+__kernel void search(
+#ifndef GOFFSET
+	const uint base,
+#endif
+	__global const uint4 * restrict input,
 volatile __global uint*restrict output, __global uint4*restrict padcache,
 const uint4 midstate0, const uint4 midstate16, const uint target)
 {
-	uint gid = get_global_id(0);
+	const uint gid = get_global_id(0)
+#ifndef GOFFSET
+		+ base
+#endif
+	;
 	uint4 X[8];
 	uint4 tstate0, tstate1, ostate0, ostate1, tmp0, tmp1;
 	uint4 data = (uint4)(input[4].x,input[4].y,input[4].z,gid);
@@ -783,7 +791,7 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 		SHA256(&pad0,&pad1, data, (uint4)(i+1,K[84],0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[87]));
 		SHA256(X+i*2,X+i*2+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
 	}
-	scrypt_core(X,padcache);
+	scrypt_core(gid, X, padcache);
 	SHA256(&tmp0,&tmp1, X[0], X[1], X[2], X[3]);
 	SHA256(&tmp0,&tmp1, X[4], X[5], X[6], X[7]);
 	SHA256_fixed(&tmp0,&tmp1);

diff --git a/bfgminer/share/bfgminer/opencl/scrypt.cl b/bfgminer/share/bfgminer/opencl/scrypt.cl
@@ -31,7 +31,7 @@
  * online backup system.
  */
 
-// kernel-interface: scrypt
+// kernel-interface: scrypt scrypt
 
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
@@ -760,13 +760,13 @@ void salsa(uint4 B[8])
 #define Coord(x,y,z) x+y*(x ## SIZE)+z*(y ## SIZE)*(x ## SIZE)
 #define CO Coord(z,x,y)
 
-void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
+void scrypt_core(const uint gid, uint4 X[8], __global uint4 * restrict lookup)
 {
 	shittify(X);
 	const uint zSIZE = 8;
 	const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
 	const uint xSIZE = CONCURRENT_THREADS;
-	uint x = get_global_id(0)%xSIZE;
+	const uint x = gid % xSIZE;
 
 	for(uint y=0; y<1024/LOOKUP_GAP; ++y)
 	{
@@ -817,11 +817,19 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 #define SETFOUND(Xnonce) output[output[SCRYPT_FOUND]++] = Xnonce
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void search(__global const uint4 * restrict input,
+__kernel void search(
+#ifndef GOFFSET
+	const uint base,
+#endif
+	__global const uint4 * restrict input,
 volatile __global uint*restrict output, __global uint4*restrict padcache,
 const uint4 midstate0, const uint4 midstate16, const uint target)
 {
-	uint gid = get_global_id(0);
+	const uint gid = get_global_id(0)
+#ifndef GOFFSET
+		+ base
+#endif
+	;
 	uint4 X[8];
 	uint4 tstate0, tstate1, ostate0, ostate1, tmp0, tmp1;
 	uint4 data = (uint4)(input[4].x,input[4].y,input[4].z,gid);
@@ -846,7 +854,7 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 		SHA256(&pad0,&pad1, data, (uint4)(i+1,K[84],0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[87]));
 		SHA256(X+i*2,X+i*2+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
 	}
-	scrypt_core(X,padcache);
+	scrypt_core(gid, X, padcache);
 	SHA256(&tmp0,&tmp1, X[0], X[1], X[2], X[3]);
 	SHA256(&tmp0,&tmp1, X[4], X[5], X[6], X[7]);
 	SHA256_fixed(&tmp0,&tmp1);

diff --git a/bfgminer/share/bfgminer/opencl/zuikkis.cl b/bfgminer/share/bfgminer/opencl/zuikkis.cl
@@ -28,7 +28,7 @@
  * online backup system.
  */
 
-// kernel-interface: scrypt
+// kernel-interface: scrypt scrypt
 
 __constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
 __constant uint K[] = {
@@ -748,13 +748,13 @@ void salsa(uint4 B[8])
 #define Coord(x,y,z) x+y*(x ## SIZE)+z*(y ## SIZE)*(x ## SIZE)
 #define CO Coord(z,x,y)
 
-void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
+void scrypt_core(const uint gid, uint4 X[8], __global uint4 * restrict lookup)
 {
 	shittify(X);
 	const uint zSIZE = 8;
 	const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
 	const uint xSIZE = CONCURRENT_THREADS;
-	uint x = get_global_id(0)%xSIZE;
+	const uint x = gid % xSIZE;
 
 	for(uint y=0; y<1024/LOOKUP_GAP; ++y)
 	{
@@ -792,11 +792,19 @@ void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
 #define SETFOUND(Xnonce) output[output[SCRYPT_FOUND]++] = Xnonce
 
 __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
-__kernel void search(__global const uint4 * restrict input,
+__kernel void search(
+#ifndef GOFFSET
+	const uint base,
+#endif
+	__global const uint4 * restrict input,
 volatile __global uint*restrict output, __global uint4*restrict padcache,
 const uint4 midstate0, const uint4 midstate16, const uint target)
 {
-	uint gid = get_global_id(0);
+	const uint gid = get_global_id(0)
+#ifndef GOFFSET
+		+ base
+#endif
+	;
 	uint4 X[8];
 	uint4 tstate0, tstate1, ostate0, ostate1, tmp0, tmp1;
 	uint4 data = (uint4)(input[4].x,input[4].y,input[4].z,gid);
@@ -820,7 +828,7 @@ const uint4 midstate0, const uint4 midstate16, const uint target)
 		SHA256(&pad0,&pad1, data, (uint4)(i+1,K[84],0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[87]));
 		SHA256(X+i*2,X+i*2+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
 	}
-	scrypt_core(X,padcache);
+	scrypt_core(gid, X, padcache);
 	SHA256(&tmp0,&tmp1, X[0], X[1], X[2], X[3]);
 	SHA256(&tmp0,&tmp1, X[4], X[5], X[6], X[7]);
 	SHA256_fixed(&tmp0,&tmp1);