Topic: Gateless Gate Sharp 1.3.8: 30Mh/s (Ethash) on RX 480! - page 43. (Read 214410 times)

  { "x16r", ALGO_X16R, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 32, 8 * 16 * 4194304, 0, x16r_regenhash, NULL, queue_x16r_kernel, gen_hash, append_x13_compiler_options, enqueue_x16r_kernels },

static cl_int queue_x16r_kernel(struct __clState *clState, struct _dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
  cl_kernel *kernel;
  unsigned int num;
  cl_ulong le_target;
  cl_int status = 0;
  uint8_t hashOrder[X16R_HASH_FUNC_COUNT];

  le_target = *(cl_ulong *)(blk->work->device_target + 24);
  flip80(clState->cldata, blk->work->data);
  x16r_getalgolist(&clState->cldata[4], hashOrder);

  status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL, NULL);
  if (status != CL_SUCCESS)
    return -1;

  for (int i = 0; i < X16R_HASH_FUNC_COUNT; i++) {
    kernel = &clState->extra_kernels[i];
    CL_SET_ARG_0(clState->padbuffer8);
  }

  kernel = &clState->extra_kernels[hashOrder[0] + X16R_HASH_FUNC_COUNT];
  num = 0;
  CL_SET_ARG(clState->CLbuffer0);
  CL_SET_ARG(clState->padbuffer8);

  kernel = &clState->kernel;
  num = 0;
  CL_SET_ARG(clState->padbuffer8);
  CL_SET_ARG(clState->outputBuffer);
  CL_SET_ARG(le_target);

  return status;
}

static cl_int enqueue_x16r_kernels(struct __clState *clState,
                                   size_t *p_global_work_offset, size_t *globalThreads, size_t *localThreads)
{
  cl_int status;
  uint8_t hashOrder[X16R_HASH_FUNC_COUNT];

  x16r_getalgolist(&clState->cldata[4], hashOrder);

  status = clEnqueueNDRangeKernel(clState->commandQueue,
      clState->extra_kernels[hashOrder[0] + X16R_HASH_FUNC_COUNT],
      1, p_global_work_offset,
      globalThreads, localThreads, 0, NULL, NULL);
  if (unlikely(status != CL_SUCCESS)) {
    applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)", status);
    return status;
  }

  for (int i = 1; i < X16R_HASH_FUNC_COUNT; i++) {
    status = clEnqueueNDRangeKernel(clState->commandQueue,
        clState->extra_kernels[hashOrder[i]],
        1, p_global_work_offset,
        globalThreads, localThreads, 0, NULL, NULL);
    if (unlikely(status != CL_SUCCESS)) {
      applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)", status);
      return status;
    }
  }

  status = clEnqueueNDRangeKernel(clState->commandQueue,
      clState->kernel,
      1, p_global_work_offset,
      globalThreads, localThreads, 0, NULL, NULL);
  if (unlikely(status != CL_SUCCESS)) {
    applog(LOG_ERR, "Error %d: Enqueueing kernel onto command queue. (clEnqueueNDRangeKernel)", status);
    return status;
  }

  return 0;
}