title: Fastest way to simulate recurrent equations in OpenCL
body: <p>I am trying to simulate a recurrent equation of the type</p>
  
  <p>s<sub>i</sub>(t+1) = f[Σ<sub>j</sub> W<sub>ij</sub>s<sub>j</sub>(t) + v<sub>i</sub>*input(t)]</p>
  
  <p>in OpenCL, where f(.) is some non-linear function (in the code below it is just a step-function with threshold th) and s(t) is some external input. Naturally, I implemented one worker for every x<sub>i</sub>. In every time-step every worker calculates the result of the equation above and subsequently this result is shared with all other workers. Therefore, all workers have to be in the same workgroup.</p>
  
  <p>My current OpenCL kernel looks like this</p>
  
  <pre><code>__kernel void part1(__global int* s, __global float* W, __global float* Th, __global float* V, __global float* input, int N, int T)
      {
          unsigned int i = get_global_id(0);
  
          float value = 0;
          float v = V[i];
          float th = Th[i];  
  
          for(int t = 0; t &lt; T; t++){
              value = v*input[t];
              for(int j = 0; j &lt; N; j++){
                  value = value + W[i*N + j]*s[j];
              }
              barrier(CLK_GLOBAL_MEM_FENCE);
              if (value &gt;= th){
                  s[i] = 1;
              } else {
                  s[i] = 0;
              }
              barrier(CLK_GLOBAL_MEM_FENCE);
          }
      }
  </code></pre>
  
  <p>Unfortunately, this code is actually three times slower than an equivalent C-implementation. Also, I expected that a change in the number of workers should not make a huge difference (because new workers are sitting on new threads that run in parallel to the others), but actually the processing time increases linearly with the number of workers. The bottleneck seems to be the writing operation after the first barrier. Eliminating this operation (but leaving the barrier in place) cuts down the processing time by a factor of 25 and eliminates the linear dependence.</p>
  
  <p>I am pretty new to OpenCL and I would appreciate any help to speed this code up!</p>
  
  <p>Thanks a lot in advance!
  Blue2script</p>