TheRock/rocm-systems/projects/hip-tests/catch/unit/memory/hipFreeAsync.cc

/*
 * Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 *
 * SPDX-License-Identifier: MIT
 */

#include <hip_test_common.hh>
#include <resource_guards.hh>
#include <utils.hh>

/**
 * @addtogroup hipFreeAsync hipFreeAsync
 * @{
 * @ingroup StreamOTest
 * `hipFreeAsync(void* dev_ptr, hipStream_t stream)`
 * - Frees memory with stream ordered semantics
 */


/**
 * Test Description
 * ------------------------
 *  - Test to verify hipFreeAsync behavior with invalid arguments:
 *    -# Nullptr dev_ptr
 *    -# Invalid stream handle
 *    -# Double hipFreeAsync
 *
 * Test source
 * ------------------------
 *  - /unit/memory/hipFreeAsync.cc
 * Test requirements
 * ------------------------
 *  - HIP_VERSION >= 6.0
 */
HIP_TEST_CASE(Unit_hipFreeAsync_Negative_Parameters) {
  int device_id = 0;
  HIP_CHECK(hipSetDevice(device_id));

  int* p = nullptr;
  size_t alloc_size = 1024;
  StreamGuard stream(Streams::created);

  SECTION("dev_ptr is nullptr") {
    HIP_CHECK(hipFreeAsync(nullptr, stream.stream()));
  }

  SECTION("Double free") {
    HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&p), alloc_size, stream.stream()));
    HIP_CHECK(hipStreamSynchronize(stream.stream()));
    HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(p), stream.stream()));
    HIP_CHECK(hipStreamSynchronize(stream.stream()));
    HIP_CHECK_ERROR(hipFreeAsync(reinterpret_cast<void*>(p), stream.stream()),
                    hipErrorInvalidValue);
  }
}

/**
 * End doxygen group StreamOTest.
 * @}
 */

/**
 * Test Description
 * ------------------------
 *  - Functional test cases to trigger capturehipFreeAsync internal api
 * Test source
 * ------------------------
 *  - unit/memory/hipFreeAsync
 * Test requirements
 * ------------------------
 *  - HIP_VERSION >= 6.0
 */
HIP_TEST_CASE(Unit_hipFreeAsync_capturehipFreeAsync) {
  HIP_CHECK(hipSetDevice(0));
  hipGraph_t graph{nullptr};
  hipGraphExec_t graphExec{nullptr};
  hipStream_t stream;
  hipMemPool_t memPool;
  int rows, cols;
  rows = GENERATE(3, 4, 1024);
  cols = GENERATE(3, 4, 1024);
  HIP_CHECK(hipDeviceGetDefaultMemPool(&memPool, 0));
  HIP_CHECK(hipStreamCreate(&stream));
  int* devMem;

  // Start Capturing
  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
  HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&devMem), sizeof(int) * rows * cols,
                                   memPool, stream));
  HIP_CHECK(hipFreeAsync(devMem, stream));
  // End Capture
  HIP_CHECK(hipStreamEndCapture(stream, &graph));

  // Create and Launch Executable Graphs
  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
  HIP_CHECK(hipGraphLaunch(graphExec, stream));
  HIP_CHECK(hipStreamSynchronize(stream));

  HIP_CHECK(hipGraphExecDestroy(graphExec));
  HIP_CHECK(hipGraphDestroy(graph));
  HIP_CHECK(hipStreamDestroy(stream));
}