Merge pull request #258 from fancyIX/feature/#247

Feature/#247
fancyIX · Dec 13, 2021 · ef5f3ee · ef5f3ee
2 parents af34e4e + 2a493a8
commit ef5f3ee
Showing 1 changed file with 24 additions and 28 deletions.
diff --git a/kernel/heavyhash.cl b/kernel/heavyhash.cl
@@ -207,16 +207,17 @@ static void keccak_f1600_no_absorb(uint2* a)
 	} 
 }
 
-__attribute__((reqd_work_group_size(64, 1, 1)))
+__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
 __kernel void search(__constant uint *header, __constant uchar* gmatrix, __global uint* output, const ulong target)
 {
     uint tid = get_local_id(0);
     uint gid = get_global_id(0);
 
 	__local ulong2 lmatrix[64 * 4];
+	uint round = 256 / WORKSIZE;
 	#pragma unroll
-	for (int i = 0; i < 4; i++) {
-		lmatrix[tid * 4 + i] = ((__constant ulong2 *) gmatrix)[tid * 4 + i];
+	for (int i = 0; i < round; i++) {
+		lmatrix[tid * round + i] = ((__constant ulong2 *) gmatrix)[tid * round + i];
 	}
 
     uint pdata[50] = {0};
@@ -226,58 +227,53 @@ __kernel void search(__constant uint *header, __constant uchar* gmatrix, __globa
     }
     pdata[19] = gid;
 
-    uchar hash_first[32];
     uchar hash_second[32];
-    uchar hash_xored[32];
 
     uint vector[64];
-    uint product[64];
 
     ((uchar *) pdata)[80] = 0x06;
     ((uchar *) pdata)[135] = 0x80;
 
     keccak_f1600_no_absorb(pdata);
 
-    for (int i = 0; i < 4; i++) {
-        ((ulong *)hash_first)[i] = ((ulong *) pdata)[i];
-    }
 	#pragma unroll
     for (int i = 0; i < 32; ++i) {
-        vector[(2*i)] = (hash_first[i] >> 4);
-        vector[(2*i+1)] = hash_first[i] & 0xF;
+        vector[(2*i)] = (((uchar *) pdata)[i] >> 4);
+        vector[(2*i+1)] = ((uchar *) pdata)[i] & 0xF;
     }
 
-    #pragma nounroll
-    for (int i = 0; i < 64; ++i) {
-        uint sum = 0;
+	volatile uint sum = 0;
+	volatile uint sum2 = 0;
+    for (int i = 0; i < 32; ++i) {
+        sum = 0;
 		#pragma unroll
 		for (int j = 0; j < 4; j++) {
 			#pragma unroll
 			for (int k = 0; k < 16; k++) {
-				uint mv = ((__local uchar *)lmatrix)[(4 * i + j) * 16 + k];
+				uint mv = ((__local uchar *)lmatrix)[(4 * (i * 2) + j) * 16 + k];
 				sum += mv * vector[j * 16 + k];
 			}
 		}
-		product[(i)] = (sum >> 10);
-    }
-
-
-    for (int i = 0; i < 32; ++i) {
-        hash_second[i] = (product[(2*i)] << 4) | (product[(2*i+1)]);
+		sum2 = 0;
+		#pragma unroll
+		for (int j = 0; j < 4; j++) {
+			#pragma unroll
+			for (int k = 0; k < 16; k++) {
+				uint mv = ((__local uchar *)lmatrix)[(4 * (i * 2 + 1) + j) * 16 + k];
+				sum2 += mv * vector[j * 16 + k];
+			}
+		}
+		hash_second[i] = ((sum >> 10) << 4) | (sum2 >> 10);
     }
-
+    #pragma unroll
     for (int i = 0; i < 32; ++i) {
-        hash_xored[i] = hash_first[i] ^ hash_second[i];
+        ((uchar *) pdata)[i] = ((uchar *) pdata)[i] ^ hash_second[i];
     }
 
-    for (int i = 0; i < 50; i++) {
+    for (int i = 8; i < 50; i++) {
 		pdata[i] = 0;
 	}
 
-    for (int i = 0; i < 32; i++) {
-        ((uchar *) pdata)[i] = hash_xored[i];
-    }
-
     ((uchar *)pdata)[32] = 0x06;
     ((uchar *)pdata)[135] = 0x80;