rescue_prime.c
cl_int
merge(cl_context ctx,
      cl_command_queue cq,
      cl_kernel krnl,
      cl_ulong* input,
      cl_ulong* output,
      size_t global_size_x,
      size_t global_size_y,
      size_t local_size_x,
      size_t local_size_y,
      cl_ulong* ts)
{
  cl_int status;

  const size_t in_width = 8ul;
  const size_t out_width = 4ul;
  const size_t in_size =
    global_size_x * global_size_y * in_width * sizeof(cl_ulong);
  const size_t out_size =
    global_size_x * global_size_y * out_width * sizeof(cl_ulong);

  // input is supplied to kernel by this buffer
  cl_mem in_buf = clCreateBuffer(ctx, CL_MEM_READ_ONLY, in_size, NULL, &status);

  // Following three buffers will be storing rescue prime hash constants
  cl_mem mds_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(MDS), NULL, &status);
  cl_mem ark1_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(ARK1), NULL, &status);
  cl_mem ark2_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(ARK2), NULL, &status);

  // output to be placed here, after kernel completes hash computation
  cl_mem out_buf =
    clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, out_size, NULL, &status);

  // input being copied to device memory
  cl_event evt_0;
  status = clEnqueueWriteBuffer(
    cq, in_buf, CL_FALSE, 0, in_size, input, 0, NULL, &evt_0);

  // scheduling rescue prime constant copying
  cl_event evt_1;
  status = clEnqueueWriteBuffer(
    cq, mds_buf, CL_FALSE, 0, sizeof(MDS), MDS, 0, NULL, &evt_1);
  cl_event evt_2;
  status = clEnqueueWriteBuffer(
    cq, ark1_buf, CL_FALSE, 0, sizeof(ARK1), ARK1, 0, NULL, &evt_2);
  cl_event evt_3;
  status = clEnqueueWriteBuffer(
    cq, ark2_buf, CL_FALSE, 0, sizeof(ARK2), ARK2, 0, NULL, &evt_3);

  // setting kernel arguments for
  // https://github.com/itzmeanjan/vectorized-rescue-prime/blob/bf40c7e41431487633288b7f64ebd804245fd8eb/kernel.cl#L320
  status = clSetKernelArg(krnl, 0, sizeof(cl_mem), &in_buf);
  status = clSetKernelArg(krnl, 1, sizeof(cl_mem), &mds_buf);
  status = clSetKernelArg(krnl, 2, sizeof(cl_mem), &ark1_buf);
  status = clSetKernelArg(krnl, 3, sizeof(cl_mem), &ark2_buf);
  status = clSetKernelArg(krnl, 4, sizeof(cl_mem), &out_buf);

  // preparing for creating dependency in compute execution graph
  cl_event evts[] = { evt_0, evt_1, evt_2, evt_3 };
  size_t global_size[] = { global_size_x, global_size_y };
  size_t local_size[] = { local_size_x, local_size_y };

  // kernel being dispatched for execution on device
  cl_event evt_4;
  status = clEnqueueNDRangeKernel(
    cq, krnl, 2, NULL, global_size, local_size, 4, evts, &evt_4);

  // hash output being copied back to host
  cl_event evt_5;
  status = clEnqueueReadBuffer(
    cq, out_buf, CL_FALSE, 0, out_size, output, 1, &evt_4, &evt_5);

  status = clWaitForEvents(1, &evt_5);

  // figure out how long kernel execution originally took
  // only when function caller wants this function to figure out
  if (ts != NULL) {
    cl_ulong start, end;
    status = clGetEventProfilingInfo(
      evt_4, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
    check(status);
    status = clGetEventProfilingInfo(
      evt_4, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
    check(status);

    *ts = (end - start);
  }

  clReleaseEvent(evt_0);
  clReleaseEvent(evt_1);
  clReleaseEvent(evt_2);
  clReleaseEvent(evt_3);
  clReleaseEvent(evt_4);
  clReleaseEvent(evt_5);

  clReleaseMemObject(in_buf);
  clReleaseMemObject(mds_buf);
  clReleaseMemObject(ark1_buf);
  clReleaseMemObject(ark2_buf);
  clReleaseMemObject(out_buf);

  return status;
}
// Merges two rescue prime hashes into single one by dispatching
// https://github.com/itzmeanjan/vectorized-rescue-prime/blob/bf40c7e41431487633288b7f64ebd804245fd8eb/kernel.cl#L378
// kernel (global_size_x x global_size_y)-many times
//
// It should produce same number of rescue prime digests, each of width 4,
// while in input each was of width 8, because two input rescue prime digests
// are being merged


It should produce same number of rescue prime digests, each of width 4,
while in input each was of width 8, because two input rescue prime digests
are being merged
It should produce same number of rescue prime digests each of width 4 while in input each was of width 8 because two input rescue prime digests are being merged
cl_int
merge(cl_context ctx,
      cl_command_queue cq,
      cl_kernel krnl,
      cl_ulong* input,
      cl_ulong* output,
      size_t global_size_x,
      size_t global_size_y,
      size_t local_size_x,
      size_t local_size_y,
      cl_ulong* ts)
{
  cl_int status;
  const size_t in_width = 8ul;
  const size_t out_width = 4ul;
  const size_t in_size =
    global_size_x * global_size_y * in_width * sizeof(cl_ulong);
  const size_t out_size =
    global_size_x * global_size_y * out_width * sizeof(cl_ulong);
  cl_mem in_buf = clCreateBuffer(ctx, CL_MEM_READ_ONLY, in_size, NULL, &status);
  cl_mem mds_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(MDS), NULL, &status);
  cl_mem ark1_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(ARK1), NULL, &status);
  cl_mem ark2_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(ARK2), NULL, &status);
  cl_mem out_buf =
    clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, out_size, NULL, &status);
  cl_event evt_0;
  status = clEnqueueWriteBuffer(
    cq, in_buf, CL_FALSE, 0, in_size, input, 0, NULL, &evt_0);
  cl_event evt_1;
  status = clEnqueueWriteBuffer(
    cq, mds_buf, CL_FALSE, 0, sizeof(MDS), MDS, 0, NULL, &evt_1);
  cl_event evt_2;
  status = clEnqueueWriteBuffer(
    cq, ark1_buf, CL_FALSE, 0, sizeof(ARK1), ARK1, 0, NULL, &evt_2);
  cl_event evt_3;
  status = clEnqueueWriteBuffer(
    cq, ark2_buf, CL_FALSE, 0, sizeof(ARK2), ARK2, 0, NULL, &evt_3);
  status = clSetKernelArg(krnl, 0, sizeof(cl_mem), &in_buf);
  status = clSetKernelArg(krnl, 1, sizeof(cl_mem), &mds_buf);
  status = clSetKernelArg(krnl, 2, sizeof(cl_mem), &ark1_buf);
  status = clSetKernelArg(krnl, 3, sizeof(cl_mem), &ark2_buf);
  status = clSetKernelArg(krnl, 4, sizeof(cl_mem), &out_buf);
  cl_event evts[] = { evt_0, evt_1, evt_2, evt_3 };
  size_t global_size[] = { global_size_x, global_size_y };
  size_t local_size[] = { local_size_x, local_size_y };
  cl_event evt_4;
  status = clEnqueueNDRangeKernel(
    cq, krnl, 2, NULL, global_size, local_size, 4, evts, &evt_4);
  cl_event evt_5;
  status = clEnqueueReadBuffer(
    cq, out_buf, CL_FALSE, 0, out_size, output, 1, &evt_4, &evt_5);
  status = clWaitForEvents(1, &evt_5);
  if (ts != NULL) {
    cl_ulong start, end;
    status = clGetEventProfilingInfo(
      evt_4, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
    check(status);
    status = clGetEventProfilingInfo(
      evt_4, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
    check(status);
    *ts = (end - start);
  }
  clReleaseEvent(evt_0);
  clReleaseEvent(evt_1);
  clReleaseEvent(evt_2);
  clReleaseEvent(evt_3);
  clReleaseEvent(evt_4);
  clReleaseEvent(evt_5);
  clReleaseMemObject(in_buf);
  clReleaseMemObject(mds_buf);
  clReleaseMemObject(ark1_buf);
  clReleaseMemObject(ark2_buf);
  clReleaseMemObject(out_buf);
  return status;
}
cl_int merge ( cl_context ctx , cl_command_queue cq , cl_kernel krnl , cl_ulong * input , cl_ulong * output , size_t global_size_x , size_t global_size_y , size_t local_size_x , size_t local_size_y , cl_ulong * ts ) { cl_int status ; const size_t in_width = 8ul ; const size_t out_width = 4ul ; const size_t in_size = global_size_x * global_size_y * in_width * sizeof ( cl_ulong ) ; const size_t out_size = global_size_x * global_size_y * out_width * sizeof ( cl_ulong ) ; cl_mem in_buf = clCreateBuffer ( ctx , CL_MEM_READ_ONLY , in_size , NULL , & status ) ; cl_mem mds_buf = clCreateBuffer ( ctx , CL_MEM_READ_ONLY , sizeof ( MDS ) , NULL , & status ) ; cl_mem ark1_buf = clCreateBuffer ( ctx , CL_MEM_READ_ONLY , sizeof ( ARK1 ) , NULL , & status ) ; cl_mem ark2_buf = clCreateBuffer ( ctx , CL_MEM_READ_ONLY , sizeof ( ARK2 ) , NULL , & status ) ; cl_mem out_buf = clCreateBuffer ( ctx , CL_MEM_WRITE_ONLY , out_size , NULL , & status ) ; cl_event evt_0 ; status = clEnqueueWriteBuffer ( cq , in_buf , CL_FALSE , 0 , in_size , input , 0 , NULL , & evt_0 ) ; cl_event evt_1 ; status = clEnqueueWriteBuffer ( cq , mds_buf , CL_FALSE , 0 , sizeof ( MDS ) , MDS , 0 , NULL , & evt_1 ) ; cl_event evt_2 ; status = clEnqueueWriteBuffer ( cq , ark1_buf , CL_FALSE , 0 , sizeof ( ARK1 ) , ARK1 , 0 , NULL , & evt_2 ) ; cl_event evt_3 ; status = clEnqueueWriteBuffer ( cq , ark2_buf , CL_FALSE , 0 , sizeof ( ARK2 ) , ARK2 , 0 , NULL , & evt_3 ) ; status = clSetKernelArg ( krnl , 0 , sizeof ( cl_mem ) , & in_buf ) ; status = clSetKernelArg ( krnl , 1 , sizeof ( cl_mem ) , & mds_buf ) ; status = clSetKernelArg ( krnl , 2 , sizeof ( cl_mem ) , & ark1_buf ) ; status = clSetKernelArg ( krnl , 3 , sizeof ( cl_mem ) , & ark2_buf ) ; status = clSetKernelArg ( krnl , 4 , sizeof ( cl_mem ) , & out_buf ) ; cl_event evts [ ] = { evt_0 , evt_1 , evt_2 , evt_3 } ; size_t global_size [ ] = { global_size_x , global_size_y } ; size_t local_size [ ] = { local_size_x , local_size_y } ; cl_event evt_4 ; status = clEnqueueNDRangeKernel ( cq , krnl , 2 , NULL , global_size , local_size , 4 , evts , & evt_4 ) ; cl_event evt_5 ; status = clEnqueueReadBuffer ( cq , out_buf , CL_FALSE , 0 , out_size , output , 1 , & evt_4 , & evt_5 ) ; status = clWaitForEvents ( 1 , & evt_5 ) ; if ( ts != NULL ) { cl_ulong start , end ; status = clGetEventProfilingInfo ( evt_4 , CL_PROFILING_COMMAND_START , sizeof ( cl_ulong ) , & start , NULL ) ; check ( status ) ; status = clGetEventProfilingInfo ( evt_4 , CL_PROFILING_COMMAND_END , sizeof ( cl_ulong ) , & end , NULL ) ; check ( status ) ; * ts = ( end - start ) ; } clReleaseEvent ( evt_0 ) ; clReleaseEvent ( evt_1 ) ; clReleaseEvent ( evt_2 ) ; clReleaseEvent ( evt_3 ) ; clReleaseEvent ( evt_4 ) ; clReleaseEvent ( evt_5 ) ; clReleaseMemObject ( in_buf ) ; clReleaseMemObject ( mds_buf ) ; clReleaseMemObject ( ark1_buf ) ; clReleaseMemObject ( ark2_buf ) ; clReleaseMemObject ( out_buf ) ; return status ; }
Merges two rescue prime hashes into single one by dispatching
https://github.com/itzmeanjan/vectorized-rescue-prime/blob/bf40c7e41431487633288b7f64ebd804245fd8eb/kernel.cl#L378
kernel (global_size_x x global_size_y)-many times
Merges two rescue prime hashes into single one by dispatching https : // github . com / itzmeanjan / vectorized - rescue - prime / blob / bf40c7e41431487633288b7f64ebd804245fd8eb / kernel . cl#L378 kernel ( global_size_x x global_size_y ) - many times
// input is supplied to kernel by this buffer // Following three buffers will be storing rescue prime hash constants // output to be placed here, after kernel completes hash computation // input being copied to device memory // scheduling rescue prime constant copying // setting kernel arguments for // https://github.com/itzmeanjan/vectorized-rescue-prime/blob/bf40c7e41431487633288b7f64ebd804245fd8eb/kernel.cl#L320 // preparing for creating dependency in compute execution graph // kernel being dispatched for execution on device // hash output being copied back to host // figure out how long kernel execution originally took // only when function caller wants this function to figure out