rescue_prime.c
cl_int
hash_elements(cl_context ctx,
              cl_command_queue cq,
              cl_kernel krnl,
              cl_ulong* input,
              const size_t input_width,
              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 output_width = 4ul;
  const size_t in_size =
    global_size_x * global_size_y * input_width * sizeof(cl_ulong);
  const size_t out_size =
    global_size_x * global_size_y * output_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);

  // buffer for keeping width of input to hash_elements kernel, stored in
  // constant memory
  cl_mem in_width_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(size_t), 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);
  // input width being copied to device memory
  cl_event evt_1;
  status = clEnqueueWriteBuffer(cq,
                                in_width_buf,
                                CL_FALSE,
                                0,
                                sizeof(size_t),
                                &input_width,
                                0,
                                NULL,
                                &evt_1);

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

  // 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), &in_width_buf);
  status = clSetKernelArg(krnl, 2, sizeof(cl_mem), &mds_buf);
  status = clSetKernelArg(krnl, 3, sizeof(cl_mem), &ark1_buf);
  status = clSetKernelArg(krnl, 4, sizeof(cl_mem), &ark2_buf);
  status = clSetKernelArg(krnl, 5, sizeof(cl_mem), &out_buf);

  // preparing for creating dependency in compute execution graph
  cl_event evts[] = { evt_0, evt_1, evt_2, evt_3, evt_4 };
  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_5;
  status = clEnqueueNDRangeKernel(
    cq, krnl, 2, NULL, global_size, local_size, 5, evts, &evt_5);

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

  status = clWaitForEvents(1, &evt_6);

  // figure out kernel execution time only when asked to
  if (ts != NULL) {
    cl_ulong start, end;
    status = clGetEventProfilingInfo(
      evt_5, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
    check(status);
    status = clGetEventProfilingInfo(
      evt_5, 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);
  clReleaseEvent(evt_6);

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

  return status;
}
// This function is used for computing (global_size_x x global_size_y)-many
// rescue prime hashes, each on equal sized input of width N, and produces same
// number of outputs, each of width 4
//
// Note when I mentioned about width of N or 4, I mean input/ output
// will have those many 64-bit prime field elements
//
// local_size_{x, y} denotes work-group size vertically/ horizontally
//
// I'm going to use this function for testing `merge` kernel
// https://github.com/itzmeanjan/vectorized-rescue-prime/blob/bf40c7e41431487633288b7f64ebd804245fd8eb/kernel.cl#L378
//
// Now I've also started using this generic function for benchmarking
// `hash_elements` kernel
This function is used for computing (global_size_x x global_size_y)-many
rescue prime hashes, each on equal sized input of width N, and produces same
number of outputs, each of width 4

Note when I mentioned about width of N or 4, I mean input/ output
will have those many 64-bit prime field elements

local_size_{x, y} denotes work-group size vertically/ horizontally



Now I've also started using this generic function for benchmarking
`hash_elements` kernel
This function is used for computing ( global_size_x x global_size_y ) - many rescue prime hashes each on equal sized input of width N and produces same number of outputs each of width 4 Note when I mentioned about width of N or 4 I mean input / output will have those many 64 - bit prime field elements local_size_ { x y } denotes work - group size vertically / horizontally Now I ' ve also started using this generic function for benchmarking ` hash_elements ` kernel
cl_int
hash_elements(cl_context ctx,
              cl_command_queue cq,
              cl_kernel krnl,
              cl_ulong* input,
              const size_t input_width,
              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 output_width = 4ul;
  const size_t in_size =
    global_size_x * global_size_y * input_width * sizeof(cl_ulong);
  const size_t out_size =
    global_size_x * global_size_y * output_width * sizeof(cl_ulong);
  cl_mem in_buf = clCreateBuffer(ctx, CL_MEM_READ_ONLY, in_size, NULL, &status);
  cl_mem in_width_buf =
    clCreateBuffer(ctx, CL_MEM_READ_ONLY, sizeof(size_t), 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,
                                in_width_buf,
                                CL_FALSE,
                                0,
                                sizeof(size_t),
                                &input_width,
                                0,
                                NULL,
                                &evt_1);
  cl_event evt_2;
  status = clEnqueueWriteBuffer(
    cq, mds_buf, CL_FALSE, 0, sizeof(MDS), MDS, 0, NULL, &evt_2);
  cl_event evt_3;
  status = clEnqueueWriteBuffer(
    cq, ark1_buf, CL_FALSE, 0, sizeof(ARK1), ARK1, 0, NULL, &evt_3);
  cl_event evt_4;
  status = clEnqueueWriteBuffer(
    cq, ark2_buf, CL_FALSE, 0, sizeof(ARK2), ARK2, 0, NULL, &evt_4);
  status = clSetKernelArg(krnl, 0, sizeof(cl_mem), &in_buf);
  status = clSetKernelArg(krnl, 1, sizeof(cl_mem), &in_width_buf);
  status = clSetKernelArg(krnl, 2, sizeof(cl_mem), &mds_buf);
  status = clSetKernelArg(krnl, 3, sizeof(cl_mem), &ark1_buf);
  status = clSetKernelArg(krnl, 4, sizeof(cl_mem), &ark2_buf);
  status = clSetKernelArg(krnl, 5, sizeof(cl_mem), &out_buf);
  cl_event evts[] = { evt_0, evt_1, evt_2, evt_3, evt_4 };
  size_t global_size[] = { global_size_x, global_size_y };
  size_t local_size[] = { local_size_x, local_size_y };
  cl_event evt_5;
  status = clEnqueueNDRangeKernel(
    cq, krnl, 2, NULL, global_size, local_size, 5, evts, &evt_5);
  cl_event evt_6;
  status = clEnqueueReadBuffer(
    cq, out_buf, CL_FALSE, 0, out_size, output, 1, &evt_5, &evt_6);
  status = clWaitForEvents(1, &evt_6);
  if (ts != NULL) {
    cl_ulong start, end;
    status = clGetEventProfilingInfo(
      evt_5, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
    check(status);
    status = clGetEventProfilingInfo(
      evt_5, 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);
  clReleaseEvent(evt_6);
  clReleaseMemObject(in_buf);
  clReleaseMemObject(in_width_buf);
  clReleaseMemObject(mds_buf);
  clReleaseMemObject(ark1_buf);
  clReleaseMemObject(ark2_buf);
  clReleaseMemObject(out_buf);
  return status;
}
cl_int hash_elements ( cl_context ctx , cl_command_queue cq , cl_kernel krnl , cl_ulong * input , const size_t input_width , 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 output_width = 4ul ; const size_t in_size = global_size_x * global_size_y * input_width * sizeof ( cl_ulong ) ; const size_t out_size = global_size_x * global_size_y * output_width * sizeof ( cl_ulong ) ; cl_mem in_buf = clCreateBuffer ( ctx , CL_MEM_READ_ONLY , in_size , NULL , & status ) ; cl_mem in_width_buf = clCreateBuffer ( ctx , CL_MEM_READ_ONLY , sizeof ( size_t ) , 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 , in_width_buf , CL_FALSE , 0 , sizeof ( size_t ) , & input_width , 0 , NULL , & evt_1 ) ; cl_event evt_2 ; status = clEnqueueWriteBuffer ( cq , mds_buf , CL_FALSE , 0 , sizeof ( MDS ) , MDS , 0 , NULL , & evt_2 ) ; cl_event evt_3 ; status = clEnqueueWriteBuffer ( cq , ark1_buf , CL_FALSE , 0 , sizeof ( ARK1 ) , ARK1 , 0 , NULL , & evt_3 ) ; cl_event evt_4 ; status = clEnqueueWriteBuffer ( cq , ark2_buf , CL_FALSE , 0 , sizeof ( ARK2 ) , ARK2 , 0 , NULL , & evt_4 ) ; status = clSetKernelArg ( krnl , 0 , sizeof ( cl_mem ) , & in_buf ) ; status = clSetKernelArg ( krnl , 1 , sizeof ( cl_mem ) , & in_width_buf ) ; status = clSetKernelArg ( krnl , 2 , sizeof ( cl_mem ) , & mds_buf ) ; status = clSetKernelArg ( krnl , 3 , sizeof ( cl_mem ) , & ark1_buf ) ; status = clSetKernelArg ( krnl , 4 , sizeof ( cl_mem ) , & ark2_buf ) ; status = clSetKernelArg ( krnl , 5 , sizeof ( cl_mem ) , & out_buf ) ; cl_event evts [ ] = { evt_0 , evt_1 , evt_2 , evt_3 , evt_4 } ; size_t global_size [ ] = { global_size_x , global_size_y } ; size_t local_size [ ] = { local_size_x , local_size_y } ; cl_event evt_5 ; status = clEnqueueNDRangeKernel ( cq , krnl , 2 , NULL , global_size , local_size , 5 , evts , & evt_5 ) ; cl_event evt_6 ; status = clEnqueueReadBuffer ( cq , out_buf , CL_FALSE , 0 , out_size , output , 1 , & evt_5 , & evt_6 ) ; status = clWaitForEvents ( 1 , & evt_6 ) ; if ( ts != NULL ) { cl_ulong start , end ; status = clGetEventProfilingInfo ( evt_5 , CL_PROFILING_COMMAND_START , sizeof ( cl_ulong ) , & start , NULL ) ; check ( status ) ; status = clGetEventProfilingInfo ( evt_5 , 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 ) ; clReleaseEvent ( evt_6 ) ; clReleaseMemObject ( in_buf ) ; clReleaseMemObject ( in_width_buf ) ; clReleaseMemObject ( mds_buf ) ; clReleaseMemObject ( ark1_buf ) ; clReleaseMemObject ( ark2_buf ) ; clReleaseMemObject ( out_buf ) ; return status ; }
This function is used for computing (global_size_x x global_size_y)-many
rescue prime hashes, each on equal sized input of width N, and produces same
number of outputs, each of width 4
This function is used for computing ( global_size_x x global_size_y ) - many rescue prime hashes each on equal sized input of width N and produces same number of outputs each of width 4
// input is supplied to kernel by this buffer // buffer for keeping width of input to hash_elements kernel, stored in // constant memory // 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 // input width 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 kernel execution time only when asked to