K φ° Δt/h2 It norm delay fps

Schrödinger equation on 256×256 square grid. |Ψ|2 for the travelling wave packet is plotted. Colors correspond to the wave function phase. Ψ = 0 at the square border. Data are stored in buffers. You can change number of relaxation algorithm iterations 2 It in every time step.

Reduction algorithm

The reduction algorithm is used to calculate the wave function norm
    s = ∑ii |2,   norm = s / s0 .
At first n = 256 partitial sums are calculated by rows and stored in local memory.
__kernel void kReduce(global float2* inp, global float* outp ){
  local float loc[n];
  int x = get_global_id(0),  y = n*n;
  float s = 0.0f;
  while( y > 0){
    float2 t = inp[x + y];
    s += dot(t,t);
    y -= n;
  }
  loc[x] = s;
Next we sum up in parallel these n values by pairs reducing data twice every iteration
  y = (n >>1);
  while( y > 1){
    barrier( CLK_LOCAL_MEM_FENCE );
    if( x < y) loc[x] += loc[x + y];
    y = (y >>1);
  }
  if( x == 0) outp[0] = loc[0] + loc[1];
}

WebCL Demos     updated 5 Nov 2011