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 = ∑_{i} |Ψ_{i }|^{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];
}