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

CMakeLists.txt

@@ -26,9 +26,14 @@ include(CMakeDependentOption)
 option(OPENCL_SDK_BUILD_UTILITY_LIBRARIES "Build utility libraries" ON)
 cmake_dependent_option(OPENCL_SDK_BUILD_SAMPLES "Build sample code" ON OPENCL_SDK_BUILD_UTILITY_LIBRARIES OFF)
 cmake_dependent_option(OPENCL_SDK_BUILD_OPENGL_SAMPLES "Build OpenCL-OpenGL interop sample code" OFF OPENCL_SDK_BUILD_SAMPLES OFF)
-cmake_dependent_option(OPENCL_SDK_BUILD_VULKAN_SAMPLES "Build OpenCL-Vulkan interop sample code" ON OPENCL_SDK_BUILD_SAMPLES OFF)
+cmake_dependent_option(OPENCL_SDK_BUILD_VULKAN_SAMPLES "Build OpenCL-Vulkan interop sample code" OFF OPENCL_SDK_BUILD_SAMPLES OFF)
 cmake_dependent_option(OPENCL_SDK_TEST_SAMPLES "Add CTest to samples (where applicable)" ON OPENCL_SDK_BUILD_SAMPLES OFF)
 
+find_package(Vulkan)
+if(VULKAN_FOUND)
+    cmake_dependent_option(OPENCL_SDK_BUILD_VULKAN_SAMPLES "Build OpenCL-Vulkan interop sample code" ON OPENCL_SDK_BUILD_SAMPLES OFF)
+endif()
+
 option(OPENCL_SDK_BUILD_CLINFO "Build clinfo utility" ON)
 if (("${CMAKE_SYSTEM_NAME}" STREQUAL "Darwin") AND ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU"))
   string(APPEND CMAKE_CXX_FLAGS " -stdlib=libstdc++")

README.md

@@ -61,7 +61,7 @@ If CMake is not provided by your build system or OS package manager, please cons
              -B ./OpenCL-SDK/build -S ./OpenCL-SDK
        cmake --build ./OpenCL-SDK/build --target install
 
-Samples that make use of OpenGL interop are disabled by default to reduce
+Samples that make use of OpenGL or Vulkan interop are disabled by default to reduce
 the number of dependencies for most users. They can be enabled using the
 `OPENCL_SDK_BUILD_OPENGL_SAMPLES` CMake option.
 

cmake/Dependencies.cmake

@@ -51,6 +51,20 @@ if(OPENCL_SDK_BUILD_SAMPLES)
       list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/Dependencies/${DEP}")
       include(${DEP})
     endforeach()
+      if(OPENCL_SDK_BUILD_VULKAN_SAMPLES AND SFML_VERSION VERSION_GREATER_EQUAL 2.6)
+        foreach(DEP IN ITEMS X11 SFML)
+          list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/Dependencies/${DEP}")
+          include(${DEP})
+        endforeach()
+    endif()
+  else()
+    if(OPENCL_SDK_BUILD_VULKAN_SAMPLES AND SFML_VERSION VERSION_GREATER_EQUAL 2.6)
+      foreach(DEP IN ITEMS X11 glm SFML)
+        list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/Dependencies/${DEP}")
+        include(${DEP})
+      endforeach()
+    endif()
+
   endif(OPENCL_SDK_BUILD_OPENGL_SAMPLES)
 
   if(OPENCL_SDK_BUILD_VULKAN_SAMPLES)

cmake/Dependencies/glfw/glfw.cmake

@@ -0,0 +1,46 @@
+if(NOT DEPENDENCIES_FORCE_DOWNLOAD AND NOT EXISTS "${CMAKE_CURRENT_BINARY_DIR}/_deps/glfw-external-src")
+  find_package(glfw3 CONFIG)
+  # To avoid every test depening on GLFW define their deps using
+  #
+  # add_sample(
+  # LIBS
+  #   $<$<TARGET_EXISTS:glfw>:glfw>
+  # INCLUDES
+  #   $<$<NOT:$<TARGET_EXISTS:glfw>>:"${GLFW_INCLUDE_DIRS}">
+  # )
+  #
+  # we create the INTERFACE target in case it didn't exist.
+  if(glfw3_FOUND AND NOT TARGET glfw)
+    add_library(glfw INTERFACE)
+    target_include_directories(glfw INTERFACE "${GLFW_INCLUDE_DIRS}")
+  endif()
+endif()
+
+if(NOT (glfw3_FOUND OR TARGET glfw))
+  if(NOT EXISTS "${CMAKE_CURRENT_BINARY_DIR}/_deps/glfw-external-src")
+    if(DEPENDENCIES_FORCE_DOWNLOAD)
+      message(STATUS "DEPENDENCIES_FORCE_DOWNLOAD is ON. Fetching glfw.")
+    else()
+      message(STATUS "Fetching glfw.")
+    endif()
+    message(STATUS "Adding glfw subproject: ${CMAKE_CURRENT_BINARY_DIR}/_deps/glfw-external-src")
+  endif()
+  cmake_minimum_required(VERSION 3.11)
+  include(FetchContent)
+  set(GLFW_BUILD_EXAMPLES OFF CACHE BOOL "Build the GLFW example programs.")
+  set(GLFW_BUILD_TESTS OFF CACHE BOOL "Build the GLFW test programs.")
+  FetchContent_Declare(
+    glfw-external
+    GIT_REPOSITORY      https://github.com/glfw/glfw
+    GIT_TAG             3.3.6 # 7d5a16ce714f0b5f4efa3262de22e4d948851525
+  )
+  FetchContent_MakeAvailable(glfw-external)
+  set_target_properties(glfw
+    PROPERTIES
+      RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_BINDIR}"
+      ARCHIVE_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}"
+      LIBRARY_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}"
+      INSTALL_RPATH "${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_LIBDIR}"
+      FOLDER "Dependencies"
+  )  
+endif()

