integrate gateless gate

hominoids · Jan 17, 2017 · 60f35ad · 60f35ad
1 parent b296789
commit 60f35ad
Show file tree

Hide file tree

Showing 8 changed files with 1,527 additions and 1,127 deletions.
diff --git a/algorithm.c b/algorithm.c
@@ -1147,6 +1147,7 @@ static cl_int queue_cryptonight_kernel(_clState *clState, dev_blk_ctx *blk, __ma
 }
 
 
+#define WORKSIZE clState->wsize
 
 static cl_int queue_equihash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads)
 {
@@ -1160,43 +1161,43 @@ static cl_int queue_equihash_kernel(_clState *clState, dev_blk_ctx *blk, __maybe
   uint32_t dbg[2] = {0};
   status |= clEnqueueWriteBuffer(clState->commandQueue, clState->padbuffer8, CL_TRUE, 0, sizeof(dbg), &dbg, 0, NULL, NULL);
 
-  cl_mem buf_ht[2] = {clState->CLbuffer0, clState->buffer1};
   cl_mem rowCounters[2] = {clState->buffer2, clState->buffer3};
   for (int round = 0; round < PARAM_K; round++) {
-    size_t global_ws = NR_ROWS / ROWS_PER_UINT;
+    size_t global_ws = RC_SIZE;
     size_t local_ws = 256;
     unsigned int num = 0;
     cl_kernel *kernel = &clState->extra_kernels[0];
     // Now on every round!!!!
-    CL_SET_ARG(buf_ht[round % 2]);
+    CL_SET_ARG(clState->index_buf[round]);
     CL_SET_ARG(rowCounters[round % 2]);
+    CL_SET_ARG(clState->outputBuffer);
+    CL_SET_ARG(clState->CLbuffer0);
     status |= clEnqueueNDRangeKernel(clState->commandQueue, *kernel, 1, NULL, &global_ws, &local_ws, 0, NULL, NULL);
 
-    num = 0;
     kernel = &clState->extra_kernels[1 + round];
     if (!round) {
-      CL_SET_ARG(clState->MidstateBuf);
-      CL_SET_ARG(buf_ht[round % 2]);
-      CL_SET_ARG(rowCounters[round % 2]);
-      work_items = threads;
+      worksize = LOCAL_WORK_SIZE_ROUND0;
+      work_items = NR_INPUTS / ROUND0_INPUTS_PER_WORK_ITEM;
     }
     else {
-      CL_SET_ARG(buf_ht[(round - 1) % 2]);
-      CL_SET_ARG(buf_ht[round % 2]);
-      CL_SET_ARG(rowCounters[(round - 1) % 2]);
-      CL_SET_ARG(rowCounters[round % 2]);
-      work_items = NR_ROWS;
+      worksize = LOCAL_WORK_SIZE;
+      work_items = NR_ROWS * worksize;
     }
-    CL_SET_ARG(clState->padbuffer8);
-    if (round == PARAM_K - 1)
-      CL_SET_ARG(clState->outputBuffer);
-    status |= clEnqueueNDRangeKernel(clState->commandQueue, *kernel, 1, NULL, &work_items, &worksize, 0, NULL, NULL);
+    status |= clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[1 + round], 1, NULL, &work_items, &worksize, 0, NULL, NULL);
   }
-  work_items = NR_ROWS;
+
+  worksize = LOCAL_WORK_SIZE_POTENTIAL_SOLS;
+  work_items = NR_ROWS * worksize;
+  status |= clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[1 + 9], 1, NULL, &work_items, &worksize, 0, NULL, NULL); 
+
+  worksize = LOCAL_WORK_SIZE_SOLS;
+  work_items = MAX_POTENTIAL_SOLS * worksize;
   status |= clEnqueueNDRangeKernel(clState->commandQueue, clState->kernel, 1, NULL, &work_items, &worksize, 0, NULL, NULL);
 
   return status;
 }
+#undef WORKSIZE
+
 
 static algorithm_settings_t algos[] = {
   // kernels starting from this will have difficulty calculated by using litecoin algorithm