title: opencl synchronization
body: <p>I am new to opencl and there seems to be something about the barrier function I don't understand. This is the code for my kernel. This is a standard matrix vector calculation with the output in *w. there is 1 work group with 64 work units, the same as the dimension of the vector</p>
  
  <pre><code>#pragma OPENCL EXTENSION cl_khr_fp64 : enable
  __kernel void fmin_stuff(__global double *h, __global double *g, __global double  
    *w,int n,__global int * gid) {
  
  // Get the index of the current element
  int i = get_global_id(0);
  int j;
  gid[i]=get_local_id(0);
  
  w[i]=-g[i];
  barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE);
  for (j=0;j&lt;n;j++)
  {
    if (j&lt;i)
      w[i]-=h[i+j*n]*w[j];
    barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE);
  }
  }
  </code></pre>
  
  <p>The problem is that the code fails at random. The output is correct for a while.  Here are the initial values for w for each run.</p>
  
  <pre><code>-0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.34999 2.51524 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.10141 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.72261 2.80155 
  -0.148351 -0.309007 0.133204 -1.39589 2.88335 -2.68636 2.77369 
  </code></pre>
  
  <p>The program reports that the kernel executed successfully in each case. For all runs the values in the vector w are eventually incorrect. any advice would be greatly appreciated.</p>
  
  <p>There was some confusion over whether this is a simple matrix multiplication. It is not. this is what the code is trying to accomplish where I include olnly the first 5 terms of w.</p>
  
  <pre><code>w(1)=-g(1);
  w(2)=-g(2);
  w(3)=-g(3);
  w(4)=-g(4);
  w(5)=-g(5);
  
  w(2)-=h(2)*w(1);
  w(3)-=h(3)*w(1);
  w(4)-=h(4)*w(1);
  w(5)-=h(5)*w(1);
  
  w(3)-=h(3+N)*w(2);
  w(4)-=h(4+N)*w(2);
  w(5)-=h(5+N)*w(2);
  
  w(4)-=h(4+2*N)*w(3);
  w(5)-=h(5+2*N)*w(3);
  
  w(5)-=h(5+3*N)*w(4);
  </code></pre>
  
  <p>Also the kernel is only called once per program run. The  random behaviour results from running the program mutiple times.  </p>
  
  <p>The comment led me to see what I was doing wrong. I had the work groups and items configured as</p>
  
  <pre><code>size_t global_item_size[3] = {N, 1, 1}; // Process the entire lists
  size_t local_item_size[3] = {1,1,1}; // Process in groups of 64
  ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
          global_item_size, local_item_size, 0, NULL, NULL);
  </code></pre>
  
  <p>when it should have been.</p>
  
  <pre><code>size_t global_item_size[3] = {N, 1, 1}; // Process the entire lists
  size_t local_item_size[3] = {N,1,1}; // Process in groups of 64
  ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
          global_item_size, local_item_size, 0, NULL, NULL);
  </code></pre>
  
  <p>Thanks for the help.  This is great for me but probably not of much interest
  to to others.</p>