reproduce.c

#define CL_TARGET_OPENCL_VERSION 120
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>

static void* slurp_file(const char *filename) {
  unsigned char *s;
  FILE *f = fopen(filename, "rb"); // To avoid Windows messing with linebreaks.
  if (f == NULL) return NULL;
  fseek(f, 0, SEEK_END);
  size_t src_size = ftell(f);
  fseek(f, 0, SEEK_SET);
  s = (unsigned char*) malloc(src_size + 1);
  if (fread(s, 1, src_size, f) != src_size) {
    free(s);
    s = NULL;
  } else {
    s[src_size] = '\0';
  }
  fclose(f);

  return s;
}

double double_zeroes[1024*1024*10] = { 0 };
cl_int int_zeroes[1024*1024*10] = { 0 };

int main() {
  const char *kernel_src = slurp_file("reproduce.cl");
  assert(kernel_src != NULL);

  cl_int err;

  cl_platform_id platform;
  cl_device_id device;
  cl_context ctx;
  cl_command_queue q;
  cl_program prog;

  err = clGetPlatformIDs(1, &platform, NULL);
  assert(err == CL_SUCCESS);

  err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
  assert(err == CL_SUCCESS);

  ctx = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
  assert(err == CL_SUCCESS);

  prog = clCreateProgramWithSource(ctx, 1, &kernel_src, NULL, &err);
  assert(err == CL_SUCCESS);

  err = clBuildProgram(prog, 1, &device, "", NULL, NULL);
  assert(err == CL_SUCCESS || err == CL_BUILD_PROGRAM_FAILURE);

  cl_build_status build_status;
  err = clGetProgramBuildInfo(prog,
                              device,
                              CL_PROGRAM_BUILD_STATUS,
                              sizeof(cl_build_status),
                              &build_status,
                              NULL);
  assert(err == CL_SUCCESS);

  if (build_status != CL_SUCCESS) {
    char *build_log;
    size_t ret_val_size;
    err = clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
    assert(err == CL_SUCCESS);

    build_log = (char*) malloc(ret_val_size+1);
    err = clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
    assert(err == CL_SUCCESS);

    build_log[ret_val_size] = '\0';

    fprintf(stderr, "Build log:\n%s\n", build_log);

    free(build_log);
  }

  cl_kernel k = clCreateKernel(prog, "mainzisegred_nonseg_169", &err);
  assert(err == CL_SUCCESS);

  q = clCreateCommandQueue(ctx, device, 0, &err);
  assert(err == CL_SUCCESS);

  int MiB = 1024*1024;

  cl_mem datapoints_mem_171 = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
                                             MiB, &double_zeroes, &err);
  assert(err == CL_SUCCESS);

  cl_mem binb_mem_172 = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
                                        MiB, &double_zeroes, &err);
  assert(err == CL_SUCCESS);

  cl_mem mem_175 = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
                                  MiB, &int_zeroes, &err);
  assert(err == CL_SUCCESS);

  cl_int numD_115 = 487;
  cl_int numBBins_116 = 21;

  int i = 0;

  err = clSetKernelArg(k, i++, sizeof(numD_115), &numD_115);
  assert(err == CL_SUCCESS);
  err = clSetKernelArg(k, i++, sizeof(numBBins_116), &numBBins_116);
  assert(err == CL_SUCCESS);
  err = clSetKernelArg(k, i++, sizeof(datapoints_mem_171), &datapoints_mem_171);
  assert(err == CL_SUCCESS);
  err = clSetKernelArg(k, i++, sizeof(binb_mem_172), &binb_mem_172);
  assert(err == CL_SUCCESS);
  err = clSetKernelArg(k, i++, sizeof(mem_175), &mem_175);
  assert(err == CL_SUCCESS);

  size_t global_work_size[1] = { 256 };
  size_t local_work_size[1] = { 256 };
  err = clEnqueueNDRangeKernel(q, k,
                               1, NULL, global_work_size, local_work_size,
                               0, NULL, 0);
  assert(err == CL_SUCCESS);

  err = clFinish(q);
  assert(err == CL_SUCCESS);

  cl_int res;
  err = clEnqueueReadBuffer(q, mem_175, 1, 0, sizeof(cl_int), &res, 0, NULL, NULL);
  printf("%d\n", (int)res);
}