132 lines
4.3 KiB
Common Lisp
132 lines
4.3 KiB
Common Lisp
//#pragma OPENCL EXTENSION cl_khr_gl_sharing : enable
|
|
|
|
__kernel void CreateTextures(__global const float *inputs, __write_only image2d_t displacement, __write_only image2d_t slopeFoam, uint2 dim, float2 twoCellsSize,
|
|
float chopScale, float dt, __global float *foamBuffer)
|
|
{
|
|
__global const float *H = inputs;
|
|
__global const float *chopX = inputs + (dim.x * dim.y);
|
|
__global const float *chopZ = chopX + (dim.x * dim.y);
|
|
|
|
int2 gridPos;
|
|
gridPos.x = get_global_id(0);
|
|
gridPos.y = get_global_id(1);
|
|
|
|
int idx = gridPos.y * dim.x + gridPos.x;
|
|
|
|
int prevX = gridPos.x > 0 ? gridPos.x-1 : dim.x-1;
|
|
int nextX = gridPos.x < dim.x-1 ? gridPos.x + 1 : 0;
|
|
int prevY = gridPos.y > 0 ? gridPos.y - 1 : dim.y-1;
|
|
int nextY = gridPos.y < dim.y - 1 ? gridPos.y + 1 : 0;
|
|
|
|
float thisChopX = chopX[idx] * chopScale;
|
|
float nextChopX = chopX[gridPos.y * dim.x + nextX];
|
|
float prevChopX = chopX[gridPos.y * dim.x + prevX];
|
|
float thisChopZ = chopZ[idx] * chopScale;
|
|
float nextChopZ = chopZ[nextY * dim.x + gridPos.x];
|
|
float prevChopZ = chopZ[prevY * dim.x + gridPos.x];
|
|
|
|
float4 color;
|
|
color.x = thisChopX;
|
|
color.y = thisChopZ;
|
|
color.z = H[idx];
|
|
color.w = 1.0f;
|
|
write_imagef(displacement, gridPos, color);
|
|
|
|
float xWidth = twoCellsSize.x + nextChopX - prevChopX;
|
|
float yDepth = twoCellsSize.y + nextChopZ - prevChopZ;
|
|
|
|
float xDelta = (H[gridPos.y * dim.x + nextX] - H[gridPos.y * dim.x + prevX]);
|
|
float yDelta = (H[nextY * dim.x + gridPos.x] - H[prevY * dim.x + gridPos.x]);
|
|
float dx = xDelta / xWidth;
|
|
float dy = yDelta / yDepth;
|
|
|
|
|
|
float sxy = (chopX[nextY * dim.x + gridPos.x] - chopX[prevY * dim.x + gridPos.x]) / yDepth;
|
|
|
|
float syy = (chopZ[nextY * dim.x + gridPos.x] - chopZ[prevY * dim.x + gridPos.x]) / yDepth;
|
|
|
|
float syx = (chopZ[gridPos.y * dim.x + nextX] - chopZ[gridPos.y * dim.x + prevX]) / xWidth;
|
|
|
|
float sxx = (chopX[gridPos.y * dim.x + nextX] - chopX[gridPos.y * dim.x + prevX]) / xWidth;
|
|
|
|
float3 sx = (float3)(1.0, 0.0, dx);
|
|
float3 sy = (float3)(0.0, 1.0, dy);
|
|
float3 norm = normalize(cross(sx, sy));
|
|
|
|
color.xyz = norm;
|
|
|
|
float Jxx = 1.0 + chopScale * sxx;
|
|
float Jyy = 1.0 + chopScale * syy;
|
|
float Jxy = chopScale * sxy;
|
|
float Jyx = chopScale * syx;
|
|
float J = Jxx * Jyy - Jxy * Jyx;
|
|
|
|
float foam = 1.0f - J;
|
|
|
|
color.w = foam;
|
|
|
|
write_imagef(slopeFoam, gridPos, color);
|
|
foamBuffer[gridPos.y * dim.x + gridPos.x] = foam;
|
|
}
|
|
|
|
__kernel void ProcessWater(__global const float *H0, __global const float *omega, __global float *fftIn, float t, uint2 dim, float2 size)
|
|
{
|
|
__global float *H = fftIn;
|
|
__global float *chopX = fftIn + (dim.x * dim.y * 2);
|
|
__global float *chopZ = chopX + (dim.x * dim.y * 2);
|
|
|
|
int2 gridPos;
|
|
gridPos.x = get_global_id(0);
|
|
gridPos.y = get_global_id(1);
|
|
|
|
int h0idx = gridPos.y*(dim.x + 1)*2 + gridPos.x * 2;
|
|
float2 h0;
|
|
h0.x = H0[h0idx];
|
|
h0.y = H0[h0idx + 1];
|
|
|
|
int h0NegKIdx = (dim.y - gridPos.y) * (dim.x + 1) * 2 + (dim.x - gridPos.x) * 2;
|
|
float2 h0NegKConj;
|
|
h0NegKConj.x = H0[h0NegKIdx];
|
|
h0NegKConj.y = H0[h0NegKIdx+1] * -1;
|
|
|
|
uint2 halfdim = dim / 2;
|
|
float2 kPos;
|
|
kPos.x = (float)gridPos.x - (float)halfdim.x;
|
|
kPos.y = (float)gridPos.y - (float)halfdim.y;
|
|
|
|
const float TWOPI = 3.14159265f * 2.0f;
|
|
float2 K = (TWOPI * kPos) / size;
|
|
|
|
float wk = omega[dim.x * gridPos.y + gridPos.x];
|
|
float wkt = wk * t;
|
|
float cwkt = native_cos(wkt);
|
|
float swkt = native_sin(wkt);
|
|
|
|
float2 term1, term2;
|
|
term1.x = h0.x * cwkt - h0.y * swkt;
|
|
term1.y = h0.x * swkt + h0.y * cwkt;
|
|
term2.x = h0NegKConj.x * cwkt - h0NegKConj.y * -swkt;
|
|
term2.y = h0NegKConj.x * -swkt + h0NegKConj.y * cwkt;
|
|
|
|
float2 Htilde = term1 + term2;
|
|
|
|
int outIdx = gridPos.y * dim.x * 2 + gridPos.x * 2;
|
|
H[outIdx] = Htilde.x;
|
|
H[outIdx+1] = Htilde.y;
|
|
|
|
if (dot(K, K) > 0) {
|
|
float2 chopImg = fast_normalize(K);
|
|
float2 cX, cZ;
|
|
cX.x = -(chopImg.x * Htilde.y);
|
|
cX.y = (chopImg.x * Htilde.x);
|
|
cZ.x = -(chopImg.y * Htilde.y);
|
|
cZ.y = (chopImg.y * Htilde.x);
|
|
|
|
chopX[outIdx] = cX.x;
|
|
chopX[outIdx+1] = cX.y;
|
|
chopZ[outIdx] = cZ.x;
|
|
chopZ[outIdx+1] = cZ.y;
|
|
} else {
|
|
chopX[outIdx] = chopX[outIdx+1] = chopZ[outIdx] = chopZ[outIdx+1] = 0.0f;
|
|
}
|
|
} |