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.
__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];
}