SWDEV-491296 - Add stream capture testcases to Virtual Memory APIs (#589)

2025-10-27 15:06:51 +01:00
@@ -75,6 +75,41 @@ TEST_CASE("Unit_hipMemAddressFree_negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemAddressFree_Capture") {
+  CTX_CREATE();
+  size_t granularity = 0;
+  size_t buffer_size = DATA_SIZE * sizeof(int);
+  int device_id = 0;
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, device_id));
+  checkVMMSupported(device);
+
+  hipMemAllocationProp alloc_prop{};
+  alloc_prop.type = hipMemAllocationTypePinned;
+  alloc_prop.location.type = hipMemLocationTypeDevice;
+  alloc_prop.location.id = device;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &alloc_prop,
+                                           hipMemAllocationGranularityMinimum));
+  REQUIRE(granularity > 0);
+
+  size_t reserved_size = ((granularity + buffer_size - 1) / granularity) * granularity;
+
+  hipDeviceptr_t reserved_ptr = nullptr;
+  HIP_CHECK(hipMemAddressReserve(&reserved_ptr, reserved_size, 0, nullptr, 0));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemAddressFree(reserved_ptr, reserved_size));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  CTX_DESTROY();
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -152,6 +152,40 @@ TEST_CASE("Unit_hipMemAddressReserve_Negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemAddressReserve_Capture") {
+  hipMemGenericAllocationHandle_t allocation_handle;
+  size_t granularity = 0;
+  constexpr size_t kAlignment = 2;
+  constexpr int kDeviceId = 0;
+  hipDevice_t device = 0;
+  hipDeviceptr_t device_ptr = nullptr;
+
+  CTX_CREATE();
+  HIP_CHECK(hipDeviceGet(&device, kDeviceId));
+
+  hipMemAllocationProp allocation_prop{};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+  HIP_CHECK(hipMemCreate(&allocation_handle, granularity, &allocation_prop, 0));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemAddressReserve(&device_ptr, granularity, kAlignment, nullptr, 0));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemAddressFree(device_ptr, granularity));
+  HIP_CHECK(hipMemRelease(allocation_handle));
+  CTX_DESTROY();
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -522,6 +522,37 @@ TEST_CASE("Unit_hipMemCreate_Negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemCreate_Capture") {
+  CTX_CREATE();
+
+  hipMemGenericAllocationHandle_t allocation_handle;
+  size_t allocation_granularity = 0;
+  constexpr int kDeviceId = 0;
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, kDeviceId));
+
+  hipMemAllocationProp allocation_prop{};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&allocation_granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemCreate(&allocation_handle, allocation_granularity, &allocation_prop, 0));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemRelease(allocation_handle));
+  CTX_DESTROY();
+}
+
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -133,6 +133,45 @@ TEST_CASE("Unit_hipMemExportToShareableHandle_Negative_Parameters") {
  HIP_CHECK(hipMemRelease(handle));
 }

+TEST_CASE("Unit_hipMemExportToShareableHandle_Capture") {
+  CTX_CREATE();
+
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  checkVMMSupported(device);
+
+  hipMemAllocationProp allocation_prop = {};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  size_t allocation_granularity;
+  HIP_CHECK(hipMemGetAllocationGranularity(&allocation_granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+
+  hipMemGenericAllocationHandle_t allocation_handle;
+  HIP_CHECK(hipMemCreate(&allocation_handle, allocation_granularity * 2, &allocation_prop, 0));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+
+  void* shareable_handle = nullptr;
+  HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, allocation_handle,
+                                          hipMemHandleTypePosixFileDescriptor, 0));
+
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  REQUIRE(shareable_handle != nullptr);
+  HIP_CHECK(hipMemRelease(allocation_handle));
+
+  CTX_DESTROY();
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -176,6 +176,34 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_NegativeTests") {
 #endif
 }

+TEST_CASE("Unit_hipMemGetAllocationGranularity_Capture") {
+  CTX_CREATE();
+
+  size_t granularity = 0;
+  constexpr int kDeviceId = 0;
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, kDeviceId));
+
+  hipMemAllocationProp allocation_prop{};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;  // Current Device
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+
+  CTX_DESTROY();
+}
+
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -121,6 +121,43 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Capture") {
+  CTX_CREATE();
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  checkVMMSupported(device);
+
+  hipMemGenericAllocationHandle_t handle;
+  hipMemAllocationProp allocation_prop = {};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  hipMemAllocationProp allocation_prop_temp = {};
+  size_t granularity = 0;
+  size_t buffer_size = DATA_SIZE * sizeof(int);
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+  REQUIRE(granularity > 0);
+
+  size_t mem_size = ((granularity + buffer_size - 1) / granularity) * granularity;
+
+  HIP_CHECK(hipMemCreate(&handle, mem_size, &allocation_prop, 0));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemGetAllocationPropertiesFromHandle(&allocation_prop_temp, handle));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemRelease(handle));
+  CTX_DESTROY();
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -547,6 +547,43 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") {
    REQUIRE(close(fdpid[0]) == 0);
  }
 }
