diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index 01731549..347aded1 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -184,6 +184,17 @@ jobs: fetch-depth: 0 submodules: recursive + - name: Install samples dependencies + run: | + if [[ "${{ matrix.BIN }}" == "64" ]]; then + apt-get update -qq; + apt-get install -y libvulkan-dev; + else + dpkg --add-architecture i386; + apt-get update -qq; + apt-get install -y libvulkan-dev:i386; + fi + - name: Configure, package & install OpenCL-Headers run: $CMAKE_EXE -G "${{matrix.CONF.GEN}}" @@ -392,6 +403,8 @@ jobs: INTEL_OCL_URL: https://github.com/intel/llvm/releases/download/2023-WW27/win-oclcpuexp-2023.16.6.0.28_rel.zip INTEL_TBB_URL: https://github.com/oneapi-src/oneTBB/releases/download/v2021.10.0/oneapi-tbb-2021.10.0-win.zip IMAGE_INTEL_PREFIX: C:\Tools\Intel + VULKAN_SDK_URL: https://sdk.lunarg.com/sdk/download/1.3.261.1/windows/VulkanSDK-1.3.261.1-Installer.exe + VULKAN_SDK: C:/VulkanSDK/1.3.261.1 steps: - name: Cache Ninja install @@ -428,6 +441,18 @@ jobs: New-Item -Type File HKLM:\SOFTWARE\Khronos\OpenCL\Vendors | Out-Null; ` Set-ItemProperty -Path HKLM:\SOFTWARE\Khronos\OpenCL\Vendors -Name ${env:IMAGE_INTEL_PREFIX}\oclcpuexp\intelocl64.dll -Type DWord -Value 0; + - name: Install samples dependencies + run: | + Invoke-WebRequest ${env:VULKAN_SDK_URL} -OutFile vulkan-sdk-installer.exe + .\vulkan-sdk-installer.exe --accept-licenses --default-answer --confirm-command install com.lunarg.vulkan.32bit + Remove-Item vulkan-sdk-installer.exe + if ('${{ matrix.BIN }}' -eq 'x64') + { + echo "Vulkan_LIB_DIR=$env:VULKAN_SDK/Lib" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append + } else { + echo "Vulkan_LIB_DIR=$env:VULKAN_SDK/Lib32" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append + } + - name: Cache dependencies (vcpkg) if: matrix.DEPS == 'vcpkg' id: vcpkg-install @@ -472,6 +497,8 @@ jobs: -D CMAKE_C_STANDARD=${{matrix.STD.C}} ` -D CMAKE_CXX_STANDARD=${{matrix.STD.CXX}} ` -D CMAKE_INSTALL_PREFIX=${env:GITHUB_WORKSPACE}\install ` + -D Vulkan_INCLUDE_DIR=${env:VULKAN_SDK}/Include ` + -D Vulkan_LIBRARY=${env:Vulkan_LIB_DIR}/vulkan-1.lib ` -S ${env:GITHUB_WORKSPACE} ` -B ${env:GITHUB_WORKSPACE}\build if ($LASTEXITCODE -ne 0) { throw "Configuring OpenCL-SDK failed." } @@ -501,6 +528,8 @@ jobs: -D CMAKE_CXX_STANDARD=${{matrix.STD.CXX}} ` -D CMAKE_EXE_LINKER_FLAGS=/INCREMENTAL ` -D CMAKE_INSTALL_PREFIX=${env:GITHUB_WORKSPACE}\install ` + -D Vulkan_INCLUDE_DIR=${env:VULKAN_SDK}/Include ` + -D Vulkan_LIBRARY=${env:Vulkan_LIB_DIR}/vulkan-1.lib ` -S ${env:GITHUB_WORKSPACE} ` -B ${env:GITHUB_WORKSPACE}\build if ($LASTEXITCODE -ne 0) { throw "Configuring OpenCL-SDK failed." } @@ -539,12 +568,14 @@ jobs: if: matrix.BIN != 'x86' working-directory: ${{runner.workspace}}/OpenCL-SDK/build run: | + $EXCLUDE_REGEX = 'externalmemory.*' foreach ($Config in 'Release','Debug') { & ctest ` --build-config ${Config} ` --output-on-failure ` --no-tests=error ` - --parallel ${env:NUMBER_OF_PROCESSORS} + --parallel ${env:NUMBER_OF_PROCESSORS} ` + --exclude-regex "$EXCLUDE_REGEX" if ($LASTEXITCODE -ne 0) { throw "Running OpenCL-SDK tests in $Config failed." } } @@ -655,6 +686,10 @@ jobs: brew install pocl cmake --version + - name: Install samples dependencies + run: | + brew install vulkan-loader + - name: Install dependencies (Homebrew) if: matrix.DEPS == 'system' run: brew install tclap glm glew sfml mesa-glu @@ -702,8 +737,10 @@ jobs: - name: Test working-directory: ${{runner.workspace}}/OpenCL-SDK/build run: | - OCL_ICD_VENDORS=/usr/local/etc/OpenCL/vendors ctest -C Debug --output-on-failure --no-tests=error --parallel `sysctl -n hw.logicalcpu` - OCL_ICD_VENDORS=/usr/local/etc/OpenCL/vendors ctest -C Release --output-on-failure --no-tests=error --parallel `sysctl -n hw.logicalcpu` + # Incompatible Vulkan Driver + export EXCLUDE_REGEX="externalmemory.*" + OCL_ICD_VENDORS=/usr/local/etc/OpenCL/vendors ctest -C Debug --output-on-failure --no-tests=error --parallel `sysctl -n hw.logicalcpu` --exclude-regex "$EXCLUDE_REGEX" + OCL_ICD_VENDORS=/usr/local/etc/OpenCL/vendors ctest -C Release --output-on-failure --no-tests=error --parallel `sysctl -n hw.logicalcpu` --exclude-regex "$EXCLUDE_REGEX" - name: Install run: | diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 4fd0ab1a..1ab2a178 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -34,7 +34,7 @@ set(BUILD_SHARED_LIBS OFF CACHE BOOL "Global flag to cause add_library() to crea # Fetch dependencies if(OPENCL_SDK_BUILD_SAMPLES) - foreach(DEP IN ITEMS cargs TCLAP Stb) + foreach(DEP IN ITEMS cargs TCLAP Stb Vulkan) list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/Dependencies/${DEP}") include(${DEP}) endforeach() diff --git a/cmake/Dependencies/Vulkan/Vulkan.cmake b/cmake/Dependencies/Vulkan/Vulkan.cmake new file mode 100644 index 00000000..ffdcdbd0 --- /dev/null +++ b/cmake/Dependencies/Vulkan/Vulkan.cmake @@ -0,0 +1 @@ +find_package(Vulkan REQUIRED) diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index 28bbe937..b3a92a4b 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -33,11 +33,12 @@ CHECK_LIBRARY_EXISTS(m sin "" HAVE_LIB_M) # KERNELS ... # optional, specifies kernel files for the sample # INCLUDES ... # optional, specifies additional include directories for the sample # LIBS ... # optional, specifies additional libraries for the sample +# DEFINITIONS # optional, specifies additional compile definitions for the sample # ) macro(add_sample) set(options TEST) set(one_value_args TARGET VERSION CATEGORY) - set(multi_value_args SOURCES KERNELS SHADERS INCLUDES LIBS) + set(multi_value_args SOURCES KERNELS SHADERS INCLUDES LIBS DEFINITIONS) cmake_parse_arguments(OPENCL_SAMPLE "${options}" "${one_value_args}" "${multi_value_args}" ${ARGN} @@ -75,6 +76,7 @@ macro(add_sample) CL_HPP_MINIMUM_OPENCL_VERSION=${OPENCL_SAMPLE_VERSION} CL_HPP_ENABLE_EXCEPTIONS $<$:_CRT_SECURE_NO_WARNINGS> # TODO: remove + ${OPENCL_SAMPLE_DEFINITIONS} ) set_target_properties(${OPENCL_SAMPLE_TARGET} diff --git a/samples/extensions/khr/CMakeLists.txt b/samples/extensions/khr/CMakeLists.txt index 82d4426f..4bf194d7 100644 --- a/samples/extensions/khr/CMakeLists.txt +++ b/samples/extensions/khr/CMakeLists.txt @@ -12,6 +12,7 @@ # See the License for the specific language governing permissions and # limitations under the License. +add_subdirectory(externalmemory) add_subdirectory(histogram) if(OPENCL_SDK_BUILD_OPENGL_SAMPLES) add_subdirectory(conway) diff --git a/samples/extensions/khr/externalmemory/CMakeLists.txt b/samples/extensions/khr/externalmemory/CMakeLists.txt new file mode 100644 index 00000000..78ea8a8c --- /dev/null +++ b/samples/extensions/khr/externalmemory/CMakeLists.txt @@ -0,0 +1,31 @@ +# 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. + +add_sample( + TEST + TARGET externalmemory + VERSION 300 + SOURCES main.c + KERNELS external_saxpy.cl + LIBS Vulkan::Vulkan + DEFINITIONS $<$:VK_USE_PLATFORM_WIN32_KHR>) + +add_sample( + TEST + TARGET externalmemorycpp + VERSION 300 + SOURCES main.cpp + KERNELS external_saxpy.cl + LIBS Vulkan::Vulkan + DEFINITIONS $<$:VK_USE_PLATFORM_WIN32_KHR>) diff --git a/samples/extensions/khr/externalmemory/README.md b/samples/extensions/khr/externalmemory/README.md new file mode 100644 index 00000000..78ca63f6 --- /dev/null +++ b/samples/extensions/khr/externalmemory/README.md @@ -0,0 +1,257 @@ +# External Memory Sample + +## Sample purpose +External devices resources can be shared across GPU APIs. This can specially come in handy when developing graphical applications, as usually we have specialized APIs for graphics (like OpenGL or the lower-level-API Vulkan) that are used for rendering and the more general APIs (like OpenCL, SYCL, etc). This sample showcases an OpenCL program that interacts with the Vulkan API by sharing buffers. For one that actually does rendering, the [open_cl_interop](https://github.com/KhronosGroup/Vulkan-Samples/tree/main/samples/extensions/open_cl_interop) sample should be consulted. + +## Key APIs and Concepts +### Kernel logic +The kernel used in this sample is a saxpy, i.e. performs the vector operation $a*x+y$ where $x$ and $y$ are the input vectors and $a$ is a scalar. This simple kernel was chosen because the main purpose of the example is to showcase the buffer sharing between the OpenCL and Vulkan APIs, rather than showing off some complex kernel implementation. + +### Create Vulkan instance with the necessary extensions enabled +The Vulkan function `vkCreateInstance` creates a new Vulkan instance (object gathering the application's state), which later can be used to query the physical devices available on the system for our program. When calling to this function, a `VkInstanceCreateInfo` object must be passed in order to tell the Vulkan API some characteristics of the application. In this sample, one of the main pieces of information passed to the named function is a list of Vulkan instance extensions to be enabled: +- `VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME` for exporting non-Vulkan handles from Vulkan buffers. +- `VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_EXTENSION_NAME` for also being able to query the properties of physical devices (needed for obtaining the devices' UUIDs). + +### Find an OpenCL device Vulkan-compatible +In the context of a given OpenCL program, for a device to be compatible with the Vulkan API there are three main requirements: +- It has to be recognized by Vulkan as a physical device, that is, Vulkan must report the existence of a physical device with the same UUID than the selected OpenCL device's. In Vulkan, with `vkGetPhysicalDeviceProperties2` we can get the properties of a physical device, among which is included the `deviceUUID` attribute storing the UUID of the corresponding device. For OpenCL, we can query the device's UUID by calling `clGetDeviceInfo` (or the C++ wrapper `cl::Device::getInfo<>()`) with the `CL_DEVICE_UUID_KHR` value as `cl_device_info` parameter. + - Beware the query of the UUID in OpenCL/Vulkan cannot be done without the device supporting the `cl_khr_device_uuid`/`VK_KHR_get_physical_device_properties2`. +- It must support the Vulkan device extensions needed for the program at hand. In this occasion, we need the Vulkan device to support exporting non-Vulkan handles from Vulkan memory objects (e.g. buffers). The `vkEnumerateDeviceExtensionProperties` function is used for querying the Vulkan device extensions supported by a given physical device. +- It also needs to support the Khronos extension `cl_khr_external_memory_opaque_fd` for Linux systems or `cl_khr_external_memory_win32` for Windows. With the C API, The function `clGetDeviceInfo` called with the parameter `CL_DEVICE_EXTENSIONS` provides information about whether this extension is supported by the OpenCL device. The C++ API (Utils library) provides the function `cl::util::supports_extension`, with which this check can be done easier. + + _Note: The `cl_khr_external_memory` extension requires OpenCL 3.0, which we make sure to check that is indeed supported on the device before compiling the OpenCL kernel._ + +Once a suitable Vulkan physical device (and its correspondent OpenCL device) has been found, we can create a Vulkan device object from it with `vkCreateDevice`. We must set the `ppEnabledExtensionNames` attribute of the `VkDeviceCreateInfo` passed to the said function with the names of the required Vulkan device extensions (that we already checked the device supports) in order for them to be enabled on the device. + +### Create Vulkan buffers for external sharing +When creating the Vulkan buffer objects for our application, we must make explicit that those buffers are going to be shared with an external API. The way of doing this can be summarized into the following steps: +- Before starting to allocate Vulkan memory objects, we need to ensure that the external memory handle type needed for importing Vulkan memory objects is supported by the device, both in OpenCL and in Vulkan APIs. +The mapping between Vulkan and OpenCL handle types is as follows: + + | Vulkan external memory handle type | OpenCL external memory handle type | + | ------------------------------------------------------------------ | ------------------------------------------------------------- | + | `VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR` | `CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR` | + | `VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR` | `CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR` | + | `VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT_KHR` | `CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_BIT_KHR` | + + The first row contains the handle types used for Linux, while for Windows platforms the handle types used are either the ones from the second or third row. + + To check whether the OpenCL device supports the memory handle we use `clGetDeviceInfo` with the `CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR` value as `cl_device_info` parameter in order to get a list of supported external memory handle types. + + For Vulkan, we can request a `VkExternalBufferProperties` object containing this information by calling to `vkGetPhysicalDeviceExternalBufferProperties`. + +- We now create our Vulkan buffer objects. We first initialize a `VkExternalMemoryBufferCreateInfo` structure with the necessary information for the buffers bounded to the exported memory. It is **mandatory** when creating a Vulkan buffer that will be bound to exported/imported memory to pass a **non-null** value for the **`handleTypes`** field of this info structure. A pointer to this object is then added as the `pNext` field of a `VkBufferCreateInfo` structure, which contains the information for creating Vulkan buffers (that are not necessarily bounded to external memory). We finally create our buffers by calling `vkCreateBuffer`. + +- The next step is to allocate device memory. This is done with the function `vkAllocateMemory`, which needs a `VkMemoryAllocateInfo` parameter. The key information to set up when allocating external memory is the `pNext` field, pointing to a `VkExportMemoryAllocateInfo` structure which `handleTypes` field specifies the handle types that may be exported. + +- After allocating the device memory, it is only left to bind it to the buffer objects with `vkBindBufferMemory` and to map the latter into the application address space with `vkMapMemory`. If the buffer objects are to be mapped in their entirety, we can use `VK_WHOLE_SIZE` as the `size` parameter of `vkMapMemory`. After mapping the buffer objects we obtain host-accessible pointers to the beginning of the mapped ranges and we can just copy the contents of the host arrays to those ranges. + +### Initialize OpenCL buffers from external API +The key point when initializing OpenCL buffers from external memory is that we need a file descriptor associated to this external memory in order to access it from the OpenCL API. In the Vulkan API we can get such file descriptor by making use of the function `vkGetMemoryFdKHR` provided by the `VK_KHR_external_memory_fd` extension. + +Being provided by an extension, we need to obtain a function pointer to it by calling to `vkGetDeviceProcAddr`. We can then call `vkGetMemoryFdKHR` with a `VkMemoryGetFdInfoKHR` parameter containing the information about the memory range for which we want to obtain a file descriptor: + - `memory` field containing the pointer to the said range + - `handleType` field with the same Vulkan external memory handle type used in the `VkExportMemoryAllocateInfo` structure when memory was allocated. + +Once we have the file descriptor, we can initialize an array of `cl_mem_properties` with the following entries: +- The OpenCL external memory handle type to use. +- The file descriptor previously obtained for the Vulkan memory range. +- A list of devices to which these properties apply. This list must start with an entry containing the macro `CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR`, followed by as many entries as devices in the list containing the corresponding `cl_device_id` objects. The list must end with an entry containing the macro `CL_DEVICE_HANDLE_LIST_END_KHR`. +- A $0$ indicating the end of the array. + +_Note: With the C++ API we can obtain the `cl_device_id` object from a `cl::Device device` wrapper by using the `()` operator._ + +This array of properties is then passed to `clCreateBufferWithProperties` (or to the C++ constructor of `cl::Buffer`). When creating OpenCL buffer objects from external memory there are a couple of restrictions in the parameters allowed for `clCreateBufferWithProperties`/`cl::Buffer::Buffer()`, namely: +- The `flags` parameter used to specify usage information for the buffer must not include `CL_MEM_USE_HOST_PTR`, `CL_MEM_ALLOC_HOST_PTR`, or `CL_MEM_COPY_HOST_PTR`. +- The `host_ptr` argument must be null. + +From this point on the OpenCL API functions are called as usual. + +## Application flow +### Overview +1. Parse user options. +2. Initialize Vulkan instance. +3. Find an OpenCL Vulkan-compatible device. +4. Create a Vulkan device object from the physical device selected enabling the required extensions on it. +5. Check that the OpenCL device supports the necessary Khronos extensions. +6. Create Vulkan's buffer objects for sharing them with an external API. +7. Query the requirements for memory to be exportable. Allocate memory, bind buffers to memory and map the former to the Vulkan address space. Copy input from host to Vulkan memory objects. +8. Query the file descriptors correspondent to Vulkan's memory ranges mapped and initialize OpenCL buffers from them. +9. Enqueue kernel call to saxpy. +10. Fetch and validate result. +11. Free resources. + +## Used API surface +### C +```c +CL_BLOCKING +CL_CONTEXT_PLATFORM +CL_DEVICE_EXTENSIONS +CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR +CL_DEVICE_HANDLE_LIST_KHR +CL_DEVICE_HANDLE_LIST_END_KHR +CL_DEVICE_NAME +CL_DEVICE_PLATFORM +CL_DEVICE_TYPE_ALL +CL_HPP_TARGET_OPENCL_VERSION +CL_INVALID_ARG_VALUE +CL_INVALID_VALUE +CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR +CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR +CL_KERNEL_WORK_GROUP_SIZE +CL_KHR_EXTERNAL_MEMORY_OPAQUE_FD_EXTENSION_NAME +CL_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME +CL_MEM_READ_ONLY +CL_MEM_READ_WRITE +CL_PLATFORM_VENDOR +CL_PROFILING_COMMAND_END +CL_PROFILING_COMMAND_START +CL_QUEUE_PROFILING_ENABLE +CL_QUEUE_PROPERTIES +CL_SUCCESS +CL_UUID_SIZE_KHR +cl_command_queue +cl_command_queue_properties +cl_context +cl_context_properties +cl_device_id +cl_event +cl_float +cl_int +cl_kernel +cl_external_memory_handle_type_khr +cl_khr_external_memory_opaque_fd +cl_khr_external_memory_win32 +cl_mem +cl_mem_properties +cl_platform_id +cl_program +cl_sdk_fill_with_random_ints_range(pcg32_random_t*, cl_int*, size_t, cl_int, cl_int) +cl_sdk_options_Diagnostic +cl_sdk_options_SingleDevice +cl_uint +cl_uchar +cl_ulong +cl_util_build_program(cl_program, cl_device_id, char*) +cl_util_get_device(cl_uint, cl_uint, cl_device_type, cl_int*) +cl_util_get_event_duration(cl_event, cl_profiling_info, cl_profiling_info, cl_int*) +cl_util_print_device_info*(cl_device_id) +cl_util_print_error(cl_int) +cl_util_read_text_file(char*const, size_t*const, cl_int*) +clCreateBufferWithProperties(cl_context, cl_mem_properties*, cl_mem_flags, size_t, void*, cl_int*) +clCreateCommandQueueWithProperties(cl_context, cl_device_id, cl_queue_properties*, cl_int*) -> OpenCL >= 2.0 +clCreateContext(cl_context_properties*, cl_uint, cl_device_id*, void *(char*, void*,size_t, void*), void*, cl_int*) +clCreateKernel(cl_program, char*, cl_int*) +clGetKernelWorkGroupInfo(cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void*, size_t*) +clCreateProgramWithSource(cl_context, cl_uint, char**, size_t*, cl_int*) +clEnqueueNDRangeKernel(cl_command_queue, cl_kernel, cl_uint, size_t*, size_t*, size_t*, cl_uint, cl_event*, cl_event*) +clEnqueueReadBuffer(cl_command_queue, cl_mem, cl_bool, size_t, size_t, void*, cl_uint, cl_event*, cl_event*) +clGetDeviceIDs(cl_platform_id, cl_device_type, cl_uint, cl_device_id*, cl_uint*) +clGetDeviceInfo(cl_device_id, cl_device_info, size_t, void*, size_t*) +clGetPlatformIDs(cl_uint, cl_platform_id*, cl_uint*) +clReleaseCommandQueue(cl_command_queue) +clReleaseContext(cl_context) +clReleaseKernel(cl_kernel) +clReleaseMemObject(cl_mem) +clReleaseProgram(cl_program) +clSetKernelArg(cl_kernel, cl_uint, size_t, void *) +clWaitForEvents(cl_uint, cl_event*) +``` + +### C++ +```c++ +cl::Buffer::Buffer(const Context&, const vector&, IteratorType, IteratorType, bool, bool=false, cl_int*=NULL) +cl::BuildError +cl::CommandQueue::CommandQueue(const cl::Context&, const Device&,cl::QueueProperties, cl_int*=NULL) +cl::Context +cl::Device::Device() +cl::EnqueueArgs::EnqueueArgs(cl::CommandQueue&, cl::NDRange, cl::NDRange) +cl::Error +cl::Event +cl::KernelFunctor::KernelFunctor(const Program&, const string, cl_int*=NULL) +cl::NDRange::NDRange(size_t, size_t) +cl::Platform::Platform() +cl::Platform::Platform(cl::Platform) +cl::Platform::get(vector*) +cl::Program::Program(cl::Program) +cl::WaitForEvents(const vector&) +cl::copy(const CommandQueue&, const cl::Buffer&, IteratorType, IteratorType) +cl::sdk::comprehend() +cl::sdk::fill_with_random() +cl::sdk::get_context(cl_uint, cl_uint, cl_device_type, cl_int*) +cl::sdk::parse() +cl::sdk::parse_cli() +cl::sdk::options::Diagnostic +cl::sdk::options::SingleDevice +cl::string::string(cl::string) +cl::util::Error +cl::util::get_duration(cl::Event&) +cl::util::supports_extension(const cl::Device&, const cl::string&) +``` + +### Vulkan +```c +PFN_vkCreateDevice(VkPhysicalDevice, const VkDeviceCreateInfo*, const VkAllocationCallbacks*, VkDevice*) +VK_BUFFER_USAGE_TRANSFER_DST_BIT +VK_BUFFER_USAGE_TRANSFER_SRC_BIT +VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR +VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR +VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME +VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME +VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME +VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME +VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_EXTENSION_NAME +VK_MAKE_VERSION +VK_MEMORY_PROPERTY_HOST_COHERENT_BIT +VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT +VK_SHARING_MODE_EXCLUSIVE +VK_STRUCTURE_TYPE_APPLICATION_INFO +VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO +VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO +VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO +VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO +VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO +VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO +VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR +VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES_KHR +VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2_KHR +VK_SUCCESS +VK_WHOLE_SIZE +VkApplicationInfo +VkBuffer +VkBufferCreateInfo +VkDevice +VkDeviceMemory +VkDeviceQueueCreateInfo +VkExportMemoryAllocateInfo +VkExtensionProperties +VkExternalMemoryBufferCreateInfo +VkExternalMemoryHandleTypeFlagBits +VkInstance +VkInstanceCreateInfo +VkMemoryAllocateInfo +VkMemoryGetFdInfoKHR +VkMemoryPropertyFlags +VkMemoryRequirements +VkPhysicalDevice +VkPhysicalDeviceIDPropertiesKHR +VkPhysicalDeviceMemoryProperties +VkPhysicalDeviceProperties2KHR +VkPhysicalDeviceProperties +VkResult +vkAllocateMemory(VkDevice, const VkMemoryAllocateInfo*, const VkAllocationCallbacks*, VkDeviceMemory*) +vkBindBufferMemory(VkDevice, VkBuffer, VkDeviceMemory, VkDeviceSize) +vkCreateBuffer(VkDevice, const VkBufferCreateInfo*, const VkAllocationCallbacks*, VkBuffer*) +vkCreateInstance(const VkInstanceCreateInfo*, const VkAllocationCallbacks*, VkInstance*) +vkDestroyBuffer(VkDevice, VkBuffer, const VkAllocationCallbacks*) +vkEnumerateDeviceExtensionProperties(VkPhysicalDevice, const char*, uint32_t*, VkExtensionProperties*) +vkEnumeratePhysicalDevices(VkInstance, uint32_t*, VkPhysicalDevice*) +vkFreeMemory(VkDevice, VkDeviceMemory, const VkAllocationCallbacks*) +vkGetBufferMemoryRequirements(VkDevice, VkBuffer, VkMemoryRequirements*) +vkGetDeviceProcAddr(VkDevice, const char*) +vkGetMemoryFdKHR(VkDevice, const VkMemoryGetFdInfoKHR*, int*) +vkGetPhysicalDeviceMemoryProperties(VkPhysicalDevice, VkPhysicalDeviceMemoryProperties*) +vkGetPhysicalDeviceProperties2(VkPhysicalDevice, VkPhysicalDeviceProperties2) +vkMapMemory(VkDevice, VkDeviceMemory, VkDeviceSize, VkDeviceSize, VkMemoryMapFlags, void**) +vkUnmapMemory(VkDevice, VkDeviceMemory) +``` diff --git a/samples/extensions/khr/externalmemory/external_saxpy.cl b/samples/extensions/khr/externalmemory/external_saxpy.cl new file mode 100644 index 00000000..c5780872 --- /dev/null +++ b/samples/extensions/khr/externalmemory/external_saxpy.cl @@ -0,0 +1,8 @@ +__kernel void saxpy(float a, + __global float* x, + __global float* y) +{ + int gid = get_global_id(0); + + y[gid] = fma(a, x[gid], y[gid]); +} diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c new file mode 100644 index 00000000..c49ff2c6 --- /dev/null +++ b/samples/extensions/khr/externalmemory/main.c @@ -0,0 +1,654 @@ +/* + * Copyright (c) 2023 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. + */ + +// OpenCL SDK includes. +#include +#include +#include +#include + +// OpenCL Utils includes. +#include +#include +#include + +// Vulkan includes. +#include + +// Vulkan utils includes. +#include "vulkan_utils.h" + +// Standard header includes. +#include +#include +#include +#include + +// Sample-specific option. +struct options_Saxpy +{ + size_t length; +}; + +// Add option to CLI-parsing SDK utility for input length. +cag_option SaxpyOptions[] = { { .identifier = 'l', + .access_letters = "l", + .access_name = "length", + .value_name = "(positive integer)", + .description = "Length of input" } }; + +ParseState parse_SaxpyOptions(const char identifier, + cag_option_context* cag_context, + struct options_Saxpy* opts) +{ + const char* value; + switch (identifier) + { + case 'l': + if (0 != (value = cag_option_get_value(cag_context))) + { + opts->length = strtoul(value, NULL, 0); + return ParsedOK; + } + else + return ParseError; + } + return NotParsed; +} + +cl_int parse_options(int argc, char* argv[], + struct cl_sdk_options_Diagnostic* diag_opts, + struct options_Saxpy* saxpy_opts) +{ + cl_int error = CL_SUCCESS; + struct cag_option *opts = NULL, *tmp = NULL; + size_t n = 0; + + // Prepare options array. + MEM_CHECK(opts = add_CLI_options(opts, &n, DiagnosticOptions, + CAG_ARRAY_SIZE(DiagnosticOptions)), + error, end); + opts = tmp; + MEM_CHECK(tmp = add_CLI_options(opts, &n, SaxpyOptions, + CAG_ARRAY_SIZE(SaxpyOptions)), + error, end); + opts = tmp; + + char identifier; + cag_option_context cag_context; + + // Prepare the context and iterate over all options. + cag_option_prepare(&cag_context, opts, n, argc, argv); + while (cag_option_fetch(&cag_context)) + { + ParseState state = NotParsed; + identifier = cag_option_get(&cag_context); + + PARS_OPTIONS(parse_DiagnosticOptions(identifier, diag_opts), state); + PARS_OPTIONS(parse_SaxpyOptions(identifier, &cag_context, saxpy_opts), + state); + + if (identifier == 'h') + { + printf("Usage: externalmemory [OPTION]...\n"); + printf("Option name and value should be separated by '=' or a " + "space\n"); + printf("Demonstrates OpenCL--Vulkan interop.\n\n"); + cag_option_print(opts, n, stdout); + exit((state == ParseError) ? CL_INVALID_ARG_VALUE : CL_SUCCESS); + } + } +end: + free(opts); + return error; +} + +// Host-side saxpy implementation. +void host_saxpy(const cl_float* x, cl_float* y, const float a, size_t length) +{ + for (size_t i = 0; i < length; ++i) + { + y[i] = fmaf(a, x[i], y[i]); + } +} + +// Vulkan instance extensions required for sharing OpenCL and Vulkan types: +// - VK_KHR_EXTERNAL_MEMORY_CAPABILITIES required for sharing buffers. +// - VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2 required for the previous one +// and for querying the device's UUID. +const char* const required_instance_extensions[] = { + VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME, /*VK_KHR_external_memory_capabilities*/ + VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_EXTENSION_NAME /*VK_KHR_get_physical_device_properties2*/ +}; +const size_t required_instance_extensions_count = + sizeof(required_instance_extensions) / sizeof(const char*); + +// General Vulkan extensions that a device needs to support for exporting +// memory. +const char* required_device_extensions[] = { + VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME, /*VK_KHR_external_memory*/ +#ifdef _WIN32 + VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME /*VK_KHR_external_memory_win32*/ +#else + VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME /*VK_KHR_external_memory_fd*/ +#endif +}; +const size_t required_device_extensions_count = + sizeof(required_device_extensions) / sizeof(const char*); + +// Khronos extensions that a device needs to support memory sharing with Vulkan. +const char* required_khronos_extensions[] = { +#ifdef _WIN32 + CL_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME /*cl_khr_external_memory_win32*/ +#else + CL_KHR_EXTERNAL_MEMORY_OPAQUE_FD_EXTENSION_NAME /*cl_khr_external_memory_opaque_fd*/ +#endif +}; +const size_t required_khronos_extensions_count = + sizeof(required_khronos_extensions) / sizeof(const char*); + +// Required Vulkan external memory handle. +const VkExternalMemoryHandleTypeFlagBits vk_external_memory_handle_type = +#ifdef _WIN32 + VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR; +#else + VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR; +#endif + +// Required OpenCL external memory handle. +const cl_external_memory_handle_type_khr cl_external_memory_handle_type = +#ifdef _WIN32 + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR; +#else + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR; +#endif + +// Check if a given OpenCL device supports a particular external memory handle +// type. +bool cl_check_external_memory_handle_type( + const cl_device_id cl_device, + cl_external_memory_handle_type_khr external_memory_handle_type) +{ + cl_external_memory_handle_type_khr* supported_handle_types = NULL; + size_t supported_handle_types_count = 0; + cl_int error = CL_SUCCESS; + + OCLERROR_RET( + clGetDeviceInfo(cl_device, + CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, 0, + NULL, &supported_handle_types_count), + error, err); + supported_handle_types = (cl_external_memory_handle_type_khr*)malloc( + supported_handle_types_count); + + OCLERROR_RET( + clGetDeviceInfo( + cl_device, CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, + supported_handle_types_count, supported_handle_types, NULL), + error, err); + for (size_t i = 0; i < supported_handle_types_count; ++i) + { + if (external_memory_handle_type == supported_handle_types[i]) + { + free(supported_handle_types); + return true; + } + } + free(supported_handle_types); + return false; +err: + fprintf(stderr, + "Error: OpenCL could not query supported external memory handle " + "types\n"); + free(supported_handle_types); + exit(EXIT_FAILURE); +} + +int main(int argc, char* argv[]) +{ + cl_int error = CL_SUCCESS; + cl_int end_error = CL_SUCCESS; + cl_platform_id cl_platform; + cl_device_id cl_device; + VkPhysicalDevice vk_physical_device; + VkDevice vk_device; + cl_context context = NULL; + cl_command_queue queue = NULL; + + cl_program program; + + // Parse command-line options. + struct cl_sdk_options_Diagnostic diag_opts = { .quiet = false, + .verbose = false }; + // Define as default length 1048576 = 4 * 262144 = sizeof(cl_float) * 2^18. + struct options_Saxpy saxpy_opts = { .length = 1048576 }; + + OCLERROR_RET(parse_options(argc, argv, &diag_opts, &saxpy_opts), error, + end); + + // Fill in Vulkan application info. + VkApplicationInfo app_info = { 0 }; + app_info.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO; + app_info.pApplicationName = "OpenCL-Vulkan interop example"; + app_info.applicationVersion = VK_MAKE_VERSION(3, 0, 0); + app_info.pEngineName = "OpenCL-SDK samples"; + app_info.engineVersion = VK_MAKE_VERSION(3, 0, 0); + app_info.apiVersion = VK_MAKE_VERSION(3, 0, 0); + + // Initialize Vulkan instance info and create Vulkan instance. + VkInstanceCreateInfo instance_create_info = { + VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO + }; + instance_create_info.pApplicationInfo = &app_info; + instance_create_info.enabledExtensionCount = + (uint32_t)required_instance_extensions_count; + instance_create_info.ppEnabledExtensionNames = required_instance_extensions; + + VkInstance instance; + VK_CHECK(vkCreateInstance(&instance_create_info, NULL, &instance)); + + // Find a suitable (Vulkan-compatible) OpenCL device for the sample. + struct device_candidate candidate = find_suitable_device( + instance, required_device_extensions, required_device_extensions_count); + + // OpenCL device object for the selected device. + cl_device = candidate.cl_candidate.device; + + // Vulkan physical device object for the selected device. + vk_physical_device = candidate.vk_candidate; + + // Set up necessary info and create Vulkan device from physical device. + const float default_queue_priority = 1.0f; + VkDeviceQueueCreateInfo queue_create_info = { + VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO + }; + queue_create_info.queueFamilyIndex = 0; + queue_create_info.queueCount = 1; + queue_create_info.pQueuePriorities = &default_queue_priority; + + VkDeviceCreateInfo device_create_info = { + VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO + }; + device_create_info.queueCreateInfoCount = 1; + device_create_info.pQueueCreateInfos = &queue_create_info; + device_create_info.enabledExtensionCount = + (uint32_t)required_device_extensions_count; + device_create_info.ppEnabledExtensionNames = required_device_extensions; + + VK_CHECK(vkCreateDevice(vk_physical_device, &device_create_info, NULL, + &vk_device)); + + if (!diag_opts.quiet) + { + cl_util_print_device_info(cl_device); + } + + // Create OpenCL runtime objects. + OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_PLATFORM, + sizeof(cl_platform_id), &cl_platform, NULL), + error, cont); + cl_context_properties context_props[] = { + CL_CONTEXT_PLATFORM, (cl_context_properties)cl_platform, 0 + }; + OCLERROR_PAR(context = clCreateContext(context_props, 1, &cl_device, NULL, + NULL, &error), + error, end); + + // Check if the device supports the Khronos extensions needed before + // attempting to compile the kernel. + if (diag_opts.verbose) + { + printf("\nChecking Khronos extensions support... "); + fflush(stdout); + } + + if (!check_khronos_extensions(cl_device, required_khronos_extensions, + required_khronos_extensions_count)) + { + fprintf(stdout, + "OpenCL device does not support the required Khronos " + "extensions\n"); + exit(EXIT_SUCCESS); + } + + // Compile kernel. + if (diag_opts.verbose) + { + printf("done.\nCompiling OpenCL kernel... "); + fflush(stdout); + } + const char* kernel_location = "./external_saxpy.cl"; + char *kernel = NULL, *tmp = NULL; + size_t program_size = 0; + OCLERROR_PAR( + kernel = cl_util_read_text_file(kernel_location, &program_size, &error), + error, que); + MEM_CHECK(tmp = (char*)realloc(kernel, program_size), error, ker); + kernel = tmp; + OCLERROR_PAR(program = clCreateProgramWithSource( + context, 1, (const char**)&kernel, &program_size, &error), + error, ker); + + // The Khronos extension showcased requires OpenCL 3.0 version. + char compiler_options[1023] = ""; +#if CL_HPP_TARGET_OPENCL_VERSION >= 300 + strcat(compiler_options, "-cl-std=CL3.0 "); +#else + fprintf(stderr, "\nError: OpenCL version must be at least 3.0\n"); + exit(EXIT_FAILURE); +#endif + + OCLERROR_RET(cl_util_build_program(program, cl_device, compiler_options), + error, prg); + + // Query maximum workgroup size (WGS) supported based on private mem + // (registers) constraints. + size_t wgs; + cl_kernel saxpy; + OCLERROR_PAR(saxpy = clCreateKernel(program, "saxpy", &error), error, prg); + OCLERROR_RET(clGetKernelWorkGroupInfo(saxpy, cl_device, + CL_KERNEL_WORK_GROUP_SIZE, + sizeof(size_t), &wgs, NULL), + error, ker); + + // Initialize host-side storage. + const size_t length = saxpy_opts.length; + + // Random number generator. + pcg32_random_t rng; + pcg32_srandom_r(&rng, 11111, 2222); + + // Initialize input and output vectors and constant. + cl_float *arr_x, *arr_y, a; + MEM_CHECK(arr_x = (cl_float*)malloc(sizeof(cl_float) * length), error, sxp); + MEM_CHECK(arr_y = (cl_float*)malloc(sizeof(cl_float) * length), error, + arrx); + if (diag_opts.verbose) + { + printf("done.\nGenerating random scalar and %zd random numbers for " + "saxpy input vector...", + length); + fflush(stdout); + } + cl_sdk_fill_with_random_floats_range(&rng, &a, 1, -100, 100); + cl_sdk_fill_with_random_floats_range(&rng, arr_x, length, -100, 100); + cl_sdk_fill_with_random_floats_range(&rng, arr_y, length, -100, 100); + + // Check if the device supports the required OpenCL handle type. + if (diag_opts.verbose) + { + printf( + "done.\nChecking OpenCL external memory handle type support... "); + fflush(stdout); + } + + if (!cl_check_external_memory_handle_type(cl_device, + cl_external_memory_handle_type)) + { + fprintf(stderr, + "\nError: Unsupported OpenCL external memory handle type\n"); + exit(EXIT_FAILURE); + } + + if (!vk_check_external_memory_handle_type(vk_physical_device, + vk_external_memory_handle_type)) + { + fprintf(stderr, + "\nError: Unsupported Vulkan external memory handle type\n"); + exit(EXIT_FAILURE); + } + + // Initialize Vulkan device-side storage. + if (diag_opts.verbose) + { + printf("done.\nInitializing Vulkan device storage... "); + fflush(stdout); + } + + // Create Vulkan (external) buffers and assign memory to them. + VkExternalMemoryBufferCreateInfo external_memory_buffer_info = { + VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO + }; + external_memory_buffer_info.handleTypes = vk_external_memory_handle_type; + + VkBufferCreateInfo buffer_info = { 0 }; + buffer_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; + buffer_info.pNext = &external_memory_buffer_info; + buffer_info.size = sizeof(cl_float) * length; + buffer_info.usage = + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT; + ; + buffer_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + + VkBuffer vk_buf_x, vk_buf_y; + VK_CHECK(vkCreateBuffer(vk_device, &buffer_info, NULL, &vk_buf_x)); + VK_CHECK(vkCreateBuffer(vk_device, &buffer_info, NULL, &vk_buf_y)); + + // Get requirements and necessary information for (exportable) memory. + VkMemoryRequirements mem_requirements_x = { 0 }, mem_requirements_y = { 0 }; + vkGetBufferMemoryRequirements(vk_device, vk_buf_x, &mem_requirements_x); + vkGetBufferMemoryRequirements(vk_device, vk_buf_y, &mem_requirements_y); + + VkExportMemoryAllocateInfo export_memory_alloc_info = { + VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO + }; + export_memory_alloc_info.handleTypes = vk_external_memory_handle_type; + + VkMemoryAllocateInfo memory_alloc_info_x = { 0 }; + memory_alloc_info_x.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + memory_alloc_info_x.pNext = &export_memory_alloc_info; + memory_alloc_info_x.allocationSize = mem_requirements_x.size; + memory_alloc_info_x.memoryTypeIndex = find_vk_memory_type( + vk_physical_device, mem_requirements_x.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT + | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + + VkMemoryAllocateInfo memory_alloc_info_y = { 0 }; + memory_alloc_info_y.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + memory_alloc_info_y.pNext = &export_memory_alloc_info; + memory_alloc_info_y.allocationSize = mem_requirements_y.size; + memory_alloc_info_y.memoryTypeIndex = find_vk_memory_type( + vk_physical_device, mem_requirements_y.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT + | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + + // Allocate and bind memory. + VkDeviceMemory vk_buf_x_memory, vk_buf_y_memory; + VK_CHECK(vkAllocateMemory(vk_device, &memory_alloc_info_x, NULL, + &vk_buf_x_memory)); + VK_CHECK(vkAllocateMemory(vk_device, &memory_alloc_info_y, NULL, + &vk_buf_y_memory)); + + VK_CHECK(vkBindBufferMemory(vk_device, vk_buf_x, vk_buf_x_memory, 0)); + VK_CHECK(vkBindBufferMemory(vk_device, vk_buf_y, vk_buf_y_memory, 0)); + + // Map memory. + void *vk_arr_x, *vk_arr_y; + VK_CHECK(vkMapMemory(vk_device, vk_buf_x_memory, 0, VK_WHOLE_SIZE, 0, + &vk_arr_x)); + VK_CHECK(vkMapMemory(vk_device, vk_buf_y_memory, 0, VK_WHOLE_SIZE, 0, + &vk_arr_y)); + + memcpy(vk_arr_x, arr_x, sizeof(cl_float) * length); + memcpy(vk_arr_y, arr_y, sizeof(cl_float) * length); + + // Get Vulkan external memory file descriptors for accessing external memory + // with OpenCL. + VkMemoryGetFdInfoKHR fd_info_x = { 0 }; + fd_info_x.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR; + fd_info_x.pNext = NULL; + fd_info_x.memory = vk_buf_x_memory; + fd_info_x.handleType = vk_external_memory_handle_type; + int fd_x; + + VkMemoryGetFdInfoKHR fd_info_y = { 0 }; + fd_info_y.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR; + fd_info_y.pNext = NULL; + fd_info_y.memory = vk_buf_y_memory; + fd_info_y.handleType = vk_external_memory_handle_type; + int fd_y; + + // We need to get the pointer to the vkGetMemoryFdKHR function because it's + // from extension VK_KHR_external_memory_fd. + PFN_vkGetMemoryFdKHR vkGetMemoryFdKHR = + (PFN_vkGetMemoryFdKHR)vkGetDeviceProcAddr(vk_device, + "vkGetMemoryFdKHR"); + + VK_CHECK(vkGetMemoryFdKHR(vk_device, &fd_info_x, &fd_x)); + VK_CHECK(vkGetMemoryFdKHR(vk_device, &fd_info_y, &fd_y)); + + // Create OpenCL buffers from Vulkan external memory file descriptors. + cl_mem_properties ext_mem_props_x[] = { + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, + (cl_mem_properties)fd_x, + (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)(uintptr_t)cl_device, + CL_DEVICE_HANDLE_LIST_END_KHR, + 0 + }; + cl_mem_properties ext_mem_props_y[] = { + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, + (cl_mem_properties)fd_y, + (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)(uintptr_t)cl_device, + CL_DEVICE_HANDLE_LIST_END_KHR, + 0 + }; + cl_mem cl_buf_x, cl_buf_y; + OCLERROR_PAR(cl_buf_x = clCreateBufferWithProperties( + context, ext_mem_props_x, CL_MEM_READ_ONLY, + sizeof(cl_float) * length, NULL, &error), + error, vulkan); + OCLERROR_PAR(cl_buf_y = clCreateBufferWithProperties( + context, ext_mem_props_y, CL_MEM_READ_WRITE, + sizeof(cl_float) * length, NULL, &error), + error, clbufx); + + // Initialize queue for command execution. + cl_command_queue_properties queue_props[] = { CL_QUEUE_PROPERTIES, + CL_QUEUE_PROFILING_ENABLE, + 0 }; + OCLERROR_PAR(queue = clCreateCommandQueueWithProperties( + context, cl_device, queue_props, &error), + error, cont); + + // Set kernel arguments. + OCLERROR_RET(clSetKernelArg(saxpy, 0, sizeof(cl_float), &a), error, clbufy); + OCLERROR_RET(clSetKernelArg(saxpy, 1, sizeof(cl_mem), &cl_buf_x), error, + clbufy); + OCLERROR_RET(clSetKernelArg(saxpy, 2, sizeof(cl_mem), &cl_buf_y), error, + clbufy); + + // Launch kernel. + if (diag_opts.verbose) + { + printf("done.\nExecuting on device... "); + fflush(stdout); + } + + cl_event kernel_run; + GET_CURRENT_TIMER(dev_start) + OCLERROR_RET(clEnqueueNDRangeKernel(queue, saxpy, 1, NULL, &length, &wgs, 0, + NULL, &kernel_run), + error, clbufy); + OCLERROR_RET(clWaitForEvents(1, &kernel_run), error, clbufy); + GET_CURRENT_TIMER(dev_end) + + cl_ulong dev_time; + TIMER_DIFFERENCE(dev_time, dev_start, dev_end) + + // Concurrently calculate reference saxpy. + if (diag_opts.verbose) + { + printf("done.\nExecuting on host... "); + } + + GET_CURRENT_TIMER(host_start) + host_saxpy(arr_x, arr_y, a, length); + GET_CURRENT_TIMER(host_end) + cl_ulong host_time; + TIMER_DIFFERENCE(host_time, host_start, host_end) + + if (diag_opts.verbose) + { + printf("done.\n"); + } + + // Fetch results. + OCLERROR_RET(clEnqueueReadBuffer(queue, cl_buf_y, CL_BLOCKING, 0, + sizeof(cl_float) * length, (void*)arr_x, 0, + NULL, NULL), + error, clbufy); + + // Validate solution. + for (size_t i = 0; i < length; ++i) + if (arr_y[i] != arr_x[i]) + { + printf("Verification failed! %f != %f at index %zu\n", arr_y[i], + arr_x[i], i); + error = CL_INVALID_VALUE; + } + if (error == CL_SUCCESS) + { + printf("Verification passed.\n"); + } + + if (!diag_opts.quiet) + { + printf("Kernel execution time as seen by host: %llu us.\n", + (unsigned long long)(dev_time + 500) / 1000); + + printf("Kernel execution time as measured by device:\n"); + printf("\t%llu us.\n", + (unsigned long long)(cl_util_get_event_duration( + kernel_run, CL_PROFILING_COMMAND_START, + CL_PROFILING_COMMAND_END, &error) + + 500) + / 1000); + + printf("Reference execution as seen by host: %llu us.\n", + (unsigned long long)(host_time + 500) / 1000); + } + + // Release resources. +clbufy: + OCLERROR_RET(clReleaseMemObject(cl_buf_y), end_error, clbufx); +clbufx: + OCLERROR_RET(clReleaseMemObject(cl_buf_x), end_error, vulkan); +vulkan: + vkDestroyBuffer(vk_device, vk_buf_y, NULL); + vkDestroyBuffer(vk_device, vk_buf_x, NULL); + vkUnmapMemory(vk_device, vk_buf_y_memory); + vkUnmapMemory(vk_device, vk_buf_x_memory); + vkFreeMemory(vk_device, vk_buf_y_memory, NULL); + vkFreeMemory(vk_device, vk_buf_x_memory, NULL); + free(arr_y); +arrx: + free(arr_x); +sxp: + OCLERROR_RET(clReleaseKernel(saxpy), end_error, prg); +prg: + OCLERROR_RET(clReleaseProgram(program), end_error, ker); +ker: + free(kernel); +que: + OCLERROR_RET(clReleaseCommandQueue(queue), end_error, cont); +cont: + OCLERROR_RET(clReleaseContext(context), end_error, end); +end: + if (error) cl_util_print_error(error); + return error; +} diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp new file mode 100644 index 00000000..e4026c23 --- /dev/null +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -0,0 +1,559 @@ +/* + * Copyright (c) 2023 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. + */ + +// OpenCL C++ headers includes. +#include + +// OpenCL SDK includes. +#include +#include +#include +#include + +// OpenCL Utils includes. +#include +#include +#include + +// Vulkan includes. +#include + +// Vulkan utils includes. +#include "vulkan_utils.hpp" + +// Standard header includes. +#include +#include +#include +#include +#include + +// Sample-specific option. +struct SaxpyOptions +{ + size_t length; +}; + +// Add option to CLI-parsing SDK utility for input dimensions. +template <> auto cl::sdk::parse() +{ + return std::make_tuple(std::make_shared>( + "l", "length", "Length of input", false, 1'048'576, + "positive integral")); +} +template <> +SaxpyOptions cl::sdk::comprehend( + std::shared_ptr> length_arg) +{ + return SaxpyOptions{ length_arg->getValue() }; +} + +// Host-side saxpy implementation. +void host_saxpy(std::vector x, std::vector& y, const float a, + size_t length) +{ + for (size_t i = 0; i < length; ++i) + { + y[i] = std::fmaf(a, x[i], y[i]); + } +} + +// Vulkan instance extensions required for sharing OpenCL and Vulkan types: +// - VK_KHR_EXTERNAL_MEMORY_CAPABILITIES required for sharing buffers. +// - VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2 required for the previous one +// and for querying the device's UUID. +const std::vector required_instance_extensions_str = { + VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME, /*VK_KHR_external_memory_capabilities*/ + VK_KHR_GET_PHYSICAL_DEVICE_PROPERTIES_2_EXTENSION_NAME /*VK_KHR_get_physical_device_properties2*/ +}; + +// General Vulkan extensions that a device needs to support to run this +// example: +// - VK_KHR_EXTERNAL_MEMORY required for sharing memory. +const std::vector required_device_extensions_str = { + std::string{ + VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME }, /*VK_KHR_external_memory*/ +#ifdef _WIN64 + std::string{ + VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME } /*VK_KHR_external_memory_win32*/ +#else + std::string{ + VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME } /*VK_KHR_external_memory_fd*/ +#endif +}; + +// Required Vulkan external memory handle. +const VkExternalMemoryHandleTypeFlagBits vk_external_memory_handle_type = +#ifdef _WIN32 + VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR; +#else + VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT_KHR; +#endif + +// Khronos extensions that a device needs to support memory sharing with Vulkan. +const std::vector required_khronos_extensions = { +#ifdef _WIN32 + std::string{ "cl_khr_external_memory_win32" } +#else + std::string{ "cl_khr_external_memory_opaque_fd" } +#endif +}; + +// Required OpenCL external memory handle. +const cl_external_memory_handle_type_khr cl_external_memory_handle_type = +#ifdef _WIN32 + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR; +#else + CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR; +#endif + +// Check if a given OpenCL device supports a particular external memory handle +// type. +bool cl_check_external_memory_handle_type( + const cl::Device cl_device, + cl_external_memory_handle_type_khr external_memory_handle_type) +{ + std::vector supported_handle_types; + cl_device.getInfo(CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, + &supported_handle_types); + + const auto it = std::find_if( + supported_handle_types.begin(), supported_handle_types.end(), + [&](const cl_external_memory_handle_type_khr& supported_handle_type) { + return external_memory_handle_type == supported_handle_type; + }); + return it != supported_handle_types.end(); +} + +int main(int argc, char* argv[]) +{ + try + { + // Parse command-line options. + auto opts = + cl::sdk::parse_cli( + argc, argv); + const auto& diag_opts = std::get<0>(opts); + const auto& saxpy_opts = std::get<1>(opts); + + // Fill in Vulkan application info. + VkApplicationInfo app_info{}; + app_info.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO; + app_info.pApplicationName = "OpenCL-Vulkan interop example"; + app_info.applicationVersion = VK_MAKE_VERSION(3, 0, 0); + app_info.pEngineName = "OpenCL-SDK samples"; + app_info.engineVersion = VK_MAKE_VERSION(3, 0, 0); + app_info.apiVersion = VK_MAKE_VERSION(3, 0, 0); + + // Initialize Vulkan instance info and create Vulkan instance. + std::vector required_instance_extensions( + required_instance_extensions_str.size(), nullptr); + std::transform(required_instance_extensions_str.begin(), + required_instance_extensions_str.end(), + required_instance_extensions.begin(), + [&](const std::string& str) { return str.c_str(); }); + VkInstanceCreateInfo instance_create_info{}; + instance_create_info.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; + instance_create_info.pApplicationInfo = &app_info; + instance_create_info.enabledExtensionCount = + static_cast(required_instance_extensions.size()); + instance_create_info.ppEnabledExtensionNames = + required_instance_extensions.data(); + + VkInstance instance; + VK_CHECK(vkCreateInstance(&instance_create_info, nullptr, &instance)); + + // Find a suitable (Vulkan-compatible) OpenCL device for the sample. + std::vector required_device_extensions( + required_device_extensions_str.size(), nullptr); + std::transform(required_device_extensions_str.begin(), + required_device_extensions_str.end(), + required_device_extensions.begin(), + [&](const std::string& str) { return str.c_str(); }); + device_candidate candidate = + find_suitable_device(instance, required_device_extensions); + + // OpenCL device and platform objects for the selected device. + cl::Device cl_device = candidate.cl_candidate.device; + const cl::Platform cl_platform{ + cl_device.getInfo() + }; + + // Vulkan physical device object for the selected device. + const VkPhysicalDevice vk_physical_device = candidate.vk_candidate; + + // Set up necessary info and create Vulkan device from physical device. + constexpr float default_queue_priority = 1.0f; + VkDeviceQueueCreateInfo queue_create_info{}; + queue_create_info.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO; + queue_create_info.queueFamilyIndex = 0; + queue_create_info.queueCount = 1; + queue_create_info.pQueuePriorities = &default_queue_priority; + + VkDeviceCreateInfo device_create_info{}; + device_create_info.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO; + device_create_info.queueCreateInfoCount = 1; + device_create_info.pQueueCreateInfos = &queue_create_info; + device_create_info.enabledExtensionCount = + static_cast(required_device_extensions.size()); + device_create_info.ppEnabledExtensionNames = + required_device_extensions.data(); + + VkDevice vk_device; + VK_CHECK(vkCreateDevice(vk_physical_device, &device_create_info, + nullptr, &vk_device)); + + if (!diag_opts.quiet) + { + std::cout << "Selected platform: " + << cl_platform.getInfo() << "\n" + << "Selected device: " + << cl_device.getInfo() << "\n" + << std::endl; + } + + // Create OpenCL runtime objects. + cl::Context cl_context{ cl_device }; + + // Check if the device supports the Khronos extensions needed before + // attempting to compile the kernel. + if (diag_opts.verbose) + { + std::cout << "Checking Khronos extensions support... "; + std::cout.flush(); + } + + for (const auto& extension : required_khronos_extensions) + { + if (!cl::util::supports_extension(cl_device, extension)) + { + std::cout << "OpenCL device does not support the required " + "Khronos extension " + << extension << std::endl; + exit(EXIT_SUCCESS); + } + } + + // Compile kernel. + if (diag_opts.verbose) + { + std::cout << " done.\nCompiling OpenCL kernel... "; + std::cout.flush(); + } + const char* kernel_location = "./external_saxpy.cl"; + std::ifstream kernel_stream{ kernel_location }; + if (!kernel_stream.is_open()) + throw std::runtime_error{ + std::string{ "Cannot open kernel source: " } + kernel_location + }; + cl::Program cl_program{ + cl_context, + std::string{ std::istreambuf_iterator{ kernel_stream }, + std::istreambuf_iterator{} } + }; + + // The Khronos extension showcased requires OpenCL 3.0 version. + cl::string compiler_options = ""; +#if CL_HPP_TARGET_OPENCL_VERSION >= 300 + compiler_options += cl::string{ "-cl-std=CL3.0 " }; +#else + sdt::cerr << "\nError: OpenCL version must be at least 3.0" + << std::endl; + exit(EXIT_FAILURE); +#endif + + cl_program.build(cl_device, compiler_options.c_str()); + + // Query maximum workgroup size (WGS) supported based on private mem + // (registers) constraints. + auto saxpy = cl::KernelFunctor( + cl_program, "saxpy"); + auto wgs = + saxpy.getKernel().getWorkGroupInfo( + cl_device); + + // Initialize host-side storage. + const auto length = saxpy_opts.length; + + // Random number generator. + auto prng = [engine = std::default_random_engine{}, + dist = std::uniform_real_distribution{ + -1.0, 1.0 }]() mutable { return dist(engine); }; + + // Initialize input and output vectors and constant. + std::vector arr_x(length), arr_y(length); + if (diag_opts.verbose) + { + std::cout << "Generating random scalar and " << length + << " random numbers for saxpy input vector." << std::endl; + } + cl_float a = prng(); + cl::sdk::fill_with_random(prng, arr_x, arr_y); + + // Check if the device supports the required OpenCL handle type. + if (diag_opts.verbose) + { + std::cout << "done.\nChecking OpenCL external memory handle type " + "support... "; + std::cout.flush(); + } + + if (!cl_check_external_memory_handle_type( + cl_device, cl_external_memory_handle_type)) + { + std::cerr + << "\nError: Unsupported OpenCL external memory handle type" + << std::endl; + exit(EXIT_FAILURE); + } + + if (!vk_check_external_memory_handle_type( + vk_physical_device, vk_external_memory_handle_type)) + { + std::cerr + << "\nError: Unsupported Vulkan external memory handle type" + << std::endl; + exit(EXIT_FAILURE); + } + + // Initialize Vulkan device-side storage. + if (diag_opts.verbose) + { + std::cout << "done.\nInitializing Vulkan device storage... "; + std::cout.flush(); + } + + // Create Vulkan (external) buffers and assign memory to them. + VkExternalMemoryBufferCreateInfo external_memory_buffer_info{}; + external_memory_buffer_info.sType = + VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO; + external_memory_buffer_info.handleTypes = + vk_external_memory_handle_type; + + VkBufferCreateInfo buffer_info{}; + buffer_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; + buffer_info.pNext = &external_memory_buffer_info; + buffer_info.size = sizeof(cl_float) * length; + buffer_info.usage = + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT; + ; + buffer_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + + VkBuffer vk_buf_x, vk_buf_y; + VK_CHECK(vkCreateBuffer(vk_device, &buffer_info, nullptr, &vk_buf_x)); + VK_CHECK(vkCreateBuffer(vk_device, &buffer_info, nullptr, &vk_buf_y)); + + // Get requirements and necessary information for (exportable) memory. + VkMemoryRequirements mem_requirements_x{}, mem_requirements_y{}; + vkGetBufferMemoryRequirements(vk_device, vk_buf_x, &mem_requirements_x); + vkGetBufferMemoryRequirements(vk_device, vk_buf_y, &mem_requirements_y); + + VkExportMemoryAllocateInfo export_memory_alloc_info{}; + export_memory_alloc_info.sType = + VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO; + export_memory_alloc_info.handleTypes = vk_external_memory_handle_type; + + VkMemoryAllocateInfo memory_alloc_info_x{}; + memory_alloc_info_x.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + memory_alloc_info_x.pNext = &export_memory_alloc_info; + memory_alloc_info_x.allocationSize = mem_requirements_x.size; + memory_alloc_info_x.memoryTypeIndex = find_vk_memory_type( + vk_physical_device, mem_requirements_x.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT + | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + + VkMemoryAllocateInfo memory_alloc_info_y{}; + memory_alloc_info_y.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + memory_alloc_info_y.pNext = &export_memory_alloc_info; + memory_alloc_info_y.allocationSize = mem_requirements_y.size; + memory_alloc_info_y.memoryTypeIndex = find_vk_memory_type( + vk_physical_device, mem_requirements_y.memoryTypeBits, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT + | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT); + + // Allocate and bind memory. + VkDeviceMemory vk_buf_x_memory, vk_buf_y_memory; + VK_CHECK(vkAllocateMemory(vk_device, &memory_alloc_info_x, nullptr, + &vk_buf_x_memory)); + VK_CHECK(vkAllocateMemory(vk_device, &memory_alloc_info_y, nullptr, + &vk_buf_y_memory)); + + VK_CHECK(vkBindBufferMemory(vk_device, vk_buf_x, vk_buf_x_memory, 0)); + VK_CHECK(vkBindBufferMemory(vk_device, vk_buf_y, vk_buf_y_memory, 0)); + + // Map memory. + void *vk_arr_x, *vk_arr_y; + VK_CHECK(vkMapMemory(vk_device, vk_buf_x_memory, 0, VK_WHOLE_SIZE, 0, + &vk_arr_x)); + VK_CHECK(vkMapMemory(vk_device, vk_buf_y_memory, 0, VK_WHOLE_SIZE, 0, + &vk_arr_y)); + + memcpy(vk_arr_x, arr_x.data(), sizeof(cl_float) * length); + memcpy(vk_arr_y, arr_y.data(), sizeof(cl_float) * length); + + // Get Vulkan external memory file descriptors for accessing external + // memory with OpenCL. + VkMemoryGetFdInfoKHR fd_info_x{}; + fd_info_x.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR; + fd_info_x.pNext = nullptr; + fd_info_x.memory = vk_buf_x_memory; + fd_info_x.handleType = vk_external_memory_handle_type; + int fd_x; + + VkMemoryGetFdInfoKHR fd_info_y{}; + fd_info_y.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR; + fd_info_y.pNext = nullptr; + fd_info_y.memory = vk_buf_y_memory; + fd_info_y.handleType = vk_external_memory_handle_type; + int fd_y; + + // We need to get the pointer to the vkGetMemoryFdKHR function because + // it's from extension VK_KHR_external_memory_fd. + PFN_vkGetMemoryFdKHR vkGetMemoryFdKHR = + (PFN_vkGetMemoryFdKHR)vkGetDeviceProcAddr(vk_device, + "vkGetMemoryFdKHR"); + + VK_CHECK(vkGetMemoryFdKHR(vk_device, &fd_info_x, &fd_x)); + VK_CHECK(vkGetMemoryFdKHR(vk_device, &fd_info_y, &fd_y)); + + // Create OpenCL buffers from Vulkan external memory file descriptors. + std::vector ext_mem_props_x = { + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, + (cl_mem_properties)fd_x, + (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)cl_device(), + CL_DEVICE_HANDLE_LIST_END_KHR, + 0 + }; + std::vector ext_mem_props_y = { + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, + (cl_mem_properties)fd_y, + (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)cl_device(), + CL_DEVICE_HANDLE_LIST_END_KHR, + 0 + }; + + cl::Buffer cl_buf_x{ cl_context, ext_mem_props_x, CL_MEM_READ_ONLY, + sizeof(cl_float) * length }; + cl::Buffer cl_buf_y{ cl_context, ext_mem_props_y, CL_MEM_READ_WRITE, + sizeof(cl_float) * length }; + + // Initialize queue for command execution. + cl_command_queue_properties queue_props[] = { CL_QUEUE_PROFILING_ENABLE, + 0 }; + cl::CommandQueue queue{ cl_context, cl_device, *queue_props }; + + // Launch kernel. + if (diag_opts.verbose) + { + std::cout << "done.\nExecuting on device... "; + std::cout.flush(); + } + + std::vector kernel_run; + auto dev_start = std::chrono::high_resolution_clock::now(); + kernel_run.push_back( + saxpy(cl::EnqueueArgs{ queue, cl::NDRange{ length }, wgs }, a, + cl_buf_x, cl_buf_y)); + cl::WaitForEvents(kernel_run); + auto dev_end = std::chrono::high_resolution_clock::now(); + + // Concurrently calculate reference saxpy. + if (diag_opts.verbose) + { + std::cout << "done.\nExecuting on host... "; + std::cout.flush(); + } + + auto host_start = std::chrono::high_resolution_clock::now(); + host_saxpy(arr_x, arr_y, a, length); + auto host_end = std::chrono::high_resolution_clock::now(); + + if (diag_opts.verbose) + { + std::cout << "done.\n"; + std::cout.flush(); + } + + // Fetch results. + cl::copy(queue, cl_buf_y, arr_x.begin(), arr_x.end()); + + // Validate solution. + if (std::equal(std::begin(arr_x), std::end(arr_x), std::begin(arr_y), + std::end(arr_y))) + std::cout << "Verification passed." << std::endl; + else + throw std::runtime_error{ "Verification failed!" }; + + if (!diag_opts.quiet) + { + std::cout << "Kernel execution time as seen by host: " + << std::chrono::duration_cast( + dev_end - dev_start) + .count() + << " us." << std::endl; + + std::cout << "Kernel execution time as measured by device: "; + std::cout << cl::util::get_duration( + kernel_run[0]) + .count() + << " us." << std::endl; + + std::cout << "Reference execution as seen by host: " + << std::chrono::duration_cast( + host_end - host_start) + .count() + << " us." << std::endl; + } + + // Release resources. + vkDestroyBuffer(vk_device, vk_buf_y, nullptr); + vkDestroyBuffer(vk_device, vk_buf_x, nullptr); + vkUnmapMemory(vk_device, vk_buf_y_memory); + vkUnmapMemory(vk_device, vk_buf_x_memory); + vkFreeMemory(vk_device, vk_buf_y_memory, nullptr); + vkFreeMemory(vk_device, vk_buf_x_memory, nullptr); + + } catch (cl::BuildError& e) + { + std::cerr << "OpenCL build error: " << e.what() << std::endl; + for (auto& build_log : e.getBuildLog()) + { + std::cerr << "\tBuild log for device: " + << build_log.first.getInfo() << "\n" + << std::endl; + std::cerr << build_log.second << "\n" << std::endl; + } + std::exit(e.err()); + } catch (cl::util::Error& e) + { + std::cerr << "OpenCL utils error: " << e.what() << std::endl; + std::exit(e.err()); + } catch (cl::Error& e) + { + std::cerr << "OpenCL runtime error: " << e.what() << std::endl; + std::exit(e.err()); + } catch (std::exception& e) + { + std::cerr << "Error: " << e.what() << std::endl; + std::exit(EXIT_FAILURE); + } + + return 0; +} diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.h b/samples/extensions/khr/externalmemory/vulkan_utils.h new file mode 100644 index 00000000..2b718564 --- /dev/null +++ b/samples/extensions/khr/externalmemory/vulkan_utils.h @@ -0,0 +1,382 @@ +/* + * Copyright (c) 2023 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. + */ + +#ifndef _SAMPLES_CORE_EXTERNALMEMORY_UTILS_H +#define _SAMPLES_CORE_EXTERNALMEMORY_UTILS_H + +// OpenCL C headers includes. +#include + +// OpenCL Utils includes. +#include + +// Vulkan includes. +#include + +// Standard header includes. +#include +#include +#include +#include + +// Check if the provided Vulkan error code is \p VK_SUCCESS. If not, prints an +// error message to the standard error output and terminates the program with an +// error code. +#define VK_CHECK(condition) \ + { \ + const VkResult _error = condition; \ + if (_error != VK_SUCCESS) \ + { \ + fprintf(stderr, "A vulkan error encountered: %d at %s: %d\n", \ + _error, __FILE__, __LINE__); \ + exit(EXIT_FAILURE); \ + } \ + } + +// OpenCL device that is suitable for this example. +struct cl_device_candidate +{ + /// The OpenCL device id representing the device. + cl_device_id device; + + /// The Vulkan-compatible device UUID. + cl_uchar uuid[CL_UUID_SIZE_KHR]; +}; + +// OpenCL and Vulkan physical device suitable for the sample. +struct device_candidate +{ + /// The Vulkan physical device handle of the device to be used. + VkPhysicalDevice vk_candidate; + + /// The candidate device's Vulkan device properties. + VkPhysicalDeviceProperties vk_props; + + /// The OpenCL device candidate that this Vulkan device corresponds to. + struct cl_device_candidate cl_candidate; +}; + +// Check if the extensions supported by a Vulkan device includes a given set of +// required extensions. +bool extensions_supported( + const VkExtensionProperties* supported_extensions_properties, + const size_t supported_extensions_count, + const char* const* required_device_extensions, + const size_t required_device_extensions_count) +{ + for (size_t i = 0; i < required_device_extensions_count; ++i) + { + size_t j = 0; + while (strcmp(required_device_extensions[i], + supported_extensions_properties[j].extensionName) + && j < supported_extensions_count) + { + ++j; + } + if (j == supported_extensions_count) + { + return false; + } + } + return true; +} + +// Check if a given Vulkan device supports all the required Vulkan extensions. +bool check_device_extensions(const VkPhysicalDevice vk_device, + const char* const* required_device_extensions, + const size_t required_device_extensions_count) +{ + uint32_t supported_extensions_count; + VK_CHECK(vkEnumerateDeviceExtensionProperties( + vk_device, NULL, &supported_extensions_count, NULL)); + VkExtensionProperties* vk_supported_extensions_properties = + (VkExtensionProperties*)malloc(supported_extensions_count + * sizeof(VkExtensionProperties)); + VK_CHECK(vkEnumerateDeviceExtensionProperties( + vk_device, NULL, &supported_extensions_count, + vk_supported_extensions_properties)); + + bool result = extensions_supported( + vk_supported_extensions_properties, supported_extensions_count, + required_device_extensions, required_device_extensions_count); + + free(vk_supported_extensions_properties); + + return result; +} + +// Check if a given Vulkan physical device is compatible with any of the OpenCL +// devices available. +bool is_vk_device_suitable(const struct cl_device_candidate* cl_candidates, + const size_t cl_candidates_count, + VkPhysicalDevice vk_device, + struct device_candidate* candidate, + const char* const* required_device_extensions, + const size_t required_device_extensions_count) +{ + // Check if the device supports OpenCL by checking if there is any device + // with the same UUID. + { + // Query the Vulkan device UUID using vkGetPhysicalDeviceProperties2. + VkPhysicalDeviceIDPropertiesKHR id_props = { 0 }; + id_props.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES_KHR; + + VkPhysicalDeviceProperties2KHR props2 = { 0 }; + props2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2_KHR; + props2.pNext = &id_props; + + vkGetPhysicalDeviceProperties2(vk_device, &props2); + + // Look for an OpenCL device which UUID matches the UUID reported by + // Vulkan. + size_t compatible_dev_index = cl_candidates_count; + for (size_t i = 0; i < cl_candidates_count + && compatible_dev_index == cl_candidates_count; + ++i) + { + compatible_dev_index = i; + for (uint32_t j = 0; j < CL_UUID_SIZE_KHR; ++j) + { + if (cl_candidates[i].uuid[j] != id_props.deviceUUID[j]) + { + compatible_dev_index = cl_candidates_count; + break; + } + } + } + + if (compatible_dev_index == cl_candidates_count) + { + return false; + } + + candidate->vk_props = props2.properties; + candidate->cl_candidate = cl_candidates[compatible_dev_index]; + } + + // Check if the device supports the required extensions. + if (!check_device_extensions(vk_device, required_device_extensions, + required_device_extensions_count)) + { + return false; + } + + candidate->vk_candidate = vk_device; + return true; +} + +// Check if a given OpenCL device supports a particular set of Khronos +// extensions. +bool check_khronos_extensions( + const cl_device_id cl_device, + const char* const* const required_khronos_extensions, + const size_t required_khronos_extensions_count) +{ + cl_int error = CL_SUCCESS; + size_t supported_extensions_count; + OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_EXTENSIONS, 0, NULL, + &supported_extensions_count), + error, ret); + char* supported_extensions = + (char*)malloc(supported_extensions_count * sizeof(char)); + OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_EXTENSIONS, + supported_extensions_count, + supported_extensions, NULL), + error, err); + + for (size_t i = 0; i < required_khronos_extensions_count; ++i) + { + if (!strstr(supported_extensions, required_khronos_extensions[i])) + { + free(supported_extensions); + return false; + } + } + free(supported_extensions); + return true; +err: + free(supported_extensions); +ret: + return false; +} + +// Find a suitable device for the example, that is, an OpenCL +// device that is also Vulkan-compatible and that supports the required +// Vulkan device extensions. +struct device_candidate +find_suitable_device(VkInstance instance, + const char* const* required_device_extensions, + const size_t required_device_extensions_count) +{ + // Query OpenCL devices available. + cl_int error = CL_SUCCESS; + bool candidate_found = false; + cl_uint cl_platform_count = 0; + struct device_candidate found_candidate = {0}; + OCLERROR_RET(clGetPlatformIDs(0, NULL, &cl_platform_count), error, ret); + + cl_platform_id* platforms = + (cl_platform_id*)malloc(cl_platform_count * sizeof(cl_platform_id)); + OCLERROR_RET(clGetPlatformIDs(cl_platform_count, platforms, NULL), error, + platforms); + + size_t cl_device_count = 0; + const char* uuid_khronos_extension[] = { + CL_KHR_DEVICE_UUID_EXTENSION_NAME + }; + for (cl_uint cl_platform_id = 0; cl_platform_id < cl_platform_count; + ++cl_platform_id) + { + cl_uint cl_platform_devices_count = 0; + OCLERROR_RET(clGetDeviceIDs(platforms[cl_platform_id], + CL_DEVICE_TYPE_ALL, 0, NULL, + &cl_platform_devices_count), + error, platforms); + for (cl_uint device_id = 0; device_id < cl_platform_devices_count; + ++device_id) + { + cl_device_id device; + OCLERROR_PAR(device = cl_util_get_device( + cl_platform_id, device_id, CL_DEVICE_TYPE_ALL, &error), error, platforms); + cl_device_count += + check_khronos_extensions(device, uuid_khronos_extension, 1); + } + } + + // For each OpenCL device, query its Vulkan-compatible device UUID and + // add it to the list of candidates. The device must support the + // cl_khr_device_uuid extension for us to be able to query the device's + // UUID. + struct cl_device_candidate* cl_candidates = + (struct cl_device_candidate*)malloc( + cl_device_count * sizeof(struct cl_device_candidate)); + cl_device_count = 0; + for (cl_uint cl_platform_id = 0; cl_platform_id < cl_platform_count; + ++cl_platform_id) + { + cl_uint cl_platform_devices_count = 0; + OCLERROR_RET(clGetDeviceIDs(platforms[cl_platform_id], + CL_DEVICE_TYPE_ALL, 0, NULL, + &cl_platform_devices_count), + error, candidates); + + for (cl_uint cl_candidate_id = 0; + cl_candidate_id < cl_platform_devices_count; + ++cl_candidate_id, ++cl_device_count) + { + cl_device_id device = cl_util_get_device( + cl_platform_id, cl_candidate_id, CL_DEVICE_TYPE_ALL, &error); + if (check_khronos_extensions(device, uuid_khronos_extension, 1)) + { + cl_uchar vk_candidate_uuid[CL_UUID_SIZE_KHR]; + OCLERROR_RET(clGetDeviceInfo(device, CL_DEVICE_UUID_KHR, + CL_UUID_SIZE_KHR, + &vk_candidate_uuid, NULL), + error, candidates); + + struct cl_device_candidate candidate; + candidate.device = device; + memcpy(candidate.uuid, &vk_candidate_uuid, + sizeof(cl_uchar) * CL_UUID_SIZE_KHR); + cl_candidates[cl_device_count] = candidate; + } + } + } + + // Query the Vulkan physical devices available. + uint32_t vk_device_count; + VK_CHECK(vkEnumeratePhysicalDevices(instance, &vk_device_count, NULL)); + + VkPhysicalDevice* vk_devices = + (VkPhysicalDevice*)malloc(vk_device_count * sizeof(VkPhysicalDevice)); + VK_CHECK( + vkEnumeratePhysicalDevices(instance, &vk_device_count, vk_devices)); + + // Find a suitable Vulkan physical device compatible with one of the OpenCL + // devices available. + for (cl_uint vk_device_id = 0; vk_device_id < vk_device_count; + ++vk_device_id) + { + VkPhysicalDevice vk_device = vk_devices[vk_device_id]; + if (is_vk_device_suitable(cl_candidates, cl_device_count, vk_device, + &found_candidate, required_device_extensions, + required_device_extensions_count)) + { + candidate_found = true; + break; + } + } + if (!candidate_found) + { + printf("No suitable OpenCL Vulkan-compatible devices available\n"); + } + + free(vk_devices); +candidates: + free(cl_candidates); +platforms: + free(platforms); +ret: + if (candidate_found) + { + return found_candidate; + } + exit(error); +} + + +// Check if a given Vulkan device supports a particular external memory handle +// type. +bool vk_check_external_memory_handle_type( + VkPhysicalDevice vk_physical_device, + VkExternalMemoryHandleTypeFlagBits vk_external_memory_handle_type) +{ + VkPhysicalDeviceExternalBufferInfo physical_device_external_buffer_info = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTERNAL_BUFFER_INFO + }; + physical_device_external_buffer_info.handleType = + vk_external_memory_handle_type; + + VkExternalBufferProperties external_buffer_properties; + + vkGetPhysicalDeviceExternalBufferProperties( + vk_physical_device, &physical_device_external_buffer_info, + &external_buffer_properties); + + return (vk_external_memory_handle_type + & external_buffer_properties.externalMemoryProperties + .compatibleHandleTypes); +} + +// Find Vulkan memory properties from Vulkan physical device property flags. +uint32_t find_vk_memory_type(VkPhysicalDevice vk_device, uint32_t type_filter, + VkMemoryPropertyFlags properties) +{ + VkPhysicalDeviceMemoryProperties mem_properties; + vkGetPhysicalDeviceMemoryProperties(vk_device, &mem_properties); + for (uint32_t i = 0; i < mem_properties.memoryTypeCount; i++) + { + if ((type_filter & (1 << i)) + && (mem_properties.memoryTypes[i].propertyFlags & properties) + == properties) + { + return i; + } + } + return 0; +} + +#endif // _SAMPLES_CORE_EXTERNALMEMORY_UTILS_H diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.hpp b/samples/extensions/khr/externalmemory/vulkan_utils.hpp new file mode 100644 index 00000000..90090dfd --- /dev/null +++ b/samples/extensions/khr/externalmemory/vulkan_utils.hpp @@ -0,0 +1,268 @@ +/* + * Copyright (c) 2023 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. + */ + +#ifndef _SAMPLES_CORE_EXTERNALMEMORY_UTILS_HPP +#define _SAMPLES_CORE_EXTERNALMEMORY_UTILS_HPP + +// OpenCL C++ headers includes. +#include + +// OpenCL Utils includes. +#include + +// Vulkan includes. +#include + +// Standard header includes. +#include +#include +#include +#include + +// Check if the provided Vulkan error code is \p VK_SUCCESS. If not, prints an +// error message to the standard error output and terminates the program with an +// error code. +#define VK_CHECK(condition) \ + { \ + const VkResult error = condition; \ + if (error != VK_SUCCESS) \ + { \ + std::cerr << "A vulkan error encountered: " << error << " at " \ + << __FILE__ << ':' << __LINE__ << std::endl; \ + std::exit(EXIT_FAILURE); \ + } \ + } + +// OpenCL device that is suitable for this example. +struct cl_device_candidate +{ + /// The OpenCL device id representing the device. + cl::Device device; + + /// The Vulkan-compatible device UUID. + cl_uchar uuid[CL_UUID_SIZE_KHR]; +}; + +// OpenCL and Vulkan physical device suitable for the sample. +struct device_candidate +{ + /// The Vulkan physical device handle of the device to be used. + VkPhysicalDevice vk_candidate; + + /// The candidate device's Vulkan device properties. + VkPhysicalDeviceProperties vk_props; + + /// The OpenCL device candidate that this Vulkan device corresponds to. + struct cl_device_candidate cl_candidate; +}; + +// Check if the extensions supported by a Vulkan device includes a given set of +// required extensions. +template +bool extensions_supported( + const std::vector supported_extensions_properties, + const IteratorT required_device_extensions_begin, + const IteratorT required_device_extensions_end) +{ + IteratorT it = required_device_extensions_begin; + for (; it != required_device_extensions_end; ++it) + { + const auto supported_it = + std::find_if(supported_extensions_properties.begin(), + supported_extensions_properties.end(), + [&](const VkExtensionProperties& props) { + return std::strcmp(*it, props.extensionName) == 0; + }); + if (supported_it == supported_extensions_properties.end()) + { + return false; + } + } + return true; +} + +// Check if a given Vulkan device supports all the required Vulkan extensions. +bool check_device_extensions( + const VkPhysicalDevice vk_device, + const std::vector required_device_extensions) +{ + uint32_t supported_extensions_count; + VK_CHECK(vkEnumerateDeviceExtensionProperties( + vk_device, nullptr, &supported_extensions_count, nullptr)); + std::vector vk_supported_extensions_properties( + supported_extensions_count); + VK_CHECK(vkEnumerateDeviceExtensionProperties( + vk_device, nullptr, &supported_extensions_count, + vk_supported_extensions_properties.data())); + + return extensions_supported(vk_supported_extensions_properties, + required_device_extensions.begin(), + required_device_extensions.end()); +} + +// Check if a given Vulkan physical device is compatible with any of the OpenCL +// devices available. +bool is_vk_device_suitable( + const std::vector cl_candidates, + VkPhysicalDevice vk_device, device_candidate& candidate, + const std::vector required_device_extensions) +{ + // Check if the device supports OpenCL by checking if there is any device + // with the same UUID. + { + // Query the Vulkan device UUID using vkGetPhysicalDeviceProperties2. + VkPhysicalDeviceIDPropertiesKHR id_props = {}; + id_props.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES_KHR; + + VkPhysicalDeviceProperties2KHR props2 = {}; + props2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2_KHR; + props2.pNext = &id_props; + + vkGetPhysicalDeviceProperties2(vk_device, &props2); + + // Look for an OpenCL device which UUID matches the UUID reported by + // Vulkan. + const auto cmp_device_uuid = + [&](const cl_device_candidate& cl_candidate) { + return std::equal(std::begin(cl_candidate.uuid), + std::end(cl_candidate.uuid), + std::begin(id_props.deviceUUID), + std::end(id_props.deviceUUID)); + }; + const auto it = std::find_if(cl_candidates.begin(), cl_candidates.end(), + cmp_device_uuid); + if (it == cl_candidates.end()) + { + // This device does not support HIP. + return false; + } + + candidate.vk_props = props2.properties; + candidate.cl_candidate = *it; + } + + // Check if the device supports the required extensions. + if (!check_device_extensions(vk_device, required_device_extensions)) + { + return false; + } + + candidate.vk_candidate = vk_device; + return true; +} + +// Find a suitable device for the example, that is, an OpenCL +// device that is also Vulkan-compatible and that supports the required +// Vulkan device extensions. +struct device_candidate +find_suitable_device(VkInstance instance, + std::vector required_device_extensions) +{ + // Query OpenCL devices available. + std::vector platforms; + cl::Platform::get(&platforms); + + // For each OpenCL device, query its Vulkan-compatible device UUID and + // add it to the list of candidates. + std::vector cl_candidates; + for (const auto& platform : platforms) + { + std::vector platform_devices; + platform.getDevices(CL_DEVICE_TYPE_ALL, &platform_devices); + + for (const auto& device : platform_devices) + { + if (cl::util::supports_extension(device, "cl_khr_device_uuid")) + { + cl_uchar vk_candidate_uuid[CL_UUID_SIZE_KHR]; + device.getInfo(CL_DEVICE_UUID_KHR, &vk_candidate_uuid); + + cl_device_candidate candidate; + candidate.device = device; + std::memcpy(candidate.uuid, &vk_candidate_uuid, + sizeof(cl_uchar) * CL_UUID_SIZE_KHR); + cl_candidates.push_back(candidate); + } + } + } + + // Query the Vulkan physical devices available. + uint32_t vk_device_count; + VK_CHECK(vkEnumeratePhysicalDevices(instance, &vk_device_count, nullptr)); + + std::vector vk_devices(vk_device_count); + VK_CHECK(vkEnumeratePhysicalDevices(instance, &vk_device_count, + vk_devices.data())); + + // Find a suitable Vulkan physical device compatible with one of the OpenCL + // devices available. + device_candidate candidate; + for (const auto vk_device : vk_devices) + { + if (is_vk_device_suitable(cl_candidates, vk_device, candidate, + required_device_extensions)) + { + return candidate; + } + } + + std::cout << "No suitable OpenCL Vulkan-compatible devices available" + << std::endl; + exit(EXIT_SUCCESS); +} + +// Check if a given OpenCL device supports a particular external memory handle +// type. +bool vk_check_external_memory_handle_type( + VkPhysicalDevice vk_physical_device, + VkExternalMemoryHandleTypeFlagBits vk_external_memory_handle_type) +{ + VkPhysicalDeviceExternalBufferInfo physical_device_external_buffer_info{}; + physical_device_external_buffer_info.sType = + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTERNAL_BUFFER_INFO; + physical_device_external_buffer_info.handleType = + vk_external_memory_handle_type; + + VkExternalBufferProperties external_buffer_properties; + + vkGetPhysicalDeviceExternalBufferProperties( + vk_physical_device, &physical_device_external_buffer_info, + &external_buffer_properties); + + return (vk_external_memory_handle_type + & external_buffer_properties.externalMemoryProperties + .compatibleHandleTypes); +} + +// Find Vulkan memory properties from Vulkan physical device property flags. +uint32_t find_vk_memory_type(VkPhysicalDevice vk_device, uint32_t type_filter, + VkMemoryPropertyFlags properties) +{ + VkPhysicalDeviceMemoryProperties mem_properties; + vkGetPhysicalDeviceMemoryProperties(vk_device, &mem_properties); + for (uint32_t i = 0; i < mem_properties.memoryTypeCount; i++) + { + if ((type_filter & (1 << i)) + && (mem_properties.memoryTypes[i].propertyFlags & properties) + == properties) + { + return i; + } + } + return 0; +} + +#endif // _SAMPLES_CORE_EXTERNALMEMORY_UTILS_HPP