source: CIVL/examples/translation/openclversion2.15/reduceCL.cvl@ c2a3f74

#include "cl.cvl"
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <string.h>
#include <civlc.h>

#include "reduce_int_kernel.cvl"

$input int MAX_WORKGROUP_SIZE;

/*
    __global int *output, 
    __global const int *input, 
    __local int *shared,
    const unsigned int n
*/
typedef struct
{
  void * param[4];
}args;

cl_kernel clCreateKernel(args * argument, char * function)
{
  cl_kernel kernel;
  kernel.arguments = argument;
  kernel.method = function;
  
  return kernel;
}

void clReleaseKernel(cl_kernel kernel)
{
  for (int i = 0; i < 4; i++)
  {
    free(((args*)kernel.arguments)->param[i]);
  }
}

#define MIN_ERROR       (1e-7)
#define MAX_GROUPS      (64)
#define MAX_WORK_ITEMS  (64)
#define SEPARATOR       ("----------------------------------------------------------------------\n")

//count is too large, was 1024 * 1024
static int iterations = 1000;
static int count      = 4 * 4;
static int channels   = 1;
static bool integer   = true;


void reduce_validate_int(int *data, int size, int * result)
{
    int i;
    int sum = data[0];
    int c = (int)0.0f;              
    for (i = 1; i < size; i++)
    {
        int y = data[i] - c;  
        int t = sum + y;      
        c = (t - sum) - y;  
        sum = t;            
    }
    result[0] = sum;
}

void create_reduction_pass_counts(
    int count, 
    int max_group_size,    
    int max_groups,
    int max_work_items, 
    int *pass_count, 
    size_t **group_counts, 
    size_t **work_item_counts,
    int **operation_counts,
    int **entry_counts)
{
    int work_items = (count < max_work_items * 2) ? count / 2 : max_work_items;
    if(count < 1)
        work_items = 1;
        
    int groups = count / (work_items * 2);
    groups = max_groups < groups ? max_groups : groups;

    int max_levels = 1;
    int s = groups;

    while(s > 1) 
    {
        int work_items = (s < max_work_items * 2) ? s / 2 : max_work_items;
        s = s / (work_items*2);
        max_levels++;
    }
 
    *group_counts = (size_t*)malloc(max_levels * sizeof(size_t));
    *work_item_counts = (size_t*)malloc(max_levels * sizeof(size_t));
    *operation_counts = (int*)malloc(max_levels * sizeof(int));
    *entry_counts = (int*)malloc(max_levels * sizeof(int));

    (*pass_count) = max_levels;
    (*group_counts)[0] = groups;
    (*work_item_counts)[0] = work_items;
    (*operation_counts)[0] = 1;
    (*entry_counts)[0] = count;
    if(max_group_size < work_items)
    {
        (*operation_counts)[0] = work_items;
        (*work_item_counts)[0] = max_group_size;
    }
    
    s = groups;
    int level = 1;
   
    while(s > 1) 
    {
        int work_items = (s < max_work_items * 2) ? s / 2 : max_work_items;
        int groups = s / (work_items * 2);
        groups = (max_groups < groups) ? max_groups : groups;

        (*group_counts)[level] = groups;
        (*work_item_counts)[level] = work_items;
        (*operation_counts)[level] = 1;
        (*entry_counts)[level] = s;
        if(max_group_size < work_items)
        {
            (*operation_counts)[level] = work_items;
            (*work_item_counts)[level] = max_group_size;
        }
        
        s = s / (work_items*2);
        level++;
    }
}


int main(int argc, char** argv)
{
  args * arguments;
  arguments = (args*)malloc(sizeof(args));

  int                   t1 = 0;
  int                   t2 = 0;
  int                   err;
  cl_device_id          device_id;
  cl_command_queue      commands;
  cl_context            context;
  /*
  cl_mem                        output_buffer;
  cl_mem                input_buffer;
  cl_mem                partials_buffer;
  */
  void *                        output_buffer;
  void *                input_buffer;
  void *                partials_buffer;
  size_t                typesize;
  int                   pass_count = 0;
  size_t*               group_counts = 0;
  size_t*               work_item_counts = 0;
  int*                  operation_counts = 0;
  int*                  entry_counts = 0;
  int                   use_gpu = 1;
    
  int i;
  int c;
  
  float *float_data = (float*)malloc(count * channels * sizeof(float));
  int *integer_data = (int*)malloc(count * channels * sizeof(int));
  
  //Cannot use strstr
  use_gpu = 0;
  integer = true;
  channels = 1;
  
  //can use random, but cannot check numbers
  for (i = 0; i < count * channels; i++)
  {
    float_data[i] = ((float) rand() / (float) RAND_MAX);
    integer_data[i] = (int) (255.0f * float_data[i]);
  }
  
  size_t returned_size = 0;
  size_t max_workgroup_size = 0;
  //clGetDeviceInfo happens here
  max_workgroup_size = MAX_WORKGROUP_SIZE;
  
  printf(SEPARATOR);
  typesize = integer ? (sizeof(int)) : (sizeof(float));
  
  
  size_t buffer_size = typesize * count * channels;
  input_buffer = (int*)malloc(buffer_size);
  
  void *input_data = (integer) ? (void*)integer_data : (void*)float_data;
  
  //clEnqueueWriteBuffer
  memcpy(input_buffer, input_data, buffer_size);
  
  //clCreateBuffer
  partials_buffer = (int*)malloc(buffer_size);
  output_buffer = (int*)malloc(buffer_size);
  
  create_reduction_pass_counts(
    count, max_workgroup_size, 
    MAX_GROUPS, MAX_WORK_ITEMS, 
    &pass_count, &group_counts, 
    &work_item_counts, &operation_counts,
    &entry_counts);
    
  cl_kernel *kernels = (cl_kernel*)malloc(pass_count * sizeof(cl_kernel));
  //memset(kernels, 0, pass_count * sizeof(cl_kernel));
  //no memset
  
  for(i = 0; i < pass_count; i++)
  {
    kernels[i] = clCreateKernel(arguments, "buffer");
  }
  
  void * pass_swap;
  void * pass_input = output_buffer;
  void * pass_output = input_buffer;
    
  for(i = 0; i < pass_count; i++)
  {
    size_t global = group_counts[i] * work_item_counts[i];        
    size_t local = work_item_counts[i];
    unsigned int operations = operation_counts[i];
    unsigned int entries = entry_counts[i];
    size_t shared_size = typesize * channels * local * operations;

    printf("Pass[%4d] Global[%4d] Local[%4d] Groups[%4d] WorkItems[%4d] Operations[%d] Entries[%d]\n",  i, 
      (int)global, (int)local, (int)group_counts[i], (int)work_item_counts[i], operations, entries);
      
    pass_swap = pass_input;
    pass_input = pass_output;
    pass_output = pass_swap;
/*
  __global int *output, 
  __global const int *input, 
  __local int *shared,
  const unsigned int n
*/
    
      
  }
  
  for(i = 0; i < pass_count; i++)
  {
    clReleaseKernel(kernels[i]);
  }
  
  free(arguments);
  
  free(output_buffer);
  free(partials_buffer);
  free(input_buffer);
  
  free(group_counts);
  free(work_item_counts);
  free(operation_counts);
  free(entry_counts);
  //free(computed_result);
  free(kernels);
  free(float_data);
  free(integer_data);
  
  return 0;
}