static void compiler_mandelbrot_alternate(void)
{
  const size_t global[2] = {w, h};
  const size_t local[2] = {16, 1};
  const size_t sz = w * h * sizeof(char[4]);
  const float rcpWidth = 1.f / float(w);
  const float rcpHeight = 1.f / float(h);

  OCL_CREATE_KERNEL("compiler_mandelbrot_alternate");

  OCL_CREATE_BUFFER(buf[0], 0, sz, NULL);
  OCL_CALL (clSetKernelArg, kernel, 0, sizeof(cl_mem), &buf[0]);
  OCL_CALL (clSetKernelArg, kernel, 1, sizeof(float), &rcpWidth);
  OCL_CALL (clSetKernelArg, kernel, 2, sizeof(float), &rcpHeight);
  OCL_CALL (clSetKernelArg, kernel, 3, sizeof(float), &criterium);
  OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, 2, NULL, global, local, 0, NULL, NULL);
  OCL_MAP_BUFFER(0);
  dst = (int *) buf_data[0];

  /* Save the image (for debug purpose) */
  cl_write_bmp(dst, w, h, "compiler_mandelbrot_alternate.bmp");

  /* Compare with the golden image */
  OCL_CHECK_IMAGE(dst, w, h, "compiler_mandelbrot_alternate_ref.bmp");
}
static void compiler_box_blur()
{
  OCL_CREATE_KERNEL("compiler_box_blur");

  /* Load the picture */
  src = cl_read_bmp("sample.bmp", &w, &h);
  sz = w * h * sizeof(int);

  /* Run the kernel */
  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, sz, src);
  OCL_CREATE_BUFFER(buf[1], 0, sz, NULL);
  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
  OCL_SET_ARG(2, sizeof(int), &w);
  OCL_SET_ARG(3, sizeof(int), &h);
  OCL_SET_ARG(4, sizeof(int), &chunk);
  globals[0] = size_t(w/4);
  globals[1] = h/chunk + ((h%chunk)?1:0);
  locals[0] = 16;
  locals[1] = 1;
  free(src);
  OCL_NDRANGE(2);
  OCL_MAP_BUFFER(1);
  dst = (int*) buf_data[1];

  /* Save the image (for debug purpose) */
  cl_write_bmp(dst, w, h, "compiler_box_blur.bmp");

  /* Compare with the golden image */
  OCL_CHECK_IMAGE(dst, w, h, "compiler_box_blur_ref.bmp");
}