lib/include/CL/Utils/File.hpp

@@ -32,5 +32,8 @@ namespace util {
 
     std::string UTILSCPP_EXPORT read_exe_relative_text_file(
         const char* const filename, cl_int* const error = nullptr);
+
+    std::vector<unsigned char> UTILSCPP_EXPORT read_exe_relative_binary_file(
+        const char* const filename, cl_int* const error = nullptr);
 }
 }

lib/src/Utils/File.cpp

@@ -170,3 +170,27 @@ std::string cl::util::read_exe_relative_text_file(const char* const filename,
     }
     return result;
 }
+
+std::vector<unsigned char>
+cl::util::read_exe_relative_binary_file(const char* const filename,
+                                        cl_int* const error)
+{
+    std::vector<unsigned char> result;
+    cl_int err = CL_SUCCESS;
+    std::string exe_folder = executable_folder(&err);
+    if (err != CL_SUCCESS)
+    {
+        detail::errHandler(CL_UTIL_FILE_OPERATION_ERROR, error,
+                           "Failed to query exe folder!");
+        return result;
+    }
+    result = read_binary_file((exe_folder + "/" + filename).c_str(), &err);
+    if (err != CL_SUCCESS)
+    {
+        result.clear();
+        detail::errHandler(CL_UTIL_FILE_OPERATION_ERROR, error,
+                           "Unable to read file!");
+        return result;
+    }
+    return result;
+}

samples/CMakeLists.txt

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

samples/vulkan/CMakeLists.txt

@@ -0,0 +1,18 @@
+# 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.
+
+
+if(OPENCL_SDK_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})

samples/vulkan/ocean/README.md

@@ -0,0 +1,73 @@
+# 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.
+
+Available CLI options are as follows:
+
+    --window_width, specifies initial window width
+    --window_height, specifies initial window window_height
+    --vulkan_device, requests number of vulkan physical device
+    --immediate, requests preference of VK_PRESENT_MODE_IMMEDIATE_KHR (no vsync)
+    --linear, requests use of linearly tiled images 
+    --deviceLocalImages, requests use of device local images
+    --useExternalMemory, requests use of cl_khr_external_memory
+
+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:
+
+    -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::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/vulkan/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/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 = 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));
+}