Merge pull request #1656 from pbalcer/leaks-program-kernel

fix leaks on level-zero interop program and kernel handles

Author: kbenzie

source/adapters/level_zero/kernel.cpp

@@ -849,6 +849,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelCreateWithNativeHandle(
   try {
     Kernel = new ur_kernel_handle_t_(ZeKernel, Properties->isNativeHandleOwned,
                                      Context);
+    if (Properties->isNativeHandleOwned) {
+      // If ownership is passed to the adapter we need to pass the kernel
+      // to this vector which is then used during ZeKernelRelease.
+      Kernel->ZeKernels.push_back(ZeKernel);
+    }
+
     *RetKernel = reinterpret_cast<ur_kernel_handle_t>(Kernel);
   } catch (const std::bad_alloc &) {
     return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY;

source/adapters/level_zero/program.cpp

@@ -891,10 +891,15 @@ void ur_program_handle_t_::ur_release_program_resources(bool deletion) {
     }
 
     if (ZeModule && OwnZeModule) {
-      for (auto &ZeModulePair : this->ZeModuleMap) {
-        ZE_CALL_NOCHECK(zeModuleDestroy, (ZeModulePair.second));
+      if (ZeModuleMap.empty()) {
+        // interop api
+        ZE_CALL_NOCHECK(zeModuleDestroy, (ZeModule));
+      } else {
+        for (auto &ZeModulePair : this->ZeModuleMap) {
+          ZE_CALL_NOCHECK(zeModuleDestroy, (ZeModulePair.second));
+        }
+        this->ZeModuleMap.clear();
       }
-      this->ZeModuleMap.clear();
     }
     resourcesReleased = true;
   }

test/adapters/CMakeLists.txt

@@ -48,6 +48,35 @@ function(add_adapter_test name)
         ENVIRONMENT "${args_ENVIRONMENT}")
 endfunction()
 
+find_program(VALGRIND valgrind)
+
+function(add_adapter_memcheck_test name)
+    cmake_parse_arguments(args
+        ""                      # options
+        ""                      # one value keywords
+        "ENVIRONMENT"           # multi value keywords
+        ${ARGN})
+    if(VALGRIND)
+        set(target test-adapter-${name})
+        set(test_name ${target}-memcheck)
+
+        add_test(NAME ${test_name}
+                COMMAND ${CMAKE_COMMAND}
+                -D TEST_FILE=valgrind
+                -D TEST_ARGS="--tool=memcheck --leak-check=full $<TARGET_FILE:${target}> --devices_count=${UR_TEST_DEVICES_COUNT} --platforms_count=${UR_TEST_DEVICES_COUNT}"
+                -D MODE=stderr
+                -D MATCH_FILE=${CMAKE_CURRENT_SOURCE_DIR}/${name}_memcheck.match
+                -P ${PROJECT_SOURCE_DIR}/cmake/match.cmake
+                DEPENDS ${TEST_TARGET_NAME}
+                WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
+        )
+
+        set_tests_properties(${test_name} PROPERTIES
+            LABELS "adapter-specific;${name}"
+            ENVIRONMENT "${args_ENVIRONMENT}")
+    endif()
+endfunction()
+
 if(UR_BUILD_ADAPTER_CUDA OR UR_BUILD_ADAPTER_ALL)
     add_subdirectory(cuda)
 endif()

test/adapters/level_zero/CMakeLists.txt

@@ -13,13 +13,26 @@ else()
         FIXTURE KERNELS
         SOURCES
             urProgramLink.cpp
+            urKernelCreateWithNativeHandle.cpp
         ENVIRONMENT
             "UR_ADAPTERS_FORCE_LOAD=\"$<TARGET_FILE:ur_adapter_level_zero>\""
     )