+
+TEST_CASE("Unit_hipMemImportFromShareableHandle_Capture") {
+  CTX_CREATE();
+
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, 0));
+  checkVMMSupported(device);
+
+  hipMemAllocationProp allocation_prop = {};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  size_t granularity = 0;
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+
+  hipMemGenericAllocationHandle_t allocation_handle;
+  HIP_CHECK(hipMemCreate(&allocation_handle, granularity * 2, &allocation_prop, 0));
+  ShareableHandle shareable_handle;
+  HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, allocation_handle,
+                                          hipMemHandleTypePosixFileDescriptor, 0));
+  hipMemGenericAllocationHandle_t imported_handle;
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, &shareable_handle,
+                                            hipMemHandleTypePosixFileDescriptor));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemRelease(allocation_handle));
+  CTX_DESTROY();
+}
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -684,6 +684,41 @@ TEST_CASE("Unit_hipMemMap_negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemMap_Capture") {
+  hipMemGenericAllocationHandle_t handle;
+  size_t granularity = 0;
+  constexpr size_t kAlignment = 2;
+  constexpr int kDeviceId = 0;
+  hipDevice_t device = 0;
+  hipDeviceptr_t device_ptr = nullptr;
+
+  CTX_CREATE();
+  HIP_CHECK(hipDeviceGet(&device, kDeviceId));
+
+  hipMemAllocationProp prop{};
+  prop.type = hipMemAllocationTypePinned;
+  prop.location.type = hipMemLocationTypeDevice;
+  prop.location.id = device;
+
+  HIP_CHECK(
+      hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
+  HIP_CHECK(hipMemCreate(&handle, granularity, &prop, 0));
+  HIP_CHECK(hipMemAddressReserve(&device_ptr, granularity, kAlignment, 0, 0));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemMap(device_ptr, granularity, 0, handle, 0));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemUnmap(device_ptr, granularity));
+  HIP_CHECK(hipMemRelease(handle));
+  HIP_CHECK(hipMemAddressFree(device_ptr, granularity));
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -49,6 +49,33 @@ TEST_CASE("Unit_hipMemRelease_negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemRelease_Capture") {
+  CTX_CREATE();
+
+  hipMemGenericAllocationHandle_t allocation_handle;
+  size_t allocation_granularity = 0;
+  int device_id = 0;
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, device_id));
+
+  hipMemAllocationProp allocation_prop{};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&allocation_granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+  HIP_CHECK(hipMemCreate(&allocation_handle, allocation_granularity, &allocation_prop, 0));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemRelease(allocation_handle));
+  END_CAPTURE(stream);
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -140,6 +140,47 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") {
  HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
 }

