From 9ba6f062d4f7acc01df515beaad8309cc3b59af0 Mon Sep 17 00:00:00 2001 From: Ahmed Hesham <117350656+ahesham-arm@users.noreply.github.com> Date: Wed, 26 Feb 2025 17:51:22 +0000 Subject: [PATCH] Add support for allocating DMA buffers (#2170) This adds support for allocating DMA buffers on systems that support it, i.e. Linux and Android. On mainline Linux, starting version 5.6 (equivalent to Android 12), there is a new kernel module framework available called [DMA-BUF Heaps](https://github.com/torvalds/linux/blob/master/drivers/dma-buf/dma-heap.c). The goal of this framework is to provide a standardised way for user applications to allocate and share memory buffers between different devices, subsystems, etc. The main feature of interest is that the framework provides device-agnostic allocation; it abstracts away the underlying hardware, and provides a single IOCTL, `DMA_HEAP_IOCTL_ALLOC`. Mainline implementation provides two heaps that act as character devices that can allocate DMA buffers; system, which uses the buddy allocator, and cma, which uses the [CMA](https://developer.toradex.com/software/linux-resources/linux-features/contiguous-memory-allocator-cma-linux/) (Contiguous Memory Allocator). Both of these are [kernel configuration options](https://github.com/torvalds/linux/blob/master/drivers/dma-buf/heaps/Kconfig) that need to be enabled when building the Linux kernel. Generally, any kernel module implementing this framework is made available under /dev/dma_heaps/, e.g. /dev/dma_heaps/system. The implementation currently only supports one type of DMA heaps; `system`, the default device path for which is `/dev/dma_heap/system`. The path can be overridden at runtime using an environment variable, `OCL_CTS_DMA_HEAP_PATH_SYSTEM`, if needed. Extending this in the future should be trivial (subject to platform support), by adding an entry to the enum `dma_buf_heap_type`, and an appropriate default path and overriding environment variable name. The proposed implementation will conditionally compile if the conditions are met (i.e. building for Linux or Android, using kernel headers >= 5.6.0), and will provide a compile-time warning otherwise, and return `-1` as the DMA handle in runtime if not. To demonstrate the functionality, a new test is added for the `cl_khr_external_memory_dma_buf` extension. If the extension is supported by the device, a DMA buffer will be allocated and used to create a CL buffer, that is then used by a simple kernel. This should provide a way forward for adding more tests that depend on DMA buffers. --------- Signed-off-by: Gorazd Sumkovski Signed-off-by: Ahmed Hesham Co-authored-by: Gorazd Sumkovski --- test_common/CMakeLists.txt | 1 + test_common/harness/alloc.cpp | 123 +++++++++++++++ test_common/harness/alloc.h | 39 ++++- test_conformance/extensions/CMakeLists.txt | 1 + .../CMakeLists.txt | 8 + .../cl_khr_external_memory_dma_buf/main.cpp | 23 +++ .../test_external_memory_dma_buf.cpp | 143 ++++++++++++++++++ 7 files changed, 337 insertions(+), 1 deletion(-) create mode 100644 test_common/harness/alloc.cpp create mode 100644 test_conformance/extensions/cl_khr_external_memory_dma_buf/CMakeLists.txt create mode 100644 test_conformance/extensions/cl_khr_external_memory_dma_buf/main.cpp create mode 100644 test_conformance/extensions/cl_khr_external_memory_dma_buf/test_external_memory_dma_buf.cpp diff --git a/test_common/CMakeLists.txt b/test_common/CMakeLists.txt index b0505345..3acc742c 100644 --- a/test_common/CMakeLists.txt +++ b/test_common/CMakeLists.txt @@ -1,5 +1,6 @@ set(HARNESS_SOURCES + harness/alloc.cpp harness/typeWrappers.cpp harness/mt19937.cpp harness/conversions.cpp diff --git a/test_common/harness/alloc.cpp b/test_common/harness/alloc.cpp new file mode 100644 index 00000000..685ff272 --- /dev/null +++ b/test_common/harness/alloc.cpp @@ -0,0 +1,123 @@ +// +// Copyright (c) 2024 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. +// + +#include "alloc.h" +#include "errorHelpers.h" +#include "testHarness.h" + +#if defined(linux) || defined(__linux__) || defined(__ANDROID__) +#include +#include +#include +#include +#include +#include +#include +#include + +#if LINUX_VERSION_CODE >= KERNEL_VERSION(5, 6, 0) +#include +#endif + +struct dma_buf_heap_helper_t +{ + dma_buf_heap_type heap_type; + const char* env_var = nullptr; + const char* default_path = nullptr; + + constexpr dma_buf_heap_helper_t(dma_buf_heap_type heap_type, + const char* env_var, + const char* default_path) + : heap_type(heap_type), env_var(env_var), default_path(default_path) + {} +}; + +constexpr dma_buf_heap_helper_t DMA_BUF_HEAP_TABLE[] = { + { dma_buf_heap_type::SYSTEM, "OCL_CTS_DMA_HEAP_PATH_SYSTEM", + "/dev/dma_heap/system" }, +}; + +static dma_buf_heap_helper_t lookup_dma_heap(dma_buf_heap_type heap_type) +{ + for (const auto& entry : DMA_BUF_HEAP_TABLE) + { + if (heap_type == entry.heap_type) + { + return entry; + } + } + + assert(false + && "DMA heap type does not have an entry in DMA_BUF_HEAP_TABLE"); + return DMA_BUF_HEAP_TABLE[0]; +} + +int allocate_dma_buf(uint64_t size, dma_buf_heap_type heap_type) +{ +#if LINUX_VERSION_CODE >= KERNEL_VERSION(5, 6, 0) + constexpr int DMA_HEAP_FLAGS = O_RDWR | O_CLOEXEC; + + const auto entry = lookup_dma_heap(heap_type); + const auto override_path = getenv(entry.env_var); + const auto dma_heap_path = + (override_path == nullptr) ? entry.default_path : override_path; + + const int dma_heap_fd = open(dma_heap_path, DMA_HEAP_FLAGS); + if (dma_heap_fd == -1) + { + log_error( + "Opening the DMA heap device: %s failed with error: %d (%s)\n", + dma_heap_path, errno, strerror(errno)); + + return TEST_SKIPPED_ITSELF; + } + + dma_heap_allocation_data dma_heap_data = { 0 }; + dma_heap_data.len = size; + dma_heap_data.fd_flags = O_RDWR | O_CLOEXEC; + + int result = ioctl(dma_heap_fd, DMA_HEAP_IOCTL_ALLOC, &dma_heap_data); + if (result != 0) + { + log_error("DMA heap allocation IOCTL call failed, error: %d\n", result); + + close(dma_heap_fd); + return -1; + } + + result = close(dma_heap_fd); + if (result == -1) + { + log_info("Failed to close the DMA heap device: %s\n", dma_heap_path); + } + + return dma_heap_data.fd; +#else +#warning \ + "Kernel version doesn't support DMA buffer heaps (at least v5.6.0 is required)." + return TEST_SKIPPED_ITSELF; +#endif // LINUX_VERSION_CODE >= KERNEL_VERSION(5, 6, 0) +} + +#else +int allocate_dma_buf(uint64_t size, dma_buf_heap_type heap_type) +{ + log_error( + "OS doesn't have DMA buffer heaps (only Linux and Android do).\n"); + + return TEST_SKIPPED_ITSELF; +} +#endif // defined(linux) || defined(__linux__) || defined(__ANDROID__) diff --git a/test_common/harness/alloc.h b/test_common/harness/alloc.h index 3b00d7c9..49093c71 100644 --- a/test_common/harness/alloc.h +++ b/test_common/harness/alloc.h @@ -1,5 +1,5 @@ // -// Copyright (c) 2020 The Khronos Group Inc. +// Copyright (c) 2020 - 2024 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. @@ -24,11 +24,16 @@ #include #endif #endif +#include #if defined(__MINGW32__) #include "mingw_compat.h" #endif +#if defined(_WIN32) +#include +#endif + inline void* align_malloc(size_t size, size_t alignment) { #if defined(_WIN32) && defined(_MSC_VER) @@ -66,4 +71,36 @@ inline void align_free(void* ptr) #endif } +enum class dma_buf_heap_type +{ + SYSTEM +}; + +/** + * @brief Allocate a DMA buffer. + * + * On systems that support it, use the DMA buffer heaps to allocate a DMA buffer + * of the requested size, using the requested heap type. The heap type defaults + * to using the system heap if no type is specified. + * + * A heap type will use a default path if one exists, and can be overriden using + * an environment variable for each type, as follows: + * + * SYSTEM: + * * Default path: /dev/dma_heap/system + * * Environment variable: OCL_CTS_DMA_HEAP_PATH_SYSTEM + * + * DMA buffer heaps require a minimum Linux kernel version 5.6. A compile-time + * warning is issued on older systems, as well as an error message at runtime. + * + * @param size [in] The requested buffer size in bytes. + * @param heap_type [in,opt] The heap type to use for the allocation. + * + * @retrun A file descriptor representing the allocated DMA buffer on success, + * -1 otherwise. Failure to open the DMA device returns TEST_SKIPPED_ITSELF so + * it can be handled separately to other failures. + */ +int allocate_dma_buf(uint64_t size, + dma_buf_heap_type heap_type = dma_buf_heap_type::SYSTEM); + #endif // #ifndef HARNESS_ALLOC_H_ diff --git a/test_conformance/extensions/CMakeLists.txt b/test_conformance/extensions/CMakeLists.txt index e0e790c2..d064e8a9 100644 --- a/test_conformance/extensions/CMakeLists.txt +++ b/test_conformance/extensions/CMakeLists.txt @@ -5,6 +5,7 @@ add_subdirectory( cl_ext_cxx_for_opencl ) add_subdirectory( cl_khr_command_buffer ) add_subdirectory( cl_khr_dx9_media_sharing ) +add_subdirectory( cl_khr_external_memory_dma_buf ) add_subdirectory( cl_khr_semaphore ) add_subdirectory( cl_khr_kernel_clock ) if(VULKAN_IS_SUPPORTED) diff --git a/test_conformance/extensions/cl_khr_external_memory_dma_buf/CMakeLists.txt b/test_conformance/extensions/cl_khr_external_memory_dma_buf/CMakeLists.txt new file mode 100644 index 00000000..c5a45a6c --- /dev/null +++ b/test_conformance/extensions/cl_khr_external_memory_dma_buf/CMakeLists.txt @@ -0,0 +1,8 @@ +set(MODULE_NAME CL_KHR_EXTERNAL_MEMORY_DMA_BUF) + +set(${MODULE_NAME}_SOURCES + main.cpp + test_external_memory_dma_buf.cpp +) + +include(../../CMakeCommon.txt) diff --git a/test_conformance/extensions/cl_khr_external_memory_dma_buf/main.cpp b/test_conformance/extensions/cl_khr_external_memory_dma_buf/main.cpp new file mode 100644 index 00000000..b0f8a718 --- /dev/null +++ b/test_conformance/extensions/cl_khr_external_memory_dma_buf/main.cpp @@ -0,0 +1,23 @@ +// +// Copyright (c) 2024 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. +// + +#include "harness/testHarness.h" + +int main(int argc, const char *argv[]) +{ + return runTestHarness(argc, argv, test_registry::getInstance().num_tests(), + test_registry::getInstance().definitions(), false, 0); +} diff --git a/test_conformance/extensions/cl_khr_external_memory_dma_buf/test_external_memory_dma_buf.cpp b/test_conformance/extensions/cl_khr_external_memory_dma_buf/test_external_memory_dma_buf.cpp new file mode 100644 index 00000000..dc158e30 --- /dev/null +++ b/test_conformance/extensions/cl_khr_external_memory_dma_buf/test_external_memory_dma_buf.cpp @@ -0,0 +1,143 @@ +// +// Copyright (c) 2024 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. +// + +#include + +#include "harness/typeWrappers.h" +#include "harness/testHarness.h" + +static const char* kernel_function_inc_buffer = R"( + kernel void inc_buffer(global uint *src, global uint *imp, global uint *dst) + { + uint global_id = get_global_id(0); + + imp[global_id] = src[global_id] + 1; + dst[global_id] = imp[global_id] + 1; + } + )"; + +/** + * Demonstrate the functionality of the cl_khr_external_memory_dma_buf extension + * by creating an imported buffer from a DMA buffer, then writing into, and + * reading from it. + */ + +REGISTER_TEST(external_memory_dma_buf) +{ + if (!is_extension_available(device, "cl_khr_external_memory_dma_buf")) + { + log_info("The device does not support the " + "cl_khr_external_memory_dma_buf extension.\n"); + + return TEST_SKIPPED_ITSELF; + } + + const size_t buffer_size = static_cast(num_elements); + const size_t buffer_size_bytes = sizeof(uint32_t) * buffer_size; + + clProgramWrapper program; + clKernelWrapper kernel; + cl_int error; + + error = + create_single_kernel_helper(context, &program, &kernel, 1, + &kernel_function_inc_buffer, "inc_buffer"); + test_error(error, "Failed to create program with source."); + + /* Source buffer initialisation */ + std::vector src_data(buffer_size); + // Arithmetic progression starting at 0 and incrementing by 1 + std::iota(std::begin(src_data), std::end(src_data), 0); + + clMemWrapper src_buffer = + clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + buffer_size_bytes, src_data.data(), &error); + test_error(error, "Failed to create the source buffer."); + + /* Imported buffer creation */ + int dma_buf_fd = allocate_dma_buf(buffer_size_bytes); + if (dma_buf_fd < 0) + { + if (dma_buf_fd == TEST_SKIPPED_ITSELF) + { + return TEST_SKIPPED_ITSELF; + } + + log_error( + "Failed to obtain a valid DMA buffer file descriptor, got %i.\n", + dma_buf_fd); + + return TEST_FAIL; + } + + const cl_mem_properties ext_mem_properties[] = { + CL_EXTERNAL_MEMORY_HANDLE_DMA_BUF_KHR, + static_cast(dma_buf_fd), CL_PROPERTIES_LIST_END_EXT + }; + + clMemWrapper imp_buffer = clCreateBufferWithProperties( + context, ext_mem_properties, CL_MEM_READ_WRITE, buffer_size_bytes, + nullptr, &error); + test_error(error, "Failed to create the imported buffer."); + + /* Destination buffer creation */ + clMemWrapper dst_buffer = clCreateBuffer( + context, CL_MEM_WRITE_ONLY, buffer_size_bytes, nullptr, &error); + test_error(error, "Failed to create the destination buffer."); + + /* Kernel arguments setup */ + error = clSetKernelArg(kernel, 0, sizeof(src_buffer), &src_buffer); + test_error(error, "Failed to set kernel argument 0 to src_buffer."); + + error = clSetKernelArg(kernel, 1, sizeof(imp_buffer), &imp_buffer); + test_error(error, "Failed to set kernel argument 1 to imp_buffer."); + + error = clSetKernelArg(kernel, 2, sizeof(dst_buffer), &dst_buffer); + test_error(error, "Failed to set kernel argument 2 to dst_buffer."); + + /* Kernel execution */ + error = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &buffer_size, + nullptr, 0, nullptr, nullptr); + test_error(error, "Failed to enqueue the kernel."); + + error = clFinish(queue); + test_error(error, "Failed to finish the queue."); + + /* Verification */ + std::vector dst_data(buffer_size, 0); + + error = clEnqueueReadBuffer(queue, dst_buffer, CL_BLOCKING, 0, + buffer_size_bytes, dst_data.data(), 0, nullptr, + nullptr); + test_error(error, "Failed to read the contents of the destination buffer."); + + std::vector expected_data(buffer_size); + std::iota(std::begin(expected_data), std::end(expected_data), 2); + + for (size_t i = 0; i < buffer_size; ++i) + { + if (dst_data[i] != expected_data[i]) + { + log_error( + "Verification failed at index %zu, expected %u but got %u\n", i, + expected_data[i], dst_data[i]); + + return TEST_FAIL; + } + } + + return TEST_PASS; +}