title: Opencl Reduction is not as expected
body: <p>I'm pretty a novice about opencl. I have tried about "get the summation of all cubes of every element in an array". Here's my kernel code:</p>
  
  <pre><code>kernel void cubeSum(global float *input,
                      local float *prods,
                      global float *output )
  {
      int gid = get_global_id( 0 );
      int tnum = get_local_id( 0 ); // thread number
      int wgNum = get_group_id( 0 ); // work-group number
      int numItems = get_local_size( 0 );
      prods[ tnum ] = input[ gid ] * input[ gid ] * input[gid]; // cube
  
      for (int offset = 1; offset &lt; numItems; offset *= 2) {
          int mask = 2 * offset - 1;
          barrier(CLK_LOCAL_MEM_FENCE);
          if ( (tnum &amp; mask) == 0 ) {
              prods[tnum] += prods[tnum + offset];
          }
      }
      barrier(CLK_LOCAL_MEM_FENCE);
  
      if ( tnum == 0 )
          output[wgNum] = prods[0];
  }
  </code></pre>
  
  <p>I can't figure out why my result is not the same with sequential result. When the array is from 0 to 511, my result is sequential result minus 2048; when the array is from 0 to 1023, my result is sequential result plus 16384.</p>
  
  <p>I will try to figure it out myself while I'm waiting for you answers.</p>
  
  <p>Another question is I found it is hard to debug kernel code since the dataset is quite big and it runs concurrently. Any advice for debugging?</p>
  
  <p>All the advices are appreciated =).</p>
  
  <p>By the way, here's my host code:</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;stdio.h&gt;
  #include &lt;math.h&gt;
  #include &lt;string.h&gt;
  #include &lt;stdlib.h&gt;
  #include &lt;OpenCL/opencl.h&gt;
  
  #define NUM_ELEMENTS (512)
  #define LOCAL_SIZE (512)
  #define MAX_SOURCE_SIZE (0x100000)
  
  int main(int argc, const char * argv[])
  {
      float data[NUM_ELEMENTS];           //hA
      float sum;
      float sumTest;
  
      size_t global;
      size_t local;
      size_t numWorkGroups;
      size_t dataSize;
      size_t resultsSize;
  
      cl_device_id device;
      cl_context context;
      cl_command_queue cmdQueue;
      cl_program program;
      cl_kernel kernel;
  
      cl_mem input;
      cl_mem output;
  
      FILE *fp;
      //failed to use relative path here. permission problem?
      char fileName[] = "/Users/sure/USC/590/cubeSum/cubeSum/cubeSum.cl";
      char *source_str;
      size_t source_size;
  
      /* カーネルを含むソースコードをロード */
      fp = fopen(fileName, "r");
      if (!fp) {
          fprintf(stderr, "Failed to load kernel.\n");
          exit(1);
      }
      source_str = (char*)malloc(MAX_SOURCE_SIZE);
      source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp);
      fclose( fp );
  
      //allocate the host memory buffers:
      int i = 0;
      unsigned int count = NUM_ELEMENTS;
      for (i = 0; i &lt; count; i++) {
          data[i] = i;
      }
  
      //array size in bytes (will need this later):
      dataSize = NUM_ELEMENTS * sizeof(float);
  
      //opencl function status
      cl_int status;
  
      // Connect to a compute device
      //
      int gpu = 1;
  
      status = clGetDeviceIDs(NULL, gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &amp;device, NULL);
      if (status != CL_SUCCESS)
      {
          printf("Error: Failed to create a device group!\n");
          return EXIT_FAILURE;
      }
  
      //create an Opencl context
      context = clCreateContext(NULL, 1, &amp;device, NULL, NULL, &amp;status);
  
      //create a command queue
      cmdQueue = clCreateCommandQueue( context, device, 0, &amp;status );
  
      //allocate memory buffers on the device
      input = clCreateBuffer( context, CL_MEM_READ_ONLY, dataSize, NULL, &amp;status );   //dA
  
      //TODO: at this line, I don't have the value of local which is calculated by clGetKernelWorkGroupInfo
      //need to figure out a way to avoid hardcode it.
      output = clCreateBuffer( context, CL_MEM_WRITE_ONLY, sizeof(float) * NUM_ELEMENTS / LOCAL_SIZE, NULL, &amp;status ); //dC
  
      // enqueue the 2 commands to write data into the device buffers:
      status = clEnqueueWriteBuffer( cmdQueue, input, CL_FALSE, 0, dataSize, data, 0, NULL, NULL );
  
      // create the kernel program on the device:
      program = clCreateProgramWithSource(context, 1, (const char **) &amp; source_str, (const size_t *)&amp;source_size, &amp;status);
      if (!program)
      {
          printf("Error: Failed to create compute program!\n");
          return EXIT_FAILURE;
      }
  
  
      // Build the program executable
      //
      status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
      if (status != CL_SUCCESS)
      {
          size_t len;
          char buffer[2048];
  
          printf("Error: Failed to build program executable!\n");
          clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &amp;len);
          printf("%s\n", buffer);
          exit(1);
      }
  
      //create compute kernel
      kernel = clCreateKernel( program, "cubeSum", &amp;status );
  
      // Get the maximum work group size for executing the kernel on the device
      //
      status = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &amp;local, NULL);
      if (status != CL_SUCCESS)
      {
          printf("Error: Failed to retrieve kernel work group info! %d\n", status);
          exit(1);
      }
  
      global = count;
  
      numWorkGroups = global / local;
      float results[numWorkGroups];        //hC
      resultsSize = numWorkGroups * sizeof(float);
  
      //set kernel parameter
      status = clSetKernelArg( kernel, 0, sizeof(cl_mem), &amp;input );
      status = clSetKernelArg( kernel, 1, sizeof(float), NULL );
      status = clSetKernelArg( kernel, 2, sizeof(cl_mem), &amp;output );
  
      // Execute the kernel over the entire range of our 1d input data set
      // using the maximum number of work group items for this device
      //
      status = clEnqueueNDRangeKernel(cmdQueue, kernel, 1, NULL, &amp;global, &amp;local, 0, NULL, NULL);
      if (status)
      {
          printf("Error: Failed to execute kernel!\n");
          return EXIT_FAILURE;
      }
  
      clFinish(cmdQueue);
      status = clEnqueueReadBuffer( cmdQueue, output, CL_TRUE, 0, resultsSize, results, 0, NULL, NULL );
  
      // Validate our results
      //
      sum = 0;
  
      for (int i=0; i&lt;numWorkGroups; i++) {
          sum += results[i];
      }
  
      sumTest = 0;
      for(i = 0; i &lt; count; i++)
      {
          sumTest += data[i] * data[i] * data[i];
      }
  
      // Print a brief summary detailing the results
      //
      printf("Computed '%f/%f'!\n", sum, sumTest);
  
      // Shutdown and cleanup
      //
      clReleaseMemObject(input);
      clReleaseMemObject(output);
      clReleaseProgram(program);
      clReleaseKernel(kernel);
      clReleaseCommandQueue(cmdQueue);
      clReleaseContext(context);
  
      return 0;
  
  }
  </code></pre>
  
  <p>EDIT: Just found another thing. My code is correct if I just sum all element without cube/square. Thus, I'm gonna figure out how cube affect to my program.</p>