Add Ocean OpenCL-OpenGL interop demo

samples/extensions/khr/CMakeLists.txt

@@ -16,4 +16,5 @@ add_subdirectory(histogram)
 if(OPENCL_SDK_BUILD_OPENGL_SAMPLES)
   add_subdirectory(conway)
   add_subdirectory(nbody)
+  add_subdirectory(ocean)
 endif()

samples/extensions/khr/ocean/CMakeLists.txt

@@ -0,0 +1,23 @@
+# Copyright (c) 2024 Mobica Limited, Marcin Hajder, Piotr Plebański
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+add_sample(
+    TARGET ocean_gl_ocl_interop
+    VERSION 300
+    SOURCES main.cpp ocean.cpp ocean.hpp ocean_util.hpp
+    KERNELS twiddle.cl time_spectrum.cl inversion.cl normals.cl fft_kernel.cl init_spectrum.cl
+    SHADERS ocean.vert.glsl ocean.frag.glsl
+    LIBS
+        glm::glm
+    )

samples/extensions/khr/ocean/README.md

@@ -0,0 +1,63 @@
+# Ocean surface simulation with Opencl and OpenGL interoperability
+
+## Sample Purpose
+
+This sample demonstrates how to share compute/render resources between OpenCL and OpenGL to simulate an ocean surface and compare it with the Vulkan equivalent sample. If the cl_khr_gl_sharing extension is available and requested through CLI option, some OpenCL images will be created through a clCreateFromGLTexture procesure and updated every frame through clEnqueueAcquireGLObjects and clEnqueueReleaseGLObjects procedures. These images will then be used for ocean rendering. If cl_khr_gl_sharing is not available, additional copying from OpenCL buffers to OpenGL images will be performed.
+
+## Key APIs and Concepts
+
+The primary focus of this sample is to understand how to set up shared resources between OpenCL and OpenGL interoperability and compare both performance and necessary workload between this and equivalent vulkan sample. Additionally, this sample demonstrates how to approach physical, real-time simulations in OpenCL and the API objects involved in executing an OpenCL application such as ocean surface simulation.
+
+
+### Application flow
+
+The application performs an initial setup during which:
+
+    -OpenCL and OpenGL resources such as device and context are prepared.
+    -A SFML window, camera, and related keyboard event callbacks are created.
+    -Both shared and private resources for OpenCL and OpenGL are set up.
+
+Available CLI options are as follows:
+
+    --useGLSharing, requests use of cl_khr_gl_sharing
+
+After the setup, the simulation starts with initial ocean parameters that can be modified with keyboard events in real-time:
+
+    - a/z - Increase/decrease wind magnitude.
+    - s/x - Change wind heading.
+    - d/c - Increase/decrease waving amplitude.
+    - f/v - Increase/decrease wave choppiness.
+    - g/b - Increase/decrease additional altitude scale.
+
+Additionally, the simulation and rendering can be paused with the Space key. Rendering can toggle between wireframe and filled modes using the 'w' key. Application tracks its performance in the title bar of the window, it could be toggled by pressing 'e' key.
+
+While the simulation is in progress, each frame of the application performs the following general steps:
+
+    -Uniform buffers are updated to handle camera and ocean parameters.
+    -OpenCL kernels are enqueued.
+    -The ocean grid is rendered using the previous OpenCL computation outcome.
+
+
+### Kernel logic
+
+Multiple kernels follow the general steps (with multiple optimizations) described in the publication:  [Realtime GPGPU FFT ocean water simulation](https://tore.tuhh.de/bitstream/11420/1439/1/GPGPU_FFT_Ocean_Simulation.pdf)
+
+### Used API surface
+
+```c++
+cl::util::supports_extension(cl::Device, cl::string)
+cl::util::read_exe_relative_text_file(const char*, cl_int* const)
+cl::util::read_exe_relative_binary_file(const char*, cl_int* const)
+cl::Context(cl::Device)
+cl::CommandQueue(cl::Context, cl::Device)
+cl::Platform::get(vector<Platform>)
+cl::Platform::getDevices(Type, vector<Device>)
+cl::Program::build()
+cl::Image2D(cl::Context, cl_mem_flags, ImageFormat, size_type, size_type)
+cl::Error::what()
+cl::Error::err()
+cl::NDRange(size_type, size_type)
+cl::Buffer::Buffer(cl::Context, cl_mem_flags, size_type)
+```
+
+

samples/extensions/khr/ocean/fft_kernel.cl

@@ -0,0 +1,54 @@
+/*
+ * Copyright (c) 2024 Mobica Limited, Marcin Hajder
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+constant sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;
+
+typedef float2 complex;
+
+complex mul(complex c0, complex c1)
+{
+    return (complex)(c0.x * c1.x - c0.y * c1.y, c0.x * c1.y + c0.y * c1.x);
+}
+
+complex add(complex c0, complex c1)
+{
+    return (complex)(c0.x + c1.x, c0.y + c1.y);
+}
+
+// mode.x - 0-horizontal, 1-vertical
+// mode.y - subsequent count
+
+__kernel void fft_1D( int2 mode, int2 patch_info,
+    read_only image2d_t twiddle, read_only image2d_t src, write_only image2d_t dst )
+{
+    int2 uv = (int2)((int)get_global_id(0), (int)get_global_id(1));
+
+    int2 data_coords = (int2)(mode.y, uv.x * (1-mode.x) + uv.y * mode.x);
+    float4 data = read_imagef(twiddle, sampler, data_coords);
+
+    int2 pp_coords0 = (int2)(data.z, uv.y) * (1-mode.x) + (int2)(uv.x, data.z) * mode.x;
+    float2 p = read_imagef(src, sampler, pp_coords0).xy;
+
+    int2 pp_coords1 = (int2)(data.w, uv.y) * (1-mode.x) + (int2)(uv.x, data.w) * mode.x;
+    float2 q = read_imagef(src, sampler, pp_coords1).xy;
+
+    float2 w = (float2)(data.x, data.y);
+
+    //Butterfly operation
+    complex H = add(p,mul(w,q));
+
+    write_imagef(dst, uv, (float4)(H.x, H.y, 0, 1));
+}

samples/extensions/khr/ocean/init_spectrum.cl

@@ -0,0 +1,63 @@
+/*
+ * Copyright (c) 2024 Mobica Limited, Marcin Hajder
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+constant float PI = 3.14159265359f;
+constant sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;
+constant float GRAVITY = 9.81f;
+
+float4 gaussRND(float4 rnd)
+{
+    float u0 = 2.0*PI*rnd.x;
+    float v0 = sqrt(-2.0 * log(rnd.y));
+    float u1 = 2.0*PI*rnd.z;
+    float v1 = sqrt(-2.0 * log(rnd.w));
+
+    float4 ret = (float4)(v0 * cos(u0), v0 * sin(u0), v1 * cos(u1), v1 * sin(u1));
+    return ret;
+}
+
+// patch_info.x - ocean patch size
+// patch_info.y - ocean texture unified resolution
+// params.x - wind x
+// params.y - wind.y
+// params.z - amplitude
+// params.w - capillar supress factor
+
+kernel void init_spectrum( int2 patch_info, float4 params, read_only image2d_t noise, write_only image2d_t dst )
+{
+    int2 uv = (int2)((int)get_global_id(0), (int)get_global_id(1));
+    int res = patch_info.y;
+
+    float2 fuv = convert_float2(uv) - (float2)((float)(res-1)/2.f);
+    float2 k = (2.f * PI * fuv) / patch_info.x;
+    float k_mag = length(k);
+
+    if (k_mag < 0.00001) k_mag = 0.00001;
+
+    float wind_speed = length((float2)(params.x, params.y));
+    float4 params_n = params;
+    params_n.xy = (float2)(params.x/wind_speed, params.y/wind_speed);
+    float l_phillips = (wind_speed * wind_speed) / GRAVITY;
+    float4 rnd = clamp(read_imagef(noise, sampler, uv), 0.001f, 1.f);
+
+    float magSq = k_mag * k_mag;
+    float h0k = sqrt((params.z/(magSq*magSq)) * pow(dot(normalize(k), params_n.xy), 2.f) *
+                exp(-(1.0/(magSq * l_phillips * l_phillips))) * exp(-magSq*pow(params.w, 2.f)))/ sqrt(2.0);
+    float h0minusk = sqrt((params.z/(magSq*magSq)) * pow(dot(normalize(-k), params_n.xy), 2.f) *
+                exp(-(1.0/(magSq * l_phillips * l_phillips))) * exp(-magSq*pow(params.w, 2.f)))/ sqrt(2.0);
+    float4 gauss_random = gaussRND(rnd);
+    write_imagef(dst, uv, (float4)(gauss_random.xy*h0k, gauss_random.zw*h0minusk));
+}

samples/extensions/khr/ocean/inversion.cl

@@ -0,0 +1,30 @@
+/*
+ * Copyright (c) 2024 Mobica Limited, Marcin Hajder
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+constant sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;
+
+kernel void inversion( int2 patch_info, read_only image2d_t src0,
+    read_only image2d_t src1, read_only image2d_t src2, write_only image2d_t dst )
+{
+    int2 uv = (int2)((int)get_global_id(0), (int)get_global_id(1));
+    int res2 = patch_info.y * patch_info.y;
+
+    float x = read_imagef(src0, sampler, uv).x;
+    float y = read_imagef(src1, sampler, uv).x;
+    float z = read_imagef(src2, sampler, uv).x;
+
+    write_imagef(dst, uv, (float4)(x/res2, y/res2, z/res2, 1));
+}