Added new sample to demonstrate OpenCL-Vulkan interop with ocean surface simulation

samples/CMakeLists.txt

@@ -146,3 +146,4 @@ endmacro()
 
 add_subdirectory(core)
 add_subdirectory(extensions)
+add_subdirectory(vulkan)

samples/vulkan/CMakeLists.txt

@@ -0,0 +1,30 @@
+# Copyright (c) 2021 The Khronos Group Inc.
+#
+# 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.
+
+find_package(Vulkan)
+find_package(glfw3 HINTS "../../external")
+
+set(BUILD_VULKAN_SAMPLES TRUE)
+if(NOT VULKAN_FOUND)
+    message(STATUS "Skipping Vulkan Samples - Vulkan is not found.")
+    set(BUILD_VULKAN_SAMPLES FALSE)
+endif()
+if(NOT glfw3_FOUND)
+    message(STATUS "Skipping Vulkan Samples - GLFW is not found.")
+    set(BUILD_VULKAN_SAMPLES FALSE)
+endif()
+
+if(BUILD_VULKAN_SAMPLES)
+    add_subdirectory( ocean )
+endif()

samples/vulkan/ocean/CMakeLists.txt

@@ -0,0 +1,23 @@
+# 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.
+
+add_sample(
+    TARGET ocean_vk_ocl_interop
+    VERSION 300 # clCreateImageWithProperties
+    CATEGORY vulkan
+    SOURCES main.cpp ocean.cpp ocean.hpp ocean_util.hpp
+    SHADERS ocean.vert.spv ocean.frag.spv
+    KERNELS twiddle.cl time_spectrum.cl inversion.cl normals.cl fft_kernel.cl init_spectrum.cl
+    INCLUDES ${Vulkan_INCLUDE_DIR}
+    LIBS ${Vulkan_LIBRARY} glfw glm::glm)

samples/vulkan/ocean/README.md

@@ -0,0 +1,59 @@
+# Ocean surface simulation with Opencl and Vulkan interoperability
+
+[Ocean Simulation With OpenCL and Vulkan](ocean.png)
+
+## Sample Purpose
+
+This sample demonstrates how to share compute/render resources between OpenCL and Vulkan to simulate an ocean surface. If the cl_khr_external_memory extension is available and requested (through CLI options), some OpenCL images will be created through a file descriptor handle received with vkGetMemoryFdKHR. These images will then be used for ocean rendering. If cl_khr_external_memory is not available, additional copying from OpenCL buffers to Vulkan 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 Vulkan interoperability. 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:
+
+    -An OpenCL platform and Vulkan physical device are selected based on CLI options.
+    -OpenCL and Vulkan devices are prepared.
+    -A GLFW window, camera, and related keyboard event callbacks are created.
+    -Both shared and private resources for OpenCL and Vulkan are set up.
+
+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. While the simulation is in progress, each frame of the application performs the following general steps:
+
+    -Necessary Vulkan/OpenCL semaphores are signaled/waited.
+    -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::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/vulkan/ocean/fft_kernel.cl