+TEST_CASE("Unit_hipMemRetainAllocationHandle_Capture") {
+  CTX_CREATE();
+  size_t granularity = 0;
+  size_t buffer_size = DATA_SIZE * sizeof(int);
+  int device_id = 0;
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, device_id));
+  checkVMMSupported(device);
+
+  hipMemAllocationProp allocation_prop{};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+  REQUIRE(granularity > 0);
+
+  size_t allocation_size = ((granularity + buffer_size - 1) / granularity) * granularity;
+  hipMemGenericAllocationHandle_t allocation_handle;
+  hipDeviceptr_t device_ptr;
+  HIP_CHECK(hipMemCreate(&allocation_handle, allocation_size, &allocation_prop, 0));
+  HIP_CHECK(hipMemAddressReserve(&device_ptr, allocation_size, 0, 0, 0));
+  HIP_CHECK(hipMemMap(device_ptr, allocation_size, 0, allocation_handle, 0));
+
+  hipMemGenericAllocationHandle_t retained_handle;
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemRetainAllocationHandle(&retained_handle, reinterpret_cast<void*>(device_ptr)));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemRelease(allocation_handle));
+  HIP_CHECK(hipMemUnmap(device_ptr, allocation_size));
+  HIP_CHECK(hipMemAddressFree(device_ptr, allocation_size));
+  CTX_DESTROY();
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}
@@ -1329,6 +1329,65 @@ TEST_CASE("Unit_hipMemSetAccess_negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemSetGetAccess_Capture") {
+  CTX_CREATE();
+
+  const size_t kBufferBytes = DATA_SIZE * sizeof(int);
+  const int kDeviceId = 0;
+  hipDevice_t device;
+  HIP_CHECK(hipDeviceGet(&device, kDeviceId));
+  checkVMMSupported(device);
+
+  hipMemAllocationProp alloc_prop{};
+  alloc_prop.type = hipMemAllocationTypePinned;
+  alloc_prop.location.type = hipMemLocationTypeDevice;
+  alloc_prop.location.id = device;
+
+  size_t granularity = 0;
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &alloc_prop,
+                                           hipMemAllocationGranularityMinimum));
+  REQUIRE(granularity > 0);
+
+  const size_t vmm_bytes = ((granularity + kBufferBytes - 1) / granularity) * granularity;
+
+  hipMemGenericAllocationHandle_t mem_handle;
+  HIP_CHECK(hipMemCreate(&mem_handle, vmm_bytes, &alloc_prop, 0));
+
+  hipDeviceptr_t vmm_ptr = nullptr;
+  HIP_CHECK(hipMemAddressReserve(&vmm_ptr, vmm_bytes, 0, 0, 0));
+  HIP_CHECK(hipMemMap(vmm_ptr, vmm_bytes, 0, mem_handle, 0));
+  HIP_CHECK(hipMemRelease(mem_handle));
+
+  hipMemAccessDesc access_desc{};
+  access_desc.location.type = hipMemLocationTypeDevice;
+  access_desc.location.id = device;
+  access_desc.flags = hipMemAccessFlagsProtReadWrite;
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+
+  // Test hipMemSetAccess inside stream capture
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemSetAccess(vmm_ptr, vmm_bytes, &access_desc, 1));
+  END_CAPTURE(stream);
+
+  // Test hipMemGetAccess inside stream capture
+  BEGIN_CAPTURE(stream);
+  hipMemLocation mem_location{};
+  mem_location.type = hipMemLocationTypeDevice;
+  mem_location.id = device;
+  unsigned long long access_flags = 0;
+  HIP_CHECK(hipMemGetAccess(&access_flags, &mem_location, vmm_ptr));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemUnmap(vmm_ptr, vmm_bytes));
+  HIP_CHECK(hipMemAddressFree(vmm_ptr, vmm_bytes));
+  CTX_DESTROY();
+}
+
 TEST_CASE("Unit_hipMemSetAccessHostDevice_hostalloc") {
  // Ensure device 0 is selected
  REQUIRE(hipSetDevice(0) == hipSuccess);
@@ -87,6 +87,46 @@ TEST_CASE("Unit_hipMemUnmap_negative") {
  CTX_DESTROY();
 }

+TEST_CASE("Unit_hipMemUnmap_Capture") {
+  CTX_CREATE();
+  size_t granularity = 0;
+  constexpr size_t kBufferSize = N * sizeof(int);
+  int device_id = 0;
+  hipDevice_t device;
+
+  HIP_CHECK(hipDeviceGet(&device, device_id));
+  checkVMMSupported(device);
+
+  hipMemAllocationProp allocation_prop{};
+  allocation_prop.type = hipMemAllocationTypePinned;
+  allocation_prop.location.type = hipMemLocationTypeDevice;
+  allocation_prop.location.id = device;
+
+  HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &allocation_prop,
+                                           hipMemAllocationGranularityMinimum));
+  REQUIRE(granularity > 0);
+  size_t mem_size = ((granularity + kBufferSize - 1) / granularity) * granularity;
+
+  hipMemGenericAllocationHandle_t allocation_handle;
+  hipDeviceptr_t device_ptr;
+  HIP_CHECK(hipMemCreate(&allocation_handle, mem_size, &allocation_prop, 0));
+  HIP_CHECK(hipMemAddressReserve(&device_ptr, mem_size, 0, nullptr, 0));
+  HIP_CHECK(hipMemMap(device_ptr, mem_size, 0, allocation_handle, 0));
+  HIP_CHECK(hipMemRelease(allocation_handle));
+
+  hipStream_t stream = nullptr;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  GENERATE_CAPTURE();
+  BEGIN_CAPTURE(stream);
+  HIP_CHECK(hipMemUnmap(device_ptr, mem_size));
+  END_CAPTURE(stream);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipMemAddressFree(device_ptr, mem_size));
+  CTX_DESTROY();
+}
+
 /**
 * End doxygen group VirtualMemoryManagementTest.
 * @}