From 8cc65f49c44f5b7cfac0e1cb199f24d027e54d9c Mon Sep 17 00:00:00 2001 From: "systems-assistant[bot]" <221163467+systems-assistant[bot]@users.noreply.github.com> Date: Mon, 27 Oct 2025 15:06:51 +0100 Subject: [PATCH] SWDEV-491296 - Add stream capture testcases to Virtual Memory APIs (#589) --- .../hipMemAddressFree.cc | 35 +++++++++++ .../hipMemAddressReserve.cc | 34 +++++++++++ .../virtualMemoryManagement/hipMemCreate.cc | 31 ++++++++++ .../hipMemExportToShareableHandle.cc | 39 ++++++++++++ .../hipMemGetAllocationGranularity.cc | 28 +++++++++ ...hipMemGetAllocationPropertiesFromHandle.cc | 37 ++++++++++++ .../hipMemImportFromShareableHandle.cc | 37 ++++++++++++ .../unit/virtualMemoryManagement/hipMemMap.cc | 35 +++++++++++ .../virtualMemoryManagement/hipMemRelease.cc | 27 +++++++++ .../hipMemRetainAllocationHandle.cc | 41 +++++++++++++ .../hipMemSetGetAccess.cc | 59 +++++++++++++++++++ .../virtualMemoryManagement/hipMemUnmap.cc | 40 +++++++++++++ 12 files changed, 443 insertions(+) diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc index b7c5733146..45e25a86e2 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc index 495bf9b226..66f0bc3f76 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc index 89f141cc3f..f4edbab94d 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc index f72324160b..299b8bc3a1 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc index e0fbf1c1bf..bb19773f92 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc index ca2cfa52e8..3a1af54bda 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc index e2f8c6cf8c..6af2640862 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc index 92b7f893cd..2253a27b93 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc index 649b330ac1..62ddadbf02 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc @@ -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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc index ba95ea2f51..3468a64079 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc @@ -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(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. * @} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc index 1b83ed3fb7..36ea1f1a4f 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc @@ -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); diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc index 49cc156dc2..c21b8dc17f 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc @@ -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. * @}