+    # TODO: valgrind tests require very new environment.
+    # Enable once all L0 runners are updated.
+    # add_adapter_memcheck_test(level_zero
+    #    ENVIRONMENT
+    #         "UR_ADAPTERS_FORCE_LOAD=\"$<TARGET_FILE:ur_adapter_level_zero>\""
+    # )
+
+    target_link_libraries(test-adapter-level_zero PRIVATE
+        LevelZeroLoader
+        LevelZeroLoader-Headers
+    )
 
     target_include_directories(test-adapter-level_zero PRIVATE
         ${PROJECT_SOURCE_DIR}/source
         ${PROJECT_SOURCE_DIR}/source/adapters/level_zero
+        LevelZeroLoader-Headers
     )
 
     add_dependencies(test-adapter-level_zero

test/adapters/level_zero/level_zero_memcheck.match

@@ -0,0 +1,2 @@
+{{IGNORE}}
+{{.*}} ERROR SUMMARY: 0 errors from 0 contexts {{.*}}

test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp

@@ -0,0 +1,63 @@
+// Copyright (C) 2024 Intel Corporation
+// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+// See LICENSE.TXT
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "level_zero/ze_api.h"
+#include "ur_api.h"
+#include "uur/checks.h"
+#include <uur/fixtures.h>
+
+using urLevelZeroKernelNativeHandleTest = uur::urContextTest;
+UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urLevelZeroKernelNativeHandleTest);
+
+TEST_P(urLevelZeroKernelNativeHandleTest, OwnedHandleRelease) {
+    ze_context_handle_t native_context;
+    urContextGetNativeHandle(context, (ur_native_handle_t *)&native_context);
+
+    ze_device_handle_t native_device;
+    urDeviceGetNativeHandle(device, (ur_native_handle_t *)&native_device);
+
+    std::shared_ptr<std::vector<char>> il_binary;
+    uur::KernelsEnvironment::instance->LoadSource("foo", il_binary);
+
+    auto kernel_name =
+        uur::KernelsEnvironment::instance->GetEntryPointNames("foo")[0];
+
+    ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC};
+    moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
+    moduleDesc.inputSize = il_binary->size();
+    moduleDesc.pInputModule =
+        reinterpret_cast<const uint8_t *>(il_binary->data());
+    moduleDesc.pBuildFlags = "";
+    ze_module_handle_t module;
+
+    ASSERT_EQ(zeModuleCreate(native_context, native_device, &moduleDesc,
+                             &module, NULL),
+              ZE_RESULT_SUCCESS);
+
+    ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
+    kernelDesc.pKernelName = kernel_name.c_str();
+
+    ze_kernel_handle_t native_kernel;
+
+    ASSERT_EQ(zeKernelCreate(module, &kernelDesc, &native_kernel),
+              ZE_RESULT_SUCCESS);
+
+    ur_program_native_properties_t pprops = {
+        UR_STRUCTURE_TYPE_PROGRAM_NATIVE_PROPERTIES, nullptr, 1};
+
+    ur_program_handle_t program;
+    ASSERT_SUCCESS(urProgramCreateWithNativeHandle((ur_native_handle_t)module,
+                                                   context, &pprops, &program));
+
+    ur_kernel_native_properties_t kprops = {
+        UR_STRUCTURE_TYPE_KERNEL_NATIVE_PROPERTIES, nullptr, 1};
+
+    ur_kernel_handle_t kernel;
+    ASSERT_SUCCESS(urKernelCreateWithNativeHandle(
+        (ur_native_handle_t)native_kernel, context, program, &kprops, &kernel));
+
+    ASSERT_SUCCESS(urKernelRelease(kernel));
+    ASSERT_SUCCESS(urProgramRelease(program));
+}

test/adapters/level_zero/urProgramLink.cpp

@@ -3,6 +3,7 @@
 // See LICENSE.TXT
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+#include "ur_api.h"
 #include <uur/fixtures.h>
 
 using urLevelZeroProgramLinkTest = uur::urProgramTest;
@@ -28,4 +29,6 @@ TEST_P(urLevelZeroProgramLinkTest, InvalidLinkOptionsPrintedInLog) {
                                          log.data(), nullptr));
     ASSERT_EQ(log[logSize - 1], '\0');
     ASSERT_NE(std::string{log.data()}.find("-foo"), std::string::npos);
+
+    ASSERT_SUCCESS(urProgramRelease(linked_program));
 }