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/<heap_name>, 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 <gorazd.sumkovski@arm.com>
Signed-off-by: Ahmed Hesham <ahmed.hesham@arm.com>
Co-authored-by: Gorazd Sumkovski <gorazd.sumkovski@arm.com>

test_common/CMakeLists.txt

@@ -1,5 +1,6 @@
 
 set(HARNESS_SOURCES
+    harness/alloc.cpp
     harness/typeWrappers.cpp
     harness/mt19937.cpp
     harness/conversions.cpp

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 <string.h>
+#include <errno.h>
+#include <fcntl.h>
+#include <assert.h>
+#include <stdlib.h>
+#include <unistd.h>
+#include <sys/ioctl.h>
+#include <linux/version.h>
+
+#if LINUX_VERSION_CODE >= KERNEL_VERSION(5, 6, 0)
+#include <linux/dma-heap.h>
+#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__)

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 <stdlib.h>
 #endif
 #endif
+#include <stdint.h>
 
 #if defined(__MINGW32__)
 #include "mingw_compat.h"
 #endif
 
+#if defined(_WIN32)
+#include <cstdlib>
+#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_

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)

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)

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);
+}

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 <numeric>
+
+#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<size_t>(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<uint32_t> 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<cl_mem_properties>(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<uint32_t> 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<uint32_t> 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;
+}