From 5304c711d97b69a41fdf85b1528700497b0429d0 Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Fri, 29 Sep 2023 10:19:00 +0000 Subject: [PATCH 01/13] External memory C/C++ example --- .github/workflows/presubmit.yml | 39 +- cmake/Dependencies.cmake | 2 +- cmake/Dependencies/Vulkan/Vulkan.cmake | 1 + samples/CMakeLists.txt | 4 +- samples/extensions/khr/CMakeLists.txt | 1 + .../khr/externalmemory/CMakeLists.txt | 31 + .../extensions/khr/externalmemory/README.md | 257 +++++++ .../khr/externalmemory/external_saxpy.cl | 8 + samples/extensions/khr/externalmemory/main.c | 654 ++++++++++++++++++ .../extensions/khr/externalmemory/main.cpp | 559 +++++++++++++++ .../khr/externalmemory/vulkan_utils.h | 382 ++++++++++ .../khr/externalmemory/vulkan_utils.hpp | 268 +++++++ 12 files changed, 2202 insertions(+), 4 deletions(-) create mode 100644 cmake/Dependencies/Vulkan/Vulkan.cmake create mode 100644 samples/extensions/khr/externalmemory/CMakeLists.txt create mode 100644 samples/extensions/khr/externalmemory/README.md create mode 100644 samples/extensions/khr/externalmemory/external_saxpy.cl create mode 100644 samples/extensions/khr/externalmemory/main.c create mode 100644 samples/extensions/khr/externalmemory/main.cpp create mode 100644 samples/extensions/khr/externalmemory/vulkan_utils.h create mode 100644 samples/extensions/khr/externalmemory/vulkan_utils.hpp diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index 8a24a42a..a13bf43f 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -198,6 +198,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}}" @@ -416,6 +427,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: @@ -457,6 +470,18 @@ jobs: fetch-depth: 0 submodules: recursive + - 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: Configure (MSBuild) if: matrix.GEN == 'Visual Studio 17 2022' run: | @@ -479,6 +504,8 @@ jobs: -D CMAKE_C_FLAGS="${env:CMAKE_CFLAGS}" ` -D CMAKE_CXX_FLAGS="${env:CMAKE_CXXFLAGS}" ` -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." } @@ -509,6 +536,8 @@ jobs: -D CMAKE_CXX_FLAGS="${env:CMAKE_CXXFLAGS}" ` -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." } @@ -546,12 +575,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." } } @@ -664,6 +695,10 @@ jobs: echo "OCL_ICD_VENDORS=$POCL_INSTALL_PATH/etc/OpenCL/vendors" >> $GITHUB_ENV 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 @@ -718,7 +753,7 @@ jobs: - name: Test working-directory: ${{runner.workspace}}/OpenCL-SDK/build run: | - EXCLUDE_REGEX="multidevice.*" + EXCLUDE_REGEX="(multidevice|externalmemory).*" ctest -C Debug --output-on-failure --no-tests=error --parallel `sysctl -n hw.logicalcpu` --exclude-regex "$EXCLUDE_REGEX" ctest -C Release --output-on-failure --no-tests=error --parallel `sysctl -n hw.logicalcpu` --exclude-regex "$EXCLUDE_REGEX" diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index e9777716..24383896 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -41,7 +41,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 0834853a..2eae373d 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 From 4b430eb36e71057deccba0e96e5ecf6fc1f095d1 Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Thu, 14 Dec 2023 10:44:15 +0000 Subject: [PATCH 02/13] Fixes from review --- CMakeLists.txt | 1 + cmake/Dependencies.cmake | 7 +- samples/extensions/khr/CMakeLists.txt | 4 +- samples/extensions/khr/externalmemory/main.c | 72 +++++++++++++++-- .../extensions/khr/externalmemory/main.cpp | 79 +++++++++++++++++-- .../khr/externalmemory/vulkan_utils.h | 24 +++--- 6 files changed, 160 insertions(+), 27 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 9cbab822..d54942c1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -26,6 +26,7 @@ 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" ON 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_TEST_SAMPLES "Add CTest to samples (where applicable)" ON OPENCL_SDK_BUILD_SAMPLES OFF) option(OPENCL_SDK_BUILD_CLINFO "Build clinfo utility" ON) diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 24383896..06145cb8 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -41,7 +41,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 Vulkan) + foreach(DEP IN ITEMS cargs TCLAP Stb) list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/Dependencies/${DEP}") include(${DEP}) endforeach() @@ -52,6 +52,11 @@ if(OPENCL_SDK_BUILD_SAMPLES) include(${DEP}) endforeach() endif(OPENCL_SDK_BUILD_OPENGL_SAMPLES) + + if(OPENCL_SDK_BUILD_VULKAN_SAMPLES) + list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/Dependencies/Vulkan") + include(Vulkan) + endif(OPENCL_SDK_BUILD_VULKAN_SAMPLES) endif(OPENCL_SDK_BUILD_SAMPLES) if(OPENCL_SDK_BUILD_CLINFO) diff --git a/samples/extensions/khr/CMakeLists.txt b/samples/extensions/khr/CMakeLists.txt index 4bf194d7..efc6d747 100644 --- a/samples/extensions/khr/CMakeLists.txt +++ b/samples/extensions/khr/CMakeLists.txt @@ -12,9 +12,11 @@ # 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) add_subdirectory(nbody) endif() +if(OPENCL_SDK_BUILD_VULKAN_SAMPLES) + add_subdirectory(externalmemory) +endif() diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index c49ff2c6..b7698fc6 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -430,7 +430,6 @@ int main(int argc, char* argv[]) 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; @@ -485,6 +484,35 @@ int main(int argc, char* argv[]) memcpy(vk_arr_x, arr_x, sizeof(cl_float) * length); memcpy(vk_arr_y, arr_y, sizeof(cl_float) * length); +#ifdef _WIN32 + // Get Vulkan external memory file descriptors for accessing external memory + // with OpenCL. + VkMemoryGetWin32HandleInfoKHR handle_info_x = { 0 }; + handle_info_x.sType = VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR; + handle_info_x.pNext = NULL; + handle_info_x.memory = vk_buf_x_memory; + handle_info_x.handleType = vk_external_memory_handle_type; + HANDLE handle_x; + + VkMemoryGetWin32HandleInfoKHR handle_info_y = { 0 }; + handle_info_y.sType = VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR; + handle_info_y.pNext = NULL; + handle_info_y.memory = vk_buf_y_memory; + handle_info_y.handleType = vk_external_memory_handle_type; + HANDLE handle_y; + + // We need to get the pointer to the + // vkGetMemoryFdKHR/vkGetMemoryWin32HandleKHR function because it's from + // extension VK_KHR_external_memory_fd. This Vulkan function exports a POSIX + // file descriptor/Windows handle referencing the payload of a Vulkan device + // memory object. + PFN_vkGetMemoryWin32HandleKHR vkGetMemoryWin32Handle; + *(PFN_vkGetMemoryWin32HandleKHR*)&vkGetMemoryWin32Handle = + (PFN_vkGetMemoryWin32HandleKHR)vkGetDeviceProcAddr( + vk_device, "vkGetMemoryWin32HandleKHR"); + VK_CHECK(vkGetMemoryWin32Handle(vk_device, &handle_info_x, &handle_x)); + VK_CHECK(vkGetMemoryWin32Handle(vk_device, &handle_info_y, &handle_y)); +#else // Get Vulkan external memory file descriptors for accessing external memory // with OpenCL. VkMemoryGetFdInfoKHR fd_info_x = { 0 }; @@ -501,19 +529,28 @@ int main(int argc, char* argv[]) 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 = + // We need to get the pointer to the + // vkGetMemoryFdKHR/vkGetMemoryWin32HandleKHR function because it's from + // extension VK_KHR_external_memory_fd. This Vulkan function exports a POSIX + // file descriptor/Windows handle referencing the payload of a Vulkan device + // memory object. + PFN_vkGetMemoryFdKHR vkGetMemoryFd; + *(PFN_vkGetMemoryFdKHR*)&vkGetMemoryFd = (PFN_vkGetMemoryFdKHR)vkGetDeviceProcAddr(vk_device, "vkGetMemoryFdKHR"); + VK_CHECK(vkGetMemoryFd(vk_device, &fd_info_x, &fd_x)); + VK_CHECK(vkGetMemoryFd(vk_device, &fd_info_y, &fd_y)); +#endif - 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, +#ifdef _WIN32 + (cl_mem_properties)handle_x, +#else (cl_mem_properties)fd_x, +#endif (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)(uintptr_t)cl_device, CL_DEVICE_HANDLE_LIST_END_KHR, @@ -521,7 +558,11 @@ int main(int argc, char* argv[]) }; cl_mem_properties ext_mem_props_y[] = { (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, +#ifdef _WIN32 + (cl_mem_properties)handle_y, +#else (cl_mem_properties)fd_y, +#endif (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)(uintptr_t)cl_device, CL_DEVICE_HANDLE_LIST_END_KHR, @@ -552,6 +593,16 @@ int main(int argc, char* argv[]) OCLERROR_RET(clSetKernelArg(saxpy, 2, sizeof(cl_mem), &cl_buf_y), error, clbufy); + // Acquire OpenCL memory objects created from Vulkan external memory + // handles. + cl_mem cl_mem_objects[] = { cl_buf_x, cl_buf_y }; + clEnqueueAcquireExternalMemObjectsKHR_fn + clEnqueueAcquireExternalMemObjects = + (clEnqueueAcquireExternalMemObjectsKHR_fn) + clGetExtensionFunctionAddressForPlatform( + cl_platform, "clEnqueueAcquireExternalMemObjectsKHR"); + clEnqueueAcquireExternalMemObjects(queue, 2, cl_mem_objects, 0, NULL, NULL); + // Launch kernel. if (diag_opts.verbose) { @@ -570,6 +621,15 @@ int main(int argc, char* argv[]) cl_ulong dev_time; TIMER_DIFFERENCE(dev_time, dev_start, dev_end) + // Release OpenCL memory objects created from Vulkan external memory + // handles. + clEnqueueReleaseExternalMemObjectsKHR_fn + clEnqueueReleaseExternalMemObjects = + (clEnqueueReleaseExternalMemObjectsKHR_fn) + clGetExtensionFunctionAddressForPlatform( + cl_platform, "clEnqueueReleaseExternalMemObjectsKHR"); + clEnqueueReleaseExternalMemObjects(queue, 2, cl_mem_objects, 0, NULL, NULL); + // Concurrently calculate reference saxpy. if (diag_opts.verbose) { diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp index e4026c23..c82d8994 100644 --- a/samples/extensions/khr/externalmemory/main.cpp +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -349,7 +349,6 @@ int main(int argc, char* argv[]) 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; @@ -404,6 +403,37 @@ int main(int argc, char* argv[]) memcpy(vk_arr_x, arr_x.data(), sizeof(cl_float) * length); memcpy(vk_arr_y, arr_y.data(), sizeof(cl_float) * length); +#ifdef _WIN32 + // Get Vulkan external memory file descriptors for accessing external + // memory with OpenCL. + VkMemoryGetWin32HandleInfoKHR handle_info_x{}; + handle_info_x.sType = + VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR; + handle_info_x.pNext = nullptr; + handle_info_x.memory = vk_buf_x_memory; + handle_info_x.handleType = vk_external_memory_handle_type; + HANDLE handle_x; + + VkMemoryGetWin32HandleInfoKHR handle_info_y{}; + handle_info_y.sType = + VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR; + handle_info_y.pNext = nullptr; + handle_info_y.memory = vk_buf_y_memory; + handle_info_y.handleType = vk_external_memory_handle_type; + HANDLE handle_y; + + // We need to get the pointer to the + // vkGetMemoryFdKHR/vkGetMemoryWin32HandleKHR function because it's from + // the extension VK_KHR_external_memory_fd. This Vulkan function exports + // a POSIX file descriptor/Windows handle referencing the payload of a + // Vulkan device memory object. + PFN_vkGetMemoryWin32HandleKHR vkGetMemoryWin32Handle; + *(PFN_vkGetMemoryWin32HandleKHR*)&vkGetMemoryWin32Handle = + (PFN_vkGetMemoryWin32HandleKHR)vkGetDeviceProcAddr( + vk_device, "vkGetMemoryWin32HandleKHR"); + VK_CHECK(vkGetMemoryWin32Handle(vk_device, &handle_info_x, &handle_x)); + VK_CHECK(vkGetMemoryWin32Handle(vk_device, &handle_info_y, &handle_y)); +#else // Get Vulkan external memory file descriptors for accessing external // memory with OpenCL. VkMemoryGetFdInfoKHR fd_info_x{}; @@ -420,19 +450,27 @@ int main(int argc, char* argv[]) 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 = + // We need to get the pointer to the + // vkGetMemoryFdKHR/vkGetMemoryWin32HandleKHR function because it's from + // extension VK_KHR_external_memory_fd. This Vulkan function exports a + // POSIX file descriptor/Windows handle referencing the payload of a + // Vulkan device memory object. + PFN_vkGetMemoryFdKHR vkGetMemoryFd; + *(PFN_vkGetMemoryFdKHR*)&vkGetMemoryFd = (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)); + VK_CHECK(vkGetMemoryFd(vk_device, &fd_info_x, &fd_x)); + VK_CHECK(vkGetMemoryFd(vk_device, &fd_info_y, &fd_y)); +#endif // 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, +#ifdef _WIN32 + (cl_mem_properties)handle_x, +#else (cl_mem_properties)fd_x, +#endif (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)cl_device(), CL_DEVICE_HANDLE_LIST_END_KHR, @@ -440,7 +478,11 @@ int main(int argc, char* argv[]) }; std::vector ext_mem_props_y = { (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, +#ifdef _WIN32 + (cl_mem_properties)handle_y, +#else (cl_mem_properties)fd_y, +#endif (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)cl_device(), CL_DEVICE_HANDLE_LIST_END_KHR, @@ -457,6 +499,18 @@ int main(int argc, char* argv[]) 0 }; cl::CommandQueue queue{ cl_context, cl_device, *queue_props }; + // Acquire OpenCL memory objects created from Vulkan external memory + // handles. + std::vector cl_mem_objects = { cl_buf_x(), cl_buf_y() }; + clEnqueueAcquireExternalMemObjectsKHR_fn + clEnqueueAcquireExternalMemObjects = + (clEnqueueAcquireExternalMemObjectsKHR_fn) + clGetExtensionFunctionAddressForPlatform( + cl_platform(), "clEnqueueAcquireExternalMemObjectsKHR"); + clEnqueueAcquireExternalMemObjects( + queue(), static_cast(cl_mem_objects.size()), + cl_mem_objects.data(), 0, nullptr, nullptr); + // Launch kernel. if (diag_opts.verbose) { @@ -472,6 +526,17 @@ int main(int argc, char* argv[]) cl::WaitForEvents(kernel_run); auto dev_end = std::chrono::high_resolution_clock::now(); + // Release OpenCL memory objects created from Vulkan external memory + // handles. + clEnqueueReleaseExternalMemObjectsKHR_fn + clEnqueueReleaseExternalMemObjects = + (clEnqueueReleaseExternalMemObjectsKHR_fn) + clGetExtensionFunctionAddressForPlatform( + cl_platform(), "clEnqueueReleaseExternalMemObjectsKHR"); + clEnqueueReleaseExternalMemObjects( + queue(), static_cast(cl_mem_objects.size()), + cl_mem_objects.data(), 0, nullptr, nullptr); + // Concurrently calculate reference saxpy. if (diag_opts.verbose) { diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.h b/samples/extensions/khr/externalmemory/vulkan_utils.h index 2b718564..130bc45e 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.h +++ b/samples/extensions/khr/externalmemory/vulkan_utils.h @@ -224,24 +224,24 @@ find_suitable_device(VkInstance instance, // Query OpenCL devices available. cl_int error = CL_SUCCESS; bool candidate_found = false; - cl_uint cl_platform_count = 0; + cl_uint platform_count = 0; struct device_candidate found_candidate = {0}; - OCLERROR_RET(clGetPlatformIDs(0, NULL, &cl_platform_count), error, ret); + OCLERROR_RET(clGetPlatformIDs(0, NULL, &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, + (cl_platform_id*)malloc(platform_count * sizeof(cl_platform_id)); + OCLERROR_RET(clGetPlatformIDs(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) + for (cl_uint platform_id = 0; platform_id < platform_count; + ++platform_id) { cl_uint cl_platform_devices_count = 0; - OCLERROR_RET(clGetDeviceIDs(platforms[cl_platform_id], + OCLERROR_RET(clGetDeviceIDs(platforms[platform_id], CL_DEVICE_TYPE_ALL, 0, NULL, &cl_platform_devices_count), error, platforms); @@ -250,7 +250,7 @@ find_suitable_device(VkInstance instance, { cl_device_id device; OCLERROR_PAR(device = cl_util_get_device( - cl_platform_id, device_id, CL_DEVICE_TYPE_ALL, &error), error, platforms); + platform_id, device_id, CL_DEVICE_TYPE_ALL, &error), error, platforms); cl_device_count += check_khronos_extensions(device, uuid_khronos_extension, 1); } @@ -264,11 +264,11 @@ find_suitable_device(VkInstance instance, (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) + for (cl_uint platform_id = 0; platform_id < platform_count; + ++platform_id) { cl_uint cl_platform_devices_count = 0; - OCLERROR_RET(clGetDeviceIDs(platforms[cl_platform_id], + OCLERROR_RET(clGetDeviceIDs(platforms[platform_id], CL_DEVICE_TYPE_ALL, 0, NULL, &cl_platform_devices_count), error, candidates); @@ -278,7 +278,7 @@ find_suitable_device(VkInstance instance, ++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); + 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]; From 14ce09fc5a6c4f266196e30a0c31750a322db034 Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 29 May 2024 10:45:54 +0200 Subject: [PATCH 03/13] Fix find_suitable_device logic to conditionally increment cl_device_count --- samples/extensions/khr/externalmemory/vulkan_utils.h | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.h b/samples/extensions/khr/externalmemory/vulkan_utils.h index 130bc45e..f2d82554 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.h +++ b/samples/extensions/khr/externalmemory/vulkan_utils.h @@ -256,6 +256,12 @@ find_suitable_device(VkInstance instance, } } + if (!cl_device_count) + { + printf("No suitable OpenCL Vulkan-compatible devices available\n"); + goto platforms; + } + // 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 @@ -275,7 +281,7 @@ find_suitable_device(VkInstance instance, for (cl_uint cl_candidate_id = 0; cl_candidate_id < cl_platform_devices_count; - ++cl_candidate_id, ++cl_device_count) + ++cl_candidate_id) { cl_device_id device = cl_util_get_device( platform_id, cl_candidate_id, CL_DEVICE_TYPE_ALL, &error); @@ -292,6 +298,7 @@ find_suitable_device(VkInstance instance, memcpy(candidate.uuid, &vk_candidate_uuid, sizeof(cl_uchar) * CL_UUID_SIZE_KHR); cl_candidates[cl_device_count] = candidate; + cl_device_count++; } } } From 5ba0d26d18b73e5f4d3d0559ae05dccf35e5ee8a Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 29 May 2024 11:23:22 +0200 Subject: [PATCH 04/13] Conform to OpenCL v3.0.15 specification release --- samples/extensions/khr/externalmemory/main.c | 8 ++++---- samples/extensions/khr/externalmemory/main.cpp | 8 ++++---- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index b7698fc6..f00b7036 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -551,9 +551,9 @@ int main(int argc, char* argv[]) #else (cl_mem_properties)fd_x, #endif - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)(uintptr_t)cl_device, - CL_DEVICE_HANDLE_LIST_END_KHR, + CL_MEM_DEVICE_HANDLE_LIST_END_KHR, 0 }; cl_mem_properties ext_mem_props_y[] = { @@ -563,9 +563,9 @@ int main(int argc, char* argv[]) #else (cl_mem_properties)fd_y, #endif - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)(uintptr_t)cl_device, - CL_DEVICE_HANDLE_LIST_END_KHR, + CL_MEM_DEVICE_HANDLE_LIST_END_KHR, 0 }; cl_mem cl_buf_x, cl_buf_y; diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp index c82d8994..97b99a29 100644 --- a/samples/extensions/khr/externalmemory/main.cpp +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -471,9 +471,9 @@ int main(int argc, char* argv[]) #else (cl_mem_properties)fd_x, #endif - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)cl_device(), - CL_DEVICE_HANDLE_LIST_END_KHR, + CL_MEM_DEVICE_HANDLE_LIST_END_KHR, 0 }; std::vector ext_mem_props_y = { @@ -483,9 +483,9 @@ int main(int argc, char* argv[]) #else (cl_mem_properties)fd_y, #endif - (cl_mem_properties)CL_DEVICE_HANDLE_LIST_KHR, + (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, (cl_mem_properties)cl_device(), - CL_DEVICE_HANDLE_LIST_END_KHR, + CL_MEM_DEVICE_HANDLE_LIST_END_KHR, 0 }; From 9717d64d612862c309d5dd796e0435c58205a26c Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 5 Jun 2024 06:53:44 +0000 Subject: [PATCH 05/13] Fix kernel execution time report --- samples/extensions/khr/externalmemory/main.c | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index f00b7036..d0d3bd28 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -671,8 +671,7 @@ int main(int argc, char* argv[]) 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", + printf("Kernel execution time as measured by device: %llu us.\n", (unsigned long long)(cl_util_get_event_duration( kernel_run, CL_PROFILING_COMMAND_START, CL_PROFILING_COMMAND_END, &error) From 5d959f9d506e03da8f6490edad48f729bceb0111 Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 2 Oct 2024 10:41:34 +0200 Subject: [PATCH 06/13] Fix opencl version check --- samples/extensions/khr/externalmemory/main.c | 38 +++++++++++++++---- .../extensions/khr/externalmemory/main.cpp | 29 ++++++++++---- 2 files changed, 53 insertions(+), 14 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index d0d3bd28..1c466a25 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -217,6 +217,11 @@ bool cl_check_external_memory_handle_type( exit(EXIT_FAILURE); } +cl_int opencl_version_is_major(cl_name_version* dev_name_version, cl_uint major) +{ + return CL_VERSION_MAJOR(dev_name_version->version) == major; +} + int main(int argc, char* argv[]) { cl_int error = CL_SUCCESS; @@ -343,13 +348,32 @@ int main(int argc, char* argv[]) 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 + // Get number of versions supported. + size_t versions_size = 0; + OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, + NULL, &versions_size), + error, end); + size_t versions_count = versions_size / sizeof(cl_name_version); + + // Get and check versions. + cl_name_version* dev_versions = (cl_name_version*)malloc(versions_size); + OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, + versions_size, dev_versions, NULL), + error, end); + char compiler_options[1024] = ""; + for (cl_uint i = 0; i < versions_count; ++i) + { + if (opencl_version_is_major(&dev_versions[i], 3)) + { + strcat(compiler_options, "-cl-std=CL3.0 "); + } + } + + if (compiler_options[0] == '\0') + { + fprintf(stderr, "\nError: OpenCL version must be at least 3.0\n"); + exit(EXIT_FAILURE); + } OCLERROR_RET(cl_util_build_program(program, cl_device, compiler_options), error, prg); diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp index 97b99a29..9cbaf433 100644 --- a/samples/extensions/khr/externalmemory/main.cpp +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -138,6 +138,12 @@ bool cl_check_external_memory_handle_type( return it != supported_handle_types.end(); } +bool opencl_version_is_major(const cl_name_version& dev_name_version, + const cl_uint& major) +{ + return CL_VERSION_MAJOR(dev_name_version.version) == major; +} + int main(int argc, char* argv[]) { try @@ -267,13 +273,22 @@ int main(int argc, char* argv[]) // 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 + std::vector dev_versions = + cl_device.getInfo(); + for (cl_name_version dev_name_version : dev_versions) + { + if (opencl_version_is_major(dev_name_version, 3)) + { + compiler_options += cl::string{ "-cl-std=CL3.0 " }; + } + } + + if (compiler_options.empty()) + { + std::cerr << "\nError: OpenCL version must be at least 3.0" + << std::endl; + exit(EXIT_FAILURE); + } cl_program.build(cl_device, compiler_options.c_str()); From 989bc643e5a691cc3407782752cc0a2ee3810cab Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 2 Oct 2024 20:35:31 +0200 Subject: [PATCH 07/13] Fix logic for finding suitable opencl device --- samples/extensions/khr/externalmemory/vulkan_utils.h | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.h b/samples/extensions/khr/externalmemory/vulkan_utils.h index f2d82554..47277d43 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.h +++ b/samples/extensions/khr/externalmemory/vulkan_utils.h @@ -241,10 +241,14 @@ find_suitable_device(VkInstance instance, ++platform_id) { cl_uint cl_platform_devices_count = 0; - OCLERROR_RET(clGetDeviceIDs(platforms[platform_id], - CL_DEVICE_TYPE_ALL, 0, NULL, - &cl_platform_devices_count), - error, platforms); + error = clGetDeviceIDs(platforms[platform_id], CL_DEVICE_TYPE_ALL, 0, + NULL, &cl_platform_devices_count); + // Some platforms may not have any suitable device. Allow the CL_DEVICE_NOT_FOUND + // error so that other platforms can be checked. + if (error != CL_SUCCESS && error != CL_DEVICE_NOT_FOUND) + { + goto platforms; + } for (cl_uint device_id = 0; device_id < cl_platform_devices_count; ++device_id) { From 8bad1b0e9f019b112f69791b003bfeabb5686dca Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 6 Nov 2024 19:55:26 +0000 Subject: [PATCH 08/13] Inform when kernel source is not found --- samples/extensions/khr/externalmemory/main.c | 36 +++++++++++--------- 1 file changed, 20 insertions(+), 16 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index 1c466a25..6e3c9a48 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -338,9 +338,12 @@ int main(int argc, char* argv[]) 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); + kernel = cl_util_read_text_file(kernel_location, &program_size, &error); + if (error != CL_SUCCESS) + { + fprintf(stderr, "Cannot open kernel source: %s\n", kernel_location); + goto cont; + } MEM_CHECK(tmp = (char*)realloc(kernel, program_size), error, ker); kernel = tmp; OCLERROR_PAR(program = clCreateProgramWithSource( @@ -352,14 +355,14 @@ int main(int argc, char* argv[]) size_t versions_size = 0; OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, NULL, &versions_size), - error, end); + error, prg); size_t versions_count = versions_size / sizeof(cl_name_version); // Get and check versions. cl_name_version* dev_versions = (cl_name_version*)malloc(versions_size); OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, versions_size, dev_versions, NULL), - error, end); + error, prg); char compiler_options[1024] = ""; for (cl_uint i = 0; i < versions_count; ++i) { @@ -424,7 +427,7 @@ int main(int argc, char* argv[]) { fprintf(stderr, "\nError: Unsupported OpenCL external memory handle type\n"); - exit(EXIT_FAILURE); + goto arry; } if (!vk_check_external_memory_handle_type(vk_physical_device, @@ -432,7 +435,7 @@ int main(int argc, char* argv[]) { fprintf(stderr, "\nError: Unsupported Vulkan external memory handle type\n"); - exit(EXIT_FAILURE); + goto arry; } // Initialize Vulkan device-side storage. @@ -608,14 +611,14 @@ int main(int argc, char* argv[]) 0 }; OCLERROR_PAR(queue = clCreateCommandQueueWithProperties( context, cl_device, queue_props, &error), - error, cont); + error, clbufy); // Set kernel arguments. - OCLERROR_RET(clSetKernelArg(saxpy, 0, sizeof(cl_float), &a), error, clbufy); + OCLERROR_RET(clSetKernelArg(saxpy, 0, sizeof(cl_float), &a), error, que); OCLERROR_RET(clSetKernelArg(saxpy, 1, sizeof(cl_mem), &cl_buf_x), error, - clbufy); + que); OCLERROR_RET(clSetKernelArg(saxpy, 2, sizeof(cl_mem), &cl_buf_y), error, - clbufy); + que); // Acquire OpenCL memory objects created from Vulkan external memory // handles. @@ -638,8 +641,8 @@ int main(int argc, char* argv[]) 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); + error, que); + OCLERROR_RET(clWaitForEvents(1, &kernel_run), error, que); GET_CURRENT_TIMER(dev_end) cl_ulong dev_time; @@ -675,7 +678,7 @@ int main(int argc, char* argv[]) OCLERROR_RET(clEnqueueReadBuffer(queue, cl_buf_y, CL_BLOCKING, 0, sizeof(cl_float) * length, (void*)arr_x, 0, NULL, NULL), - error, clbufy); + error, que); // Validate solution. for (size_t i = 0; i < length; ++i) @@ -707,6 +710,8 @@ int main(int argc, char* argv[]) } // Release resources. +que: + OCLERROR_RET(clReleaseCommandQueue(queue), end_error, cont); clbufy: OCLERROR_RET(clReleaseMemObject(cl_buf_y), end_error, clbufx); clbufx: @@ -718,6 +723,7 @@ int main(int argc, char* argv[]) vkUnmapMemory(vk_device, vk_buf_x_memory); vkFreeMemory(vk_device, vk_buf_y_memory, NULL); vkFreeMemory(vk_device, vk_buf_x_memory, NULL); +arry: free(arr_y); arrx: free(arr_x); @@ -727,8 +733,6 @@ int main(int argc, char* argv[]) 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: From 168014bf94c3f7ba1104b37ed9b69ad9401cd4f2 Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 6 Nov 2024 20:17:31 +0000 Subject: [PATCH 09/13] Set VkPhysicalDeviceExternalBufferInfo->usage --- samples/extensions/khr/externalmemory/main.c | 7 +++++-- samples/extensions/khr/externalmemory/main.cpp | 9 ++++++--- samples/extensions/khr/externalmemory/vulkan_utils.h | 2 ++ samples/extensions/khr/externalmemory/vulkan_utils.hpp | 5 ++++- 4 files changed, 17 insertions(+), 6 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index 6e3c9a48..330576b4 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -430,7 +430,11 @@ int main(int argc, char* argv[]) goto arry; } + VkBufferUsageFlags vk_external_memory_usage = + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT; + if (!vk_check_external_memory_handle_type(vk_physical_device, + vk_external_memory_usage, vk_external_memory_handle_type)) { fprintf(stderr, @@ -455,8 +459,7 @@ int main(int argc, char* argv[]) 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.usage = vk_external_memory_usage; buffer_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE; VkBuffer vk_buf_x, vk_buf_y; diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp index 9cbaf433..2f6c7e8b 100644 --- a/samples/extensions/khr/externalmemory/main.cpp +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -335,8 +335,12 @@ int main(int argc, char* argv[]) exit(EXIT_FAILURE); } + VkBufferUsageFlags vk_external_memory_usage = + VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT; + if (!vk_check_external_memory_handle_type( - vk_physical_device, vk_external_memory_handle_type)) + vk_physical_device, vk_external_memory_usage, + vk_external_memory_handle_type)) { std::cerr << "\nError: Unsupported Vulkan external memory handle type" @@ -362,8 +366,7 @@ int main(int argc, char* argv[]) 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.usage = vk_external_memory_usage; buffer_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE; VkBuffer vk_buf_x, vk_buf_y; diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.h b/samples/extensions/khr/externalmemory/vulkan_utils.h index 47277d43..b8b91d24 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.h +++ b/samples/extensions/khr/externalmemory/vulkan_utils.h @@ -353,11 +353,13 @@ find_suitable_device(VkInstance instance, // type. bool vk_check_external_memory_handle_type( VkPhysicalDevice vk_physical_device, + VkBufferUsageFlags vk_external_memory_usage, 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.usage = vk_external_memory_usage; physical_device_external_buffer_info.handleType = vk_external_memory_handle_type; diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.hpp b/samples/extensions/khr/externalmemory/vulkan_utils.hpp index 90090dfd..82c5c1c0 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.hpp +++ b/samples/extensions/khr/externalmemory/vulkan_utils.hpp @@ -228,11 +228,14 @@ find_suitable_device(VkInstance instance, // type. bool vk_check_external_memory_handle_type( VkPhysicalDevice vk_physical_device, + VkBufferUsageFlags vk_external_memory_usage, VkExternalMemoryHandleTypeFlagBits vk_external_memory_handle_type) { - VkPhysicalDeviceExternalBufferInfo physical_device_external_buffer_info{}; + 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.usage = vk_external_memory_usage; physical_device_external_buffer_info.handleType = vk_external_memory_handle_type; From f168b3f6b320689aadc36d0dd1e6ba0ef39e1443 Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Wed, 6 Nov 2024 22:52:16 +0000 Subject: [PATCH 10/13] Set VkExternalBufferProperties->sType and pNext --- samples/extensions/khr/externalmemory/vulkan_utils.h | 4 +++- samples/extensions/khr/externalmemory/vulkan_utils.hpp | 5 ++++- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.h b/samples/extensions/khr/externalmemory/vulkan_utils.h index b8b91d24..ff6730e0 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.h +++ b/samples/extensions/khr/externalmemory/vulkan_utils.h @@ -363,7 +363,9 @@ bool vk_check_external_memory_handle_type( physical_device_external_buffer_info.handleType = vk_external_memory_handle_type; - VkExternalBufferProperties external_buffer_properties; + VkExternalBufferProperties external_buffer_properties = { + VK_STRUCTURE_TYPE_EXTERNAL_BUFFER_PROPERTIES, NULL + }; vkGetPhysicalDeviceExternalBufferProperties( vk_physical_device, &physical_device_external_buffer_info, diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.hpp b/samples/extensions/khr/externalmemory/vulkan_utils.hpp index 82c5c1c0..ac110193 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.hpp +++ b/samples/extensions/khr/externalmemory/vulkan_utils.hpp @@ -239,7 +239,10 @@ bool vk_check_external_memory_handle_type( physical_device_external_buffer_info.handleType = vk_external_memory_handle_type; - VkExternalBufferProperties external_buffer_properties; + VkExternalBufferProperties external_buffer_properties = {}; + external_buffer_properties.sType = + VK_STRUCTURE_TYPE_EXTERNAL_BUFFER_PROPERTIES; + external_buffer_properties.pNext = nullptr; vkGetPhysicalDeviceExternalBufferProperties( vk_physical_device, &physical_device_external_buffer_info, From b1d906a18d7e5a78b8e95f6f2a4d2a54e54165a8 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 6 Nov 2024 21:08:46 -0800 Subject: [PATCH 11/13] fixes for Windows --- samples/extensions/khr/externalmemory/main.c | 6 ++++-- samples/extensions/khr/externalmemory/main.cpp | 6 ++++-- samples/extensions/khr/externalmemory/vulkan_utils.h | 10 ++++++---- 3 files changed, 14 insertions(+), 8 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index 330576b4..59ea67c9 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -575,10 +575,11 @@ int main(int argc, char* argv[]) // 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, #ifdef _WIN32 + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR, (cl_mem_properties)handle_x, #else + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, (cl_mem_properties)fd_x, #endif (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, @@ -587,10 +588,11 @@ int main(int argc, char* argv[]) 0 }; cl_mem_properties ext_mem_props_y[] = { - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, #ifdef _WIN32 + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR, (cl_mem_properties)handle_y, #else + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, (cl_mem_properties)fd_y, #endif (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp index 2f6c7e8b..ef929819 100644 --- a/samples/extensions/khr/externalmemory/main.cpp +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -483,10 +483,11 @@ int main(int argc, char* argv[]) // 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, #ifdef _WIN32 + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR, (cl_mem_properties)handle_x, #else + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, (cl_mem_properties)fd_x, #endif (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, @@ -495,10 +496,11 @@ int main(int argc, char* argv[]) 0 }; std::vector ext_mem_props_y = { - (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, #ifdef _WIN32 + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR, (cl_mem_properties)handle_y, #else + (cl_mem_properties)CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR, (cl_mem_properties)fd_y, #endif (cl_mem_properties)CL_MEM_DEVICE_HANDLE_LIST_KHR, diff --git a/samples/extensions/khr/externalmemory/vulkan_utils.h b/samples/extensions/khr/externalmemory/vulkan_utils.h index ff6730e0..74b2c561 100644 --- a/samples/extensions/khr/externalmemory/vulkan_utils.h +++ b/samples/extensions/khr/externalmemory/vulkan_utils.h @@ -278,10 +278,12 @@ find_suitable_device(VkInstance instance, ++platform_id) { cl_uint cl_platform_devices_count = 0; - OCLERROR_RET(clGetDeviceIDs(platforms[platform_id], - CL_DEVICE_TYPE_ALL, 0, NULL, - &cl_platform_devices_count), - error, candidates); + error = clGetDeviceIDs(platforms[platform_id], CL_DEVICE_TYPE_ALL, 0, + NULL, &cl_platform_devices_count); + if (error != CL_SUCCESS && error != CL_DEVICE_NOT_FOUND) + { + goto candidates; + } for (cl_uint cl_candidate_id = 0; cl_candidate_id < cl_platform_devices_count; From 2c3b6dc01db097e2bbd00179e2c8608f23fed0ed Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Sun, 16 Mar 2025 11:35:41 +0000 Subject: [PATCH 12/13] Fix external memory handle types loop count --- samples/extensions/khr/externalmemory/main.c | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index 59ea67c9..8d22e130 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -183,22 +183,26 @@ bool cl_check_external_memory_handle_type( 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; + size_t supported_handle_types_byte_count = 0; + const size_t handle_type_size = sizeof(cl_external_memory_handle_type_khr); cl_int error = CL_SUCCESS; OCLERROR_RET( clGetDeviceInfo(cl_device, CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, 0, - NULL, &supported_handle_types_count), + NULL, &supported_handle_types_byte_count), error, err); supported_handle_types = (cl_external_memory_handle_type_khr*)malloc( - supported_handle_types_count); + supported_handle_types_byte_count); OCLERROR_RET( clGetDeviceInfo( cl_device, CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, - supported_handle_types_count, supported_handle_types, NULL), + supported_handle_types_byte_count, supported_handle_types, NULL), error, err); + + const size_t supported_handle_types_count = + supported_handle_types_byte_count / handle_type_size; for (size_t i = 0; i < supported_handle_types_count; ++i) { if (external_memory_handle_type == supported_handle_types[i]) From 5f52d0d5660409480a7d57cea5115babbd5a838b Mon Sep 17 00:00:00 2001 From: Beatriz Navidad Vilches Date: Sun, 16 Mar 2025 11:58:40 +0000 Subject: [PATCH 13/13] Remove unnecesary OpenCL kernel compilation to 3.0 --- samples/extensions/khr/externalmemory/main.c | 34 +++---------------- .../extensions/khr/externalmemory/main.cpp | 22 ++---------- 2 files changed, 6 insertions(+), 50 deletions(-) diff --git a/samples/extensions/khr/externalmemory/main.c b/samples/extensions/khr/externalmemory/main.c index 8d22e130..0ab03568 100644 --- a/samples/extensions/khr/externalmemory/main.c +++ b/samples/extensions/khr/externalmemory/main.c @@ -354,36 +354,10 @@ int main(int argc, char* argv[]) context, 1, (const char**)&kernel, &program_size, &error), error, ker); - // The Khronos extension showcased requires OpenCL 3.0 version. - // Get number of versions supported. - size_t versions_size = 0; - OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, - NULL, &versions_size), - error, prg); - size_t versions_count = versions_size / sizeof(cl_name_version); - - // Get and check versions. - cl_name_version* dev_versions = (cl_name_version*)malloc(versions_size); - OCLERROR_RET(clGetDeviceInfo(cl_device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, - versions_size, dev_versions, NULL), - error, prg); - char compiler_options[1024] = ""; - for (cl_uint i = 0; i < versions_count; ++i) - { - if (opencl_version_is_major(&dev_versions[i], 3)) - { - strcat(compiler_options, "-cl-std=CL3.0 "); - } - } - - if (compiler_options[0] == '\0') - { - fprintf(stderr, "\nError: OpenCL version must be at least 3.0\n"); - exit(EXIT_FAILURE); - } - - OCLERROR_RET(cl_util_build_program(program, cl_device, compiler_options), - error, prg); + // Build OpenCL executable. + OCLERROR_RET( + cl_util_build_program(program, cl_device, NULL /*compiler_options*/), + error, prg); // Query maximum workgroup size (WGS) supported based on private mem // (registers) constraints. diff --git a/samples/extensions/khr/externalmemory/main.cpp b/samples/extensions/khr/externalmemory/main.cpp index ef929819..6bf6de0a 100644 --- a/samples/extensions/khr/externalmemory/main.cpp +++ b/samples/extensions/khr/externalmemory/main.cpp @@ -271,26 +271,8 @@ int main(int argc, char* argv[]) std::istreambuf_iterator{} } }; - // The Khronos extension showcased requires OpenCL 3.0 version. - cl::string compiler_options = ""; - std::vector dev_versions = - cl_device.getInfo(); - for (cl_name_version dev_name_version : dev_versions) - { - if (opencl_version_is_major(dev_name_version, 3)) - { - compiler_options += cl::string{ "-cl-std=CL3.0 " }; - } - } - - if (compiler_options.empty()) - { - std::cerr << "\nError: OpenCL version must be at least 3.0" - << std::endl; - exit(EXIT_FAILURE); - } - - cl_program.build(cl_device, compiler_options.c_str()); + // Build OpenCL executable. + cl_program.build(cl_device); // Query maximum workgroup size (WGS) supported based on private mem // (registers) constraints.