@@ -0,0 +1,58 @@
+/*
+ * 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);
+
+
+    work_group_barrier(CLK_IMAGE_MEM_FENCE);
+
+
+    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).rg;
+
+    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).rg;
+
+    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/vulkan/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 = (float2)(get_global_id(0), get_global_id(1)) - (float)(res)/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/vulkan/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).r;
+    float y = read_imagef(src1, sampler, uv).r;
+    float z = read_imagef(src2, sampler, uv).r;
+
+    write_imagef(dst, uv, (float4)(x/res2, y/res2, z/res2, 1));
+}

samples/vulkan/ocean/main.cpp

@@ -0,0 +1,124 @@
+/*
+ * 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.
+ */
+
+#include "ocean.hpp"
+
+namespace {
+
+void glfw_keyboard(GLFWwindow* window, int key, int scancode, int action,
+                   int mods)
+{
+    auto app = (OceanApplication*)glfwGetWindowUserPointer(window);
+    app->keyboard(key, scancode, action, mods);
+}
+
+void glfw_mouse_event(GLFWwindow* window, int button, int action, int mods)
+{
+    auto app = (OceanApplication*)glfwGetWindowUserPointer(window);
+    app->mouse_event(button, action, mods);
+}
+
+void glfw_mouse_pos(GLFWwindow* window, double pX, double pY)
+{
+    auto app = (OceanApplication*)glfwGetWindowUserPointer(window);
+    app->mouse_pos(pX, pY);
+}
+
+void glfw_mouse_roll(GLFWwindow* window, double oX, double oY)
+{
+    auto app = (OceanApplication*)glfwGetWindowUserPointer(window);
+    app->mouse_roll(oX, oY);
+}
+
+} // anonymous namespace
+
+
+void OceanApplication::main_loop()
+{
+    glfwSetKeyCallback(window, glfw_keyboard);
+    glfwSetMouseButtonCallback(window, glfw_mouse_event);
+    glfwSetCursorPosCallback(window, glfw_mouse_pos);
+    glfwSetScrollCallback(window, glfw_mouse_roll);
+
+    while (!glfwWindowShouldClose(window))
+    {
+        draw_frame();
+        glfwPollEvents();
+    }
+
+    vkDeviceWaitIdle(device);
+}
+
+template <> auto cl::sdk::parse<CliOptions>()
+{
+    return std::make_tuple(
+        std::make_shared<TCLAP::ValueArg<size_t>>("", "window_width",
+                                                  "Window width", false, 1024,
+                                                  "positive integral"),
+        std::make_shared<TCLAP::ValueArg<size_t>>("", "window_height",
+                                                  "Window height", false, 1024,
+                                                  "positive integral"),
+        std::make_shared<TCLAP::ValueArg<bool>>(
+            "", "immediate", "Prefer VK_PRESENT_MODE_IMMEDIATE_KHR (no vsync)",
+            false, false, "boolean"),
+        std::make_shared<TCLAP::ValueArg<bool>>(
+            "", "linear", "Use linearly tiled images", false, false, "boolean"),
+        std::make_shared<TCLAP::ValueArg<bool>>("", "deviceLocalImages",
+                                                "Use device local images",
+                                                false, true, "boolean"),
+        std::make_shared<TCLAP::ValueArg<bool>>("", "useExternalMemory",
+                                                "Use cl_khr_external_memory",
+                                                false, true, "boolean"));
+}
+
+template <>
+CliOptions cl::sdk::comprehend<CliOptions>(
+    std::shared_ptr<TCLAP::ValueArg<size_t>> window_width,
+    std::shared_ptr<TCLAP::ValueArg<size_t>> window_height,
+    std::shared_ptr<TCLAP::ValueArg<bool>> immediate,
+    std::shared_ptr<TCLAP::ValueArg<bool>> linearImages,
+    std::shared_ptr<TCLAP::ValueArg<bool>> deviceLocalImages,
+    std::shared_ptr<TCLAP::ValueArg<bool>> useExternalMemory)
+{
+    return CliOptions{
+        window_width->getValue(),      window_height->getValue(),
+        immediate->getValue(),         linearImages->getValue(),
+        deviceLocalImages->getValue(), useExternalMemory->getValue()
+    };
+}
+
+int main(int argc, char** argv)
+{
+    OceanApplication app;
+
+    auto opts = cl::sdk::parse_cli<cl::sdk::options::Diagnostic,
+                                   cl::sdk::options::SingleDevice, CliOptions>(
+        argc, argv);
+
+    app.dev_opts = std::get<1>(opts);
+    app.app_opts = std::get<2>(opts);
+
+    try
+    {
+        app.run();
+    } catch (const std::exception& e)
+    {
+        fprintf(stderr, "%s\n", e.what());
+        return EXIT_FAILURE;
+    }
+
+    return EXIT_SUCCESS;